Update .github/release-debian-v3

.github/build-canary-i386

@@ -0,0 +1 @@
+1

.github/build-canary-v3

@@ -0,0 +1 @@
+1

.github/build-nest-i386

@@ -0,0 +1 @@
+1

.github/build-nest-v3

@@ -0,0 +1 @@
+1

.github/release-canary-i386

@@ -0,0 +1 @@
+2

.github/release-canary-v3

@@ -0,0 +1 @@
+4

.github/release-debian-v3

@@ -0,0 +1 @@
+2

.github/release-nest-i386

@@ -0,0 +1 @@
+1

.github/release-nest-v3

@@ -0,0 +1 @@
+1

.github/workflows/build-canaryi386.yml

@@ -0,0 +1,37 @@
+name: PikaOS Package Build Only (Canary) (i386)
+
+on:
+  push:
+    branches:
+      - main
+    paths:
+      - '.github/build-canary-i386'
+
+jobs:
+  build:
+    runs-on: ubuntu-latest
+    container:
+      image: ghcr.io/pikaos-linux/pikaos-builder:canaryi386
+      volumes:
+        - /proc:/proc
+      options: --privileged -it
+
+    steps:
+    - uses: actions/checkout@v3
+
+    - name: Install SSH key
+      uses: shimataro/ssh-key-action@v2
+      with:
+        key: ${{ secrets.SSH_KEY }}
+        name: id_rsa
+        known_hosts: ${{ secrets.KNOWN_HOSTS }}
+        if_key_exists: replace
+    
+    - name: Update APT Cache
+      run: apt-get update -y
+
+    - name: Set Build Config
+      run: cp -vf ./pika-build-config/i386.sh ./pika-build-config.sh
+
+    - name: Build Package
+      run: ./main.sh

.github/workflows/build-canaryv3.yml

@@ -0,0 +1,37 @@
+name: PikaOS Package Build Only (Canary) (amd64-v3)
+
+on:
+  push:
+    branches:
+      - main
+    paths:
+      - '.github/build-canary-v3'
+
+jobs:
+  build:
+    runs-on: ubuntu-latest
+    container:
+      image: ghcr.io/pikaos-linux/pikaos-builder:canaryv3
+      volumes:
+        - /proc:/proc
+      options: --privileged -it
+
+    steps:
+    - uses: actions/checkout@v3
+
+    - name: Install SSH key
+      uses: shimataro/ssh-key-action@v2
+      with:
+        key: ${{ secrets.SSH_KEY }}
+        name: id_rsa
+        known_hosts: ${{ secrets.KNOWN_HOSTS }}
+        if_key_exists: replace
+    
+    - name: Update APT Cache
+      run: apt-get update -y
+
+    - name: Set Build Config
+      run: cp -vf ./pika-build-config/amd64-v3.sh ./pika-build-config.sh
+
+    - name: Build Package
+      run: ./main.sh

.github/workflows/build-nesti386.yml

@@ -0,0 +1,37 @@
+name: PikaOS Package Build Only (i386)
+
+on:
+  push:
+    branches:
+      - main
+    paths:
+      - '.github/build-nest-i386'
+
+jobs:
+  build:
+    runs-on: ubuntu-latest
+    container:
+      image: ghcr.io/pikaos-linux/pikaos-builder:nesti386
+      volumes:
+        - /proc:/proc
+      options: --privileged -it
+
+    steps:
+    - uses: actions/checkout@v3
+
+    - name: Install SSH key
+      uses: shimataro/ssh-key-action@v2
+      with:
+        key: ${{ secrets.SSH_KEY }}
+        name: id_rsa
+        known_hosts: ${{ secrets.KNOWN_HOSTS }}
+        if_key_exists: replace
+    
+    - name: Update APT Cache
+      run: apt-get update -y
+
+    - name: Set Build Config
+      run: cp -vf ./pika-build-config/i386.sh ./pika-build-config.sh
+
+    - name: Build Package
+      run: ./main.sh

.github/workflows/build-nestv3.yml

@@ -0,0 +1,37 @@
+name: PikaOS Package Build Only (amd64-v3)
+
+on:
+  push:
+    branches:
+      - main
+    paths:
+      - '.github/build-nest-v3'
+
+jobs:
+  build:
+    runs-on: ubuntu-latest
+    container:
+      image: ghcr.io/pikaos-linux/pikaos-builder:nestv3
+      volumes:
+        - /proc:/proc
+      options: --privileged -it
+
+    steps:
+    - uses: actions/checkout@v3
+
+    - name: Install SSH key
+      uses: shimataro/ssh-key-action@v2
+      with:
+        key: ${{ secrets.SSH_KEY }}
+        name: id_rsa
+        known_hosts: ${{ secrets.KNOWN_HOSTS }}
+        if_key_exists: replace
+    
+    - name: Update APT Cache
+      run: apt-get update -y
+
+    - name: Set Build Config
+      run: cp -vf ./pika-build-config/amd64-v3.sh ./pika-build-config.sh
+
+    - name: Build Package
+      run: ./main.sh

.github/workflows/release-canaryi386.yml

@@ -0,0 +1,40 @@
+name: PikaOS Package Build & Release (Canary) (i386)
+
+on:
+  push:
+    branches:
+      - main
+    paths:
+      - '.github/release-canary-i386'
+
+jobs:
+  build:
+    runs-on: ubuntu-latest
+    container:
+      image: ghcr.io/pikaos-linux/pikaos-builder:canaryi386
+      volumes:
+        - /proc:/proc
+      options: --privileged -it
+
+    steps:
+    - uses: actions/checkout@v3
+
+    - name: Install SSH key
+      uses: shimataro/ssh-key-action@v2
+      with:
+        key: ${{ secrets.SSH_KEY }}
+        name: id_rsa
+        known_hosts: ${{ secrets.KNOWN_HOSTS }}
+        if_key_exists: replace
+    
+    - name: Update APT Cache
+      run: apt-get update -y
+
+    - name: Set Build Config
+      run: cp -vf ./pika-build-config/i386.sh ./pika-build-config.sh
+
+    - name: Build Package
+      run: ./main.sh
+
+    - name: Release Package
+      run: ./release.sh

.github/workflows/release-canaryv3.yml

@@ -0,0 +1,40 @@
+name: PikaOS Package Build & Release (Canary) (amd64-v3)
+
+on:
+  push:
+    branches:
+      - main
+    paths:
+      - '.github/release-canary-v3'
+
+jobs:
+  build:
+    runs-on: ubuntu-latest
+    container:
+      image: ghcr.io/pikaos-linux/pikaos-builder:canaryv3
+      volumes:
+        - /proc:/proc
+      options: --privileged -it
+
+    steps:
+    - uses: actions/checkout@v3
+
+    - name: Install SSH key
+      uses: shimataro/ssh-key-action@v2
+      with:
+        key: ${{ secrets.SSH_KEY }}
+        name: id_rsa
+        known_hosts: ${{ secrets.KNOWN_HOSTS }}
+        if_key_exists: replace
+    
+    - name: Update APT Cache
+      run: apt-get update -y
+
+    - name: Set Build Config
+      run: cp -vf ./pika-build-config/amd64-v3.sh ./pika-build-config.sh
+
+    - name: Build Package
+      run: ./main.sh
+
+    - name: Release Package
+      run: ./release.sh

.github/workflows/release-debian-v3.yml

@@ -0,0 +1,40 @@
+name: PikaOS Package Build & Release (amd64-v3 on debian)
+
+on:
+  push:
+    branches:
+      - main
+    paths:
+      - '.github/release-debian-v3'
+
+jobs:
+  build:
+    runs-on: ubuntu-latest
+    container:
+      image: debian:sid
+      volumes:
+        - /proc:/proc
+      options: --privileged -it
+
+    steps:
+    - name: Update APT Cache
+      run: apt-get update -y && apt install -y git devscripts dh-make nodejs npm sudo rsync ssh
+    
+    - uses: actions/checkout@v3
+
+    - name: Install SSH key
+      uses: shimataro/ssh-key-action@v2
+      with:
+        key: ${{ secrets.SSH_KEY }}
+        name: id_rsa
+        known_hosts: ${{ secrets.KNOWN_HOSTS }}
+        if_key_exists: replace
+
+    - name: Set Build Config
+      run: cp -vf ./pika-build-config/amd64-v3.sh ./pika-build-config.sh
+
+    - name: Build Package
+      run: ./main.sh
+
+    - name: Release Package
+      run: ./release.sh

.github/workflows/release-nesti386.yml

@@ -0,0 +1,40 @@
+name: PikaOS Package Build & Release (i386)
+
+on:
+  push:
+    branches:
+      - main
+    paths:
+      - '.github/release-nest-i386'
+
+jobs:
+  build:
+    runs-on: ubuntu-latest
+    container:
+      image: ghcr.io/pikaos-linux/pikaos-builder:nesti386
+      volumes:
+        - /proc:/proc
+      options: --privileged -it
+
+    steps:
+    - uses: actions/checkout@v3
+
+    - name: Install SSH key
+      uses: shimataro/ssh-key-action@v2
+      with:
+        key: ${{ secrets.SSH_KEY }}
+        name: id_rsa
+        known_hosts: ${{ secrets.KNOWN_HOSTS }}
+        if_key_exists: replace
+    
+    - name: Update APT Cache
+      run: apt-get update -y
+
+    - name: Set Build Config
+      run: cp -vf ./pika-build-config/i386.sh ./pika-build-config.sh
+
+    - name: Build Package
+      run: ./main.sh
+
+    - name: Release Package
+      run: ./release.sh

.github/workflows/release-nestv3.yml

@@ -0,0 +1,40 @@
+name: PikaOS Package Build & Release (amd64-v3)
+
+on:
+  push:
+    branches:
+      - main
+    paths:
+      - '.github/release-nest-v3'
+
+jobs:
+  build:
+    runs-on: ubuntu-latest
+    container:
+      image: ghcr.io/pikaos-linux/pikaos-builder:nestv3
+      volumes:
+        - /proc:/proc
+      options: --privileged -it
+
+    steps:
+    - uses: actions/checkout@v3
+
+    - name: Install SSH key
+      uses: shimataro/ssh-key-action@v2
+      with:
+        key: ${{ secrets.SSH_KEY }}
+        name: id_rsa
+        known_hosts: ${{ secrets.KNOWN_HOSTS }}
+        if_key_exists: replace
+    
+    - name: Update APT Cache
+      run: apt-get update -y
+
+    - name: Set Build Config
+      run: cp -vf ./pika-build-config/amd64-v3.sh ./pika-build-config.sh
+
+    - name: Build Package
+      run: ./main.sh
+
+    - name: Release Package
+      run: ./release.sh

.github/workflows/release.yml

@@ -1,49 +0,0 @@
-name: PikaOS Package Release
-
-on:
-  workflow_dispatch
-
-jobs:
-  build:
-    runs-on: self-hosted
-    container:
-      image: ghcr.io/pikaos-linux/pika-package-container:latest
-      volumes:
-        - /proc:/proc
-      options: --privileged -it
-
-    steps:
-    - uses: actions/checkout@v3
-
-    - name: Import GPG key
-      id: import_gpg
-      uses: crazy-max/ghaction-import-gpg@v5
-      with:
-        gpg_private_key: ${{ secrets.GPG_PRIVATE_KEY }}
-        passphrase: ${{ secrets.PASSPHRASE }}
-
-    - name: Install SSH key
-      uses: shimataro/ssh-key-action@v2
-      with:
-        key: ${{ secrets.SSH_KEY }}
-        name: id_rsa
-        known_hosts: ${{ secrets.KNOWN_HOSTS }}
-        if_key_exists: replace
-
-    - name: Update apt cache
-      run: apt-get update -y
-    
-    - name: Build Package
-      run: ./main.sh
-
-    - name: Release Package
-      run: ./release.sh
-
-    - name: Purge cache
-      uses: strrife/cloudflare-chunked-purge-action@master
-      env:
-        # Zone is required by both authentication methods
-        CLOUDFLARE_ZONE: ${{ secrets.CLOUDFLARE_ZONE }}
-
-        CLOUDFLARE_TOKEN: ${{ secrets.CLOUDFLARE_TOKEN }}
-        PURGE_URLS: ${{ vars.PURGE_URLS }}

.github/workflows/release_i386.yml

@@ -1,49 +0,0 @@
-name: PikaOS Package Release (i386)
-
-on:
-  workflow_dispatch
-
-jobs:
-  build:
-    runs-on: self-hosted
-    container:
-      image: ghcr.io/pikaos-linux/pika-i386-package-container:latest
-      volumes:
-        - /proc:/proc
-      options: --privileged -it
-
-    steps:
-    - uses: actions/checkout@v3
-
-    - name: Import GPG key
-      id: import_gpg
-      uses: crazy-max/ghaction-import-gpg@v5
-      with:
-        gpg_private_key: ${{ secrets.GPG_PRIVATE_KEY }}
-        passphrase: ${{ secrets.PASSPHRASE }}
-
-    - name: Install SSH key
-      uses: shimataro/ssh-key-action@v2
-      with:
-        key: ${{ secrets.SSH_KEY }}
-        name: id_rsa
-        known_hosts: ${{ secrets.KNOWN_HOSTS }}
-        if_key_exists: replace
-
-    - name: Update apt cache
-      run: apt-get update -y
-    
-    - name: Build Package
-      run: ./main32.sh
-
-    - name: Release Package
-      run: ./release.sh
-
-    - name: Purge cache
-      uses: strrife/cloudflare-chunked-purge-action@master
-      env:
-        # Zone is required by both authentication methods
-        CLOUDFLARE_ZONE: ${{ secrets.CLOUDFLARE_ZONE }}
-
-        CLOUDFLARE_TOKEN: ${{ secrets.CLOUDFLARE_TOKEN }}
-        PURGE_URLS: ${{ vars.PURGE_URLS }}

LICENSE.md

@@ -0,0 +1,24 @@
+MIT License (With DPKG packaging compatibility)
+
+Copyright (c) 2024 PikaOS
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in all
+copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+SOFTWARE.
+
+Notes:
+The files covered by this license are any files and directories in the root of this repository (including but not limited to: `main.sh`, `release.sh`, and `.github`), with the exception of the `debian` directory and its contents if `debian/copyright` exists, and declares any files or directories as a different LICENSE/COPYRIGHT.

debian/control

@@ -2,26 +2,28 @@ Source: mesa-git
 Section: graphics
 Priority: optional
 Maintainer: First Mate Rummey <fmrummey@gmail.com>
-XSBC-Original-Maintainer: Ubuntu X-SWAT <ubuntu-x@lists.ubuntu.com>
+XSBC-Original-Maintainer: Ubuntu Developers <ubuntu-devel-discuss@lists.ubuntu.com>
 Uploaders: Andreas Boll <aboll@debian.org>
 Standards-Version: 4.1.4
 Build-Depends:
- debhelper-compat (= 12),
- directx-headers-dev (>= 1.602.0) [linux-amd64 linux-arm64],
- glslang-tools [amd64 arm64 armel armhf i386 mips64el mipsel powerpc ppc64 ppc64el riscv64 s390x sparc64 x32],
- meson (>= 0.45),
+ debhelper-compat (= 13),
+ directx-headers-dev (>= 1.613.0) [linux-amd64 linux-arm64],
+ flatbuffers-compiler [linux-arm64],
+ glslang-tools [amd64 arm64 armel armhf i386 loong64 mips64el powerpc ppc64 ppc64el riscv64 s390x sparc64 x32],
+ meson (>= 1.4.0),
  quilt (>= 0.63-8.2~),
- pkg-config,
- libdrm-dev (>= 2.4.107-4),
+ pkgconf,
+ libdrm-dev (>= 2.4.121),
  libx11-dev,
  libxxf86vm-dev,
  libexpat1-dev,
+ libflatbuffers-dev [linux-arm64],
  libsensors-dev [!hurd-any],
  libxfixes-dev,
  libxext-dev,
- libva-dev (>= 1.6.0) [linux-any kfreebsd-any] <!pkg.mesa.nolibva>,
- libvdpau-dev (>= 1.1.1) [linux-any kfreebsd-any],
- libvulkan-dev [amd64 arm64 armel armhf i386 mips64el mipsel powerpc ppc64 ppc64el riscv64 s390x sparc64 x32],
+ libva-dev (>= 1.6.0) [linux-any] <!pkg.mesa.nolibva>,
+ libvdpau-dev (>= 1.5) [linux-any],
+ libvulkan-dev [amd64 arm64 armel armhf i386 loong64 mips64el powerpc ppc64 ppc64el riscv64 s390x sparc64 x32],
  x11proto-dev,
  linux-libc-dev (>= 2.6.31) [linux-any],
  libx11-xcb-dev,
@@ -35,31 +37,39 @@ Build-Depends:
  libxcb-sync-dev,
  libxrandr-dev,
  libxshmfence-dev (>= 1.1),
+ libxtensor-dev [linux-arm64],
  libzstd-dev,
+ lua5.4 [arm64 armel armhf],
  python3,
  python3-mako,
  python3-ply,
+ python3-yaml,
+ python3-pycparser [arm64 armhf],
  python3-setuptools,
  flex,
  bison,
- libelf-dev [amd64 arm64 armel armhf i386 mips64el mipsel powerpc ppc64 ppc64el riscv64 s390x sparc64 x32],
+ libelf-dev [amd64 arm64 armel armhf i386 loong64 mips64el powerpc ppc64 ppc64el riscv64 s390x sparc64 x32],
  libwayland-dev (>= 1.15.0) [linux-any],
  libwayland-egl-backend-dev (>= 1.15.0) [linux-any],
- llvm-15-dev [amd64 arm64 armel armhf i386 mips64el mipsel powerpc ppc64 ppc64el riscv64 s390x sparc64 x32],
- libclang-15-dev [amd64 arm64 armel armhf i386 mips64el mipsel powerpc ppc64 ppc64el riscv64 s390x sparc64 x32],
- libclang-cpp15-dev [amd64 arm64 armel armhf i386 mips64el mipsel powerpc ppc64 ppc64el riscv64 s390x sparc64 x32],
- libclc-15-dev [amd64 arm64 armel armhf i386 mips64el mipsel powerpc ppc64 ppc64el riscv64 s390x sparc64 x32],
- wayland-protocols (>= 1.24),
+ llvm-18-dev [amd64 arm64 armel armhf i386 loong64 mips64el powerpc ppc64 ppc64el riscv64 s390x sparc64 x32],
+ libclang-18-dev [amd64 arm64 armel armhf i386 loong64 mips64el powerpc ppc64 ppc64el riscv64 s390x sparc64 x32],
+ libclang-cpp18-dev [amd64 arm64 armel armhf i386 loong64 mips64el powerpc ppc64 ppc64el riscv64 s390x sparc64 x32],
+ libclc-18-dev [amd64 arm64 armel armhf i386 loong64 mips64el powerpc ppc64 ppc64el riscv64 s390x sparc64 x32],
+ libclc-18 [amd64 arm64 armel armhf i386 loong64 mips64el powerpc ppc64 ppc64el riscv64 s390x sparc64 x32],
+ wayland-protocols (>= 1.34),
  zlib1g-dev,
  libglvnd-core-dev (>= 1.3.2),
- valgrind [amd64 arm64 armhf i386 mips64el mipsel powerpc ppc64 ppc64el s390x],
- rustc [amd64 arm64 armel armhf mips64el mipsel ppc64el s390x],
- bindgen [amd64 arm64 armel armhf mips64el mipsel ppc64el s390x],
- llvm-spirv-15 [amd64 arm64 armel armhf mips64el mipsel ppc64el s390x],
- libclc-15 [amd64 arm64 armel armhf mips64el mipsel ppc64el s390x],
- libllvmspirvlib-15-dev,
+ valgrind [amd64 arm64 armhf i386 mips64el powerpc ppc64 ppc64el s390x],
+ rustc (>= 1.73) [amd64 arm64 armel armhf loong64 mips64el powerpc ppc64 ppc64el riscv64 s390x x32],
+ rustfmt [amd64 arm64 armel armhf loong64 mips64el powerpc ppc64 ppc64el riscv64 s390x x32],
+ bindgen (>= 0.66.1~) [amd64 arm64 armel armhf loong64 mips64el powerpc ppc64 ppc64el riscv64 s390x x32],
+ cbindgen [amd64 arm64 armel armhf loong64 mips64el powerpc ppc64 ppc64el riscv64 s390x x32],
+ llvm-spirv-18 [amd64 arm64 armel armhf loong64 mips64el powerpc ppc64 ppc64el riscv64 s390x x32],
+ libllvmspirvlib-18-dev [amd64 arm64 armel armhf i386 loong64 mips64el powerpc ppc64 ppc64el riscv64 s390x sparc64 x32],
+ librust-paste-dev [amd64 arm64 armhf ppc64 riscv64 x32],
+ librust-syn-dev [amd64 arm64 armhf ppc64 riscv64 x32],
 Rules-Requires-Root: no
-Vcs: https://salsa.debian.org/xorg-team/lib/mesa.git
+Vcs-Git: https://salsa.debian.org/xorg-team/lib/mesa.git
 Vcs-Browser: https://salsa.debian.org/xorg-team/lib/mesa
 Homepage: https://mesa3d.org/
 
@@ -276,6 +286,19 @@ Multi-Arch: same
 Description: transitional dummy package
  This is a transitional dummy package, it can be safely removed.
 
+Package: mesa-teflon-delegate-git
+Section: libs
+Architecture: arm64
+Depends:
+ ${shlibs:Depends},
+ ${misc:Depends},
+Pre-Depends: ${misc:Pre-Depends}
+Multi-Arch: same
+Description: Mesa TensorFlow Lite external delegate
+ TensorFlow Lite delegate which can make use of NPUs to accelerate ML
+ inference. It is implemented in the form of a external delegate, a shared
+ library which the TensorFlow Lite runtime can load at startup.
+
 Package: libegl1-mesa-git-dev
 Section: libdevel
 Architecture: any
@@ -563,7 +586,7 @@ Section: libs
 Architecture: amd64 arm64 armel armhf i386 mips64el mipsel powerpc ppc64 ppc64el riscv64 s390x sparc64 x32
 Pre-Depends: ${misc:Pre-Depends}
 Depends: libmesa-git,
- libclc-15,
+ libclc-17,
  ocl-icd-libopencl1 | libopencl1,
  ${shlibs:Depends},
  ${misc:Depends},

debian/control.in

@@ -2,25 +2,27 @@ Source: mesa-git
 Section: graphics
 Priority: optional
 Maintainer: First Mate Rummey <fmrummey@gmail.com>
-XSBC-Original-Maintainer: Ubuntu X-SWAT <ubuntu-x@lists.ubuntu.com>
+XSBC-Original-Maintainer: Ubuntu Developers <ubuntu-devel-discuss@lists.ubuntu.com>
 Uploaders: Andreas Boll <aboll@debian.org>
 Standards-Version: 4.1.4
 Build-Depends:
- debhelper-compat (= 12),
- directx-headers-dev (>= 1.602.0) [linux-amd64 linux-arm64],
+ debhelper-compat (= 13),
+ directx-headers-dev (>= 1.613.0) [linux-amd64 linux-arm64],
+ flatbuffers-compiler [linux-arm64],
  glslang-tools [@LLVM_ARCHS@],
- meson (>= 0.45),
+ meson (>= 1.4.0),
  quilt (>= 0.63-8.2~),
- pkg-config,
- libdrm-dev (>= 2.4.107-4),
+ pkgconf,
+ libdrm-dev (>= 2.4.121),
  libx11-dev,
  libxxf86vm-dev,
  libexpat1-dev,
+ libflatbuffers-dev [linux-arm64],
  libsensors-dev [!hurd-any],
  libxfixes-dev,
  libxext-dev,
- libva-dev (>= 1.6.0) [linux-any kfreebsd-any] <!pkg.mesa.nolibva>,
- libvdpau-dev (>= 1.1.1) [linux-any kfreebsd-any],
+ libva-dev (>= 1.6.0) [linux-any] <!pkg.mesa.nolibva>,
+ libvdpau-dev (>= 1.5) [linux-any],
  libvulkan-dev [@LLVM_ARCHS@],
  x11proto-dev,
  linux-libc-dev (>= 2.6.31) [linux-any],
@@ -35,10 +37,14 @@ Build-Depends:
  libxcb-sync-dev,
  libxrandr-dev,
  libxshmfence-dev (>= 1.1),
+ libxtensor-dev [linux-arm64],
  libzstd-dev,
+ lua5.4 [arm64 armel armhf],
  python3,
  python3-mako,
  python3-ply,
+ python3-yaml,
+ python3-pycparser [arm64 armhf],
  python3-setuptools,
  flex,
  bison,
@@ -49,17 +55,21 @@ Build-Depends:
  libclang-@LLVM_VERSION@-dev [@LLVM_ARCHS@],
  libclang-cpp@LLVM_VERSION@-dev [@LLVM_ARCHS@],
  libclc-@LLVM_VERSION@-dev [@LLVM_ARCHS@],
- wayland-protocols (>= 1.24),
+ libclc-@LLVM_VERSION@ [@LLVM_ARCHS@],
+ wayland-protocols (>= 1.34),
  zlib1g-dev,
  libglvnd-core-dev (>= 1.3.2),
  valgrind [@VALGRIND_ARCHS@],
- rustc [@RUSTICL_ARCHS@],
- bindgen [@RUSTICL_ARCHS@],
+ rustc (>= 1.73) [@RUSTICL_ARCHS@],
+ rustfmt [@RUSTICL_ARCHS@],
+ bindgen (>= 0.66.1~) [@RUSTICL_ARCHS@],
+ cbindgen [@RUSTICL_ARCHS@],
  llvm-spirv-@LLVM_VERSION@ [@RUSTICL_ARCHS@],
- libclc-@LLVM_VERSION@ [@RUSTICL_ARCHS@],
- libllvmspirvlib-@LLVM_VERSION@-dev,
+ libllvmspirvlib-@LLVM_VERSION@-dev [@LLVM_ARCHS@],
+ librust-paste-dev [@NVK_ARCHS@],
+ librust-syn-dev [@NVK_ARCHS@],
 Rules-Requires-Root: no
-Vcs: https://salsa.debian.org/xorg-team/lib/mesa.git
+Vcs-Git: https://salsa.debian.org/xorg-team/lib/mesa.git
 Vcs-Browser: https://salsa.debian.org/xorg-team/lib/mesa
 Homepage: https://mesa3d.org/
 

debian/libegl-mesa0-git.symbols

@@ -1,2 +0,0 @@
-libEGL_mesa.so.0 libegl-mesa0-git #MINVER#
- __egl_Main@Base 17.0.0~

debian/libgbm1-git.symbols

@@ -1,40 +0,0 @@
-libgbm.so.1 libgbm1-git #MINVER#
-| libgbm1-private
- gbm_bo_create@Base 7.11~1
- gbm_bo_create_with_modifiers2@Base 21.3.0~rc1
- gbm_bo_create_with_modifiers@Base 17.1.0~rc2
- gbm_bo_destroy@Base 7.11~1
- gbm_bo_get_bpp@Base 17.3.0~rc1
- gbm_bo_get_device@Base 8.1~0
- gbm_bo_get_fd@Base 10.2~0
- gbm_bo_get_fd_for_plane@Base 21.1.0
- gbm_bo_get_format@Base 8.1~0
- gbm_bo_get_handle@Base 7.11~1
- gbm_bo_get_handle_for_plane@Base 17.1.0~rc2
- gbm_bo_get_height@Base 7.11~1
- gbm_bo_get_modifier@Base 17.1.0~rc2
- gbm_bo_get_offset@Base 17.1.0~rc2
- gbm_bo_get_plane_count@Base 17.1.0~rc2
- gbm_bo_get_stride@Base 8.1~0
- gbm_bo_get_stride_for_plane@Base 17.1.0~rc2
- gbm_bo_get_user_data@Base 8.1~0
- gbm_bo_get_width@Base 7.11~1
- gbm_bo_import@Base 8.1~0
- gbm_bo_map@Base 12.0.0~0
- gbm_bo_set_user_data@Base 8.1~0
- gbm_bo_unmap@Base 12.0.0~0
- gbm_bo_write@Base 8.1~0
- gbm_create_device@Base 7.11~1
- gbm_device_destroy@Base 7.11~1
- gbm_device_get_backend_name@Base 7.11~1
- gbm_device_get_fd@Base 7.11~1
- gbm_device_get_format_modifier_plane_count@Base 17.3.0~rc1
- gbm_device_is_format_supported@Base 8.1~0
- gbm_format_get_name@Base 19.0.0~rc5
- gbm_surface_create@Base 8.1~0
- gbm_surface_create_with_modifiers2@Base 21.3.0~rc1
- gbm_surface_create_with_modifiers@Base 17.1.0~rc2
- gbm_surface_destroy@Base 8.1~0
- gbm_surface_has_free_buffers@Base 8.1~0
- gbm_surface_lock_front_buffer@Base 8.1~0
- gbm_surface_release_buffer@Base 8.1~0

debian/libgl1-mesa-dri-git.install

@@ -1 +1,2 @@
 usr/share/drirc.d/00-mesa-defaults.conf
+usr/lib/*/libgallium-*.so

debian/libglx-mesa0.symbols-git.hurd

@@ -1,6 +0,0 @@
-libGLX_mesa.so.0 libglx-mesa0-git
- __glx_Main@Base 17.0.0~
- glAreTexturesResidentEXT@Base 0
- glDeleteTexturesEXT@Base 0
- glGenTexturesEXT@Base 0
- glIsTextureEXT@Base 0

debian/libxatracker2-git.symbols

@@ -1,35 +0,0 @@
-libxatracker.so.2 libxatracker2-git #MINVER#
- xa_composite_allocation@Base 0
- xa_composite_check_accelerated@Base 0
- xa_composite_done@Base 0
- xa_composite_prepare@Base 0
- xa_composite_rect@Base 0
- xa_context_create@Base 0
- xa_context_default@Base 0
- xa_context_destroy@Base 0
- xa_context_flush@Base 0
- xa_copy@Base 0
- xa_copy_done@Base 0
- xa_copy_prepare@Base 0
- xa_fence_destroy@Base 0
- xa_fence_get@Base 0
- xa_fence_wait@Base 0
- xa_format_check_supported@Base 0
- xa_solid@Base 0
- xa_solid_done@Base 0
- xa_solid_prepare@Base 0
- xa_surface_create@Base 0
- xa_surface_dma@Base 0
- xa_surface_format@Base 0
- xa_surface_from_handle2@Base 11.1.0~
- xa_surface_from_handle@Base 0
- xa_surface_handle@Base 0
- xa_surface_map@Base 0
- xa_surface_redefine@Base 0
- xa_surface_ref@Base 0
- xa_surface_unmap@Base 0
- xa_surface_unref@Base 0
- xa_tracker_create@Base 0
- xa_tracker_destroy@Base 0
- xa_tracker_version@Base 0
- xa_yuv_planar_blit@Base 0

debian/mesa-vulkan-drivers-git.install

@@ -5,6 +5,3 @@ usr/share/vulkan/icd.d/*.json
 usr/share/vulkan/implicit_layer.d/*.json
 usr/lib/*/libvulkan_*.so
 usr/lib/*/libVkLayer_*.so
-# microsoft
-#usr/bin/spirv2dxil
-#usr/lib/x86_64-linux-gnu/libspirv_to_dxil.*

debian/patches/07_gallium-fix-build-failure-on-powerpcspe.diff

@@ -1,35 +0,0 @@
-From a4f14e7239780b02af8d74669c5458d4b0957d4d Mon Sep 17 00:00:00 2001
-From: Roland Stigge <stigge@antcom.de>
-Date: Sun, 2 Mar 2014 19:52:56 +0100
-Subject: [PATCH] gallium: fix build failure on powerpcspe
-
-In the case of powerpc, mesa activates some altivec instructions
-that are unknown on the powerpcspe architecture (see
-https://wiki.debian.org/PowerPCSPEPort), causing a build failure as the
-'vand' opcode is not recognized by the assembler.
-
-This patch fixes this by preventing the PPC-specialcasing in case of
-powerpcspe (__NO_FPRS__ is only defined there).
-
-https://bugs.debian.org/695746
----
- src/gallium/include/pipe/p_config.h | 2 ++
- 1 file changed, 2 insertions(+)
-
---- a/src/util/detect_arch.h
-+++ b/src/util/detect_arch.h
-@@ -70,12 +70,14 @@
- #endif
- #endif
- 
-+#ifndef __NO_FPRS__
- #if defined(__ppc__) || defined(__ppc64__) || defined(__PPC__) || defined(__PPC64__)
- #define DETECT_ARCH_PPC 1
- #if defined(__ppc64__) || defined(__PPC64__)
- #define DETECT_ARCH_PPC_64 1
- #endif
- #endif
-+#endif
- 
- #if defined(__s390x__)
- #define DETECT_ARCH_S390 1

debian/patches/21929.patch

@@ -1,679 +0,0 @@
-From ed9fb6be100cff6c2066beb0cdf8b3a17cab292c Mon Sep 17 00:00:00 2001
-From: Konstantin Seurer <konstantin.seurer@gmail.com>
-Date: Sat, 24 Jun 2023 15:49:13 +0200
-Subject: [PATCH 1/4] radv: Add rt.monolithic to radv_pipeline_key
-
----
- src/amd/vulkan/radv_shader.h | 4 ++++
- 1 file changed, 4 insertions(+)
-
-diff --git a/src/amd/vulkan/radv_shader.h b/src/amd/vulkan/radv_shader.h
-index 0c53695edae7c..6eb95fdd0a097 100644
---- a/src/amd/vulkan/radv_shader.h
-+++ b/src/amd/vulkan/radv_shader.h
-@@ -135,6 +135,10 @@ struct radv_pipeline_key {
- 
-       bool line_smooth_enabled;
-    } ps;
-+
-+   struct {
-+      bool monolithic;
-+   } rt;
- };
- 
- struct radv_nir_compiler_options {
--- 
-
-
-From 8f45cc08361f55c1e613a11198b1ae97c519406e Mon Sep 17 00:00:00 2001
-From: Konstantin Seurer <konstantin.seurer@gmail.com>
-Date: Sat, 24 Jun 2023 15:46:51 +0200
-Subject: [PATCH 2/4] radv/rt: Store NIR shaders separately
-
-In order to compile monolithic shaders with pipeline libraries, we need
-to keep the NIR around for inlining recursive stages.
----
- src/amd/vulkan/radv_pipeline_cache.c |  9 +--
- src/amd/vulkan/radv_pipeline_rt.c    | 93 +++++++++++++++++++++-------
- src/amd/vulkan/radv_private.h        |  1 +
- src/amd/vulkan/radv_rt_shader.c      |  7 +--
- 4 files changed, 79 insertions(+), 31 deletions(-)
-
-diff --git a/src/amd/vulkan/radv_pipeline_cache.c b/src/amd/vulkan/radv_pipeline_cache.c
-index 5bbbc755ae11f..7e4c6f8898130 100644
---- a/src/amd/vulkan/radv_pipeline_cache.c
-+++ b/src/amd/vulkan/radv_pipeline_cache.c
-@@ -481,11 +481,12 @@ radv_ray_tracing_pipeline_cache_search(struct radv_device *device, struct vk_pip
-       pipeline->base.base.shaders[MESA_SHADER_INTERSECTION] = radv_shader_ref(pipeline_obj->shaders[idx++]);
- 
-    for (unsigned i = 0; i < pCreateInfo->stageCount; i++) {
--      if (radv_ray_tracing_stage_is_compiled(&pipeline->stages[i])) {
-+      if (radv_ray_tracing_stage_is_compiled(&pipeline->stages[i]))
-          pipeline->stages[i].shader = &radv_shader_ref(pipeline_obj->shaders[idx++])->base;
--      } else if (is_library) {
--         pipeline->stages[i].shader = radv_pipeline_cache_search_nir(device, cache, pipeline->stages[i].sha1);
--         complete &= pipeline->stages[i].shader != NULL;
-+
-+      if (is_library) {
-+         pipeline->stages[i].nir = radv_pipeline_cache_search_nir(device, cache, pipeline->stages[i].sha1);
-+         complete &= pipeline->stages[i].nir != NULL;
-       }
-    }
- 
-diff --git a/src/amd/vulkan/radv_pipeline_rt.c b/src/amd/vulkan/radv_pipeline_rt.c
-index c86ea3a508468..85afc8cb28e1b 100644
---- a/src/amd/vulkan/radv_pipeline_rt.c
-+++ b/src/amd/vulkan/radv_pipeline_rt.c
-@@ -263,7 +263,10 @@ radv_rt_fill_stage_info(const VkRayTracingPipelineCreateInfoKHR *pCreateInfo, st
-          RADV_FROM_HANDLE(radv_pipeline, pipeline, pCreateInfo->pLibraryInfo->pLibraries[i]);
-          struct radv_ray_tracing_pipeline *library_pipeline = radv_pipeline_to_ray_tracing(pipeline);
-          for (unsigned j = 0; j < library_pipeline->stage_count; ++j) {
--            stages[idx].shader = vk_pipeline_cache_object_ref(library_pipeline->stages[j].shader);
-+            stages[idx].nir = vk_pipeline_cache_object_ref(library_pipeline->stages[j].nir);
-+            if (library_pipeline->stages[j].shader)
-+               stages[idx].shader = vk_pipeline_cache_object_ref(library_pipeline->stages[j].shader);
-+
-             stages[idx].stage = library_pipeline->stages[j].stage;
-             stages[idx].stack_size = library_pipeline->stages[j].stack_size;
-             memcpy(stages[idx].sha1, library_pipeline->stages[j].sha1, SHA1_DIGEST_LENGTH);
-@@ -462,45 +465,83 @@ radv_rt_compile_shaders(struct radv_device *device, struct vk_pipeline_cache *ca
-       return VK_PIPELINE_COMPILE_REQUIRED;
-    VkResult result = VK_SUCCESS;
- 
--   struct radv_ray_tracing_stage *stages = pipeline->stages;
-+   struct radv_ray_tracing_stage *rt_stages = pipeline->stages;
-+
-+   struct radv_shader_stage *stages = calloc(pCreateInfo->stageCount, sizeof(struct radv_shader_stage));
-+   if (!stages)
-+      return VK_ERROR_OUT_OF_HOST_MEMORY;
-+
-+   bool has_callable = false;
-+   for (uint32_t i = 0; i < pipeline->stage_count; i++) {
-+      if (pipeline->stages[i].stage == MESA_SHADER_CALLABLE) {
-+         has_callable = true;
-+         break;
-+      }
-+   }
- 
-    for (uint32_t idx = 0; idx < pCreateInfo->stageCount; idx++) {
-+      if (rt_stages[idx].shader || rt_stages[idx].nir)
-+         continue;
-+
-       int64_t stage_start = os_time_get_nano();
--      struct radv_shader_stage stage;
--      radv_pipeline_stage_init(&pCreateInfo->pStages[idx], pipeline_layout, &stage);
- 
--      if (stages[idx].shader)
--         goto feedback;
-+      struct radv_shader_stage *stage = &stages[idx];
-+      radv_pipeline_stage_init(&pCreateInfo->pStages[idx], pipeline_layout, stage);
- 
-       /* precompile the shader */
--      stage.nir = radv_parse_rt_stage(device, &pCreateInfo->pStages[idx], key, pipeline_layout);
-+      stage->nir = radv_parse_rt_stage(device, &pCreateInfo->pStages[idx], key, pipeline_layout);
-+
-+      /* Cases in which we need to keep around the NIR:
-+       *    - pipeline library: The final pipeline might be monolithic in which case it will need every NIR shader.
-+       *                        If there is a callable shader, we can be sure that the final pipeline won't be
-+       *                        monolithic.
-+       *    - non-recursive:    Non-recursive shaders are inlined into the traversal shader.
-+       *    - monolithic:       Callable shaders (chit/miss) are inlined into the raygen shader.
-+       */
-+      bool compiled = radv_ray_tracing_stage_is_compiled(&rt_stages[idx]);
-+      bool library = pCreateInfo->flags & VK_PIPELINE_CREATE_LIBRARY_BIT_KHR;
-+      bool nir_needed =
-+         (library && !has_callable) || !compiled || (key->rt.monolithic && rt_stages[idx].stage != MESA_SHADER_RAYGEN);
-+      nir_needed &= !rt_stages[idx].nir;
-+      if (nir_needed) {
-+         rt_stages[idx].stack_size = stage->nir->scratch_size;
-+         rt_stages[idx].nir = radv_pipeline_cache_nir_to_handle(device, cache, stage->nir, rt_stages[idx].sha1,
-+                                                                !key->optimisations_disabled);
-+      }
- 
--      if (radv_ray_tracing_stage_is_compiled(&stages[idx])) {
--         uint32_t stack_size = 0;
-+      stage->feedback.duration = os_time_get_nano() - stage_start;
-+   }
- 
-+   for (uint32_t idx = 0; idx < pCreateInfo->stageCount; idx++) {
-+      int64_t stage_start = os_time_get_nano();
-+      struct radv_shader_stage *stage = &stages[idx];
-+
-+      /* Cases in which we need to compile the shader (raygen/callable/chit/miss):
-+       *    TODO: - monolithic: Extend the loop to cover imported stages and force compilation of imported raygen
-+       *                        shaders since pipeline library shaders use separate compilation.
-+       *    - separate:   Compile any recursive stage if wasn't compiled yet.
-+       * TODO: Skip chit and miss shaders in the monolithic case.
-+       */
-+      bool shader_needed = radv_ray_tracing_stage_is_compiled(&rt_stages[idx]) && !rt_stages[idx].shader;
-+      if (shader_needed) {
-+         uint32_t stack_size = 0;
-          struct radv_serialized_shader_arena_block *replay_block =
-             capture_replay_handles[idx].arena_va ? &capture_replay_handles[idx] : NULL;
- 
-          struct radv_shader *shader;
-          result =
--            radv_rt_nir_to_asm(device, cache, pCreateInfo, key, pipeline, &stage, &stack_size, replay_block, &shader);
--         stages[idx].stack_size = stack_size;
--         stages[idx].shader = shader ? &shader->base : NULL;
--      } else {
--         stages[idx].stack_size = stage.nir->scratch_size;
--         stages[idx].shader =
--            radv_pipeline_cache_nir_to_handle(device, cache, stage.nir, stages[idx].sha1, !key->optimisations_disabled);
--      }
--      ralloc_free(stage.nir);
-+            radv_rt_nir_to_asm(device, cache, pCreateInfo, key, pipeline, stage, &stack_size, replay_block, &shader);
-+         if (result != VK_SUCCESS)
-+            goto cleanup;
- 
--      if (result != VK_SUCCESS)
--         return result;
-+         rt_stages[idx].stack_size = stack_size;
-+         rt_stages[idx].shader = shader ? &shader->base : NULL;
-+      }
- 
--   feedback:
-       if (creation_feedback && creation_feedback->pipelineStageCreationFeedbackCount) {
-          assert(idx < creation_feedback->pipelineStageCreationFeedbackCount);
--         stage.feedback.duration = os_time_get_nano() - stage_start;
--         creation_feedback->pPipelineStageCreationFeedbacks[idx] = stage.feedback;
-+         stage->feedback.duration += os_time_get_nano() - stage_start;
-+         creation_feedback->pPipelineStageCreationFeedbacks[idx] = stage->feedback;
-       }
-    }
- 
-@@ -527,6 +568,10 @@ radv_rt_compile_shaders(struct radv_device *device, struct vk_pipeline_cache *ca
-    result = radv_rt_nir_to_asm(device, cache, pCreateInfo, key, pipeline, &traversal_stage, NULL, NULL,
-                                &pipeline->base.base.shaders[MESA_SHADER_INTERSECTION]);
- 
-+cleanup:
-+   for (uint32_t i = 0; i < pCreateInfo->stageCount; i++)
-+      ralloc_free(stages[i].nir);
-+   free(stages);
-    return result;
- }
- 
-@@ -732,6 +777,8 @@ void
- radv_destroy_ray_tracing_pipeline(struct radv_device *device, struct radv_ray_tracing_pipeline *pipeline)
- {
-    for (unsigned i = 0; i < pipeline->stage_count; i++) {
-+      if (pipeline->stages[i].nir)
-+         vk_pipeline_cache_object_unref(&device->vk, pipeline->stages[i].nir);
-       if (pipeline->stages[i].shader)
-          vk_pipeline_cache_object_unref(&device->vk, pipeline->stages[i].shader);
-    }
-diff --git a/src/amd/vulkan/radv_private.h b/src/amd/vulkan/radv_private.h
-index 47e315488e9f9..7ab46738b46f0 100644
---- a/src/amd/vulkan/radv_private.h
-+++ b/src/amd/vulkan/radv_private.h
-@@ -2364,6 +2364,7 @@ struct radv_ray_tracing_group {
- };
- 
- struct radv_ray_tracing_stage {
-+   struct vk_pipeline_cache_object *nir;
-    struct vk_pipeline_cache_object *shader;
-    gl_shader_stage stage;
-    uint32_t stack_size;
-diff --git a/src/amd/vulkan/radv_rt_shader.c b/src/amd/vulkan/radv_rt_shader.c
-index cc92beebc3503..3def324bcccf3 100644
---- a/src/amd/vulkan/radv_rt_shader.c
-+++ b/src/amd/vulkan/radv_rt_shader.c
-@@ -1132,7 +1132,7 @@ visit_any_hit_shaders(struct radv_device *device, nir_builder *b, struct travers
-       if (is_dup)
-          continue;
- 
--      nir_shader *nir_stage = radv_pipeline_cache_handle_to_nir(device, data->pipeline->stages[shader_id].shader);
-+      nir_shader *nir_stage = radv_pipeline_cache_handle_to_nir(device, data->pipeline->stages[shader_id].nir);
-       assert(nir_stage);
- 
-       insert_rt_case(b, nir_stage, vars, sbt_idx, data->pipeline->groups[i].handle.any_hit_index);
-@@ -1262,13 +1262,12 @@ handle_candidate_aabb(nir_builder *b, struct radv_leaf_intersection *intersectio
-       if (is_dup)
-          continue;
- 
--      nir_shader *nir_stage = radv_pipeline_cache_handle_to_nir(data->device, data->pipeline->stages[shader_id].shader);
-+      nir_shader *nir_stage = radv_pipeline_cache_handle_to_nir(data->device, data->pipeline->stages[shader_id].nir);
-       assert(nir_stage);
- 
-       nir_shader *any_hit_stage = NULL;
-       if (any_hit_shader_id != VK_SHADER_UNUSED_KHR) {
--         any_hit_stage =
--            radv_pipeline_cache_handle_to_nir(data->device, data->pipeline->stages[any_hit_shader_id].shader);
-+         any_hit_stage = radv_pipeline_cache_handle_to_nir(data->device, data->pipeline->stages[any_hit_shader_id].nir);
-          assert(any_hit_stage);
- 
-          /* reserve stack size for any_hit before it is inlined */
--- 
-GitLab
-
-
-From bba42cbc235e75a5c7ed05e55e48f71640c68ad4 Mon Sep 17 00:00:00 2001
-From: Konstantin Seurer <konstantin.seurer@gmail.com>
-Date: Mon, 21 Aug 2023 13:32:53 +0200
-Subject: [PATCH 3/4] radv/rt: Add monolithic raygen lowering
-
-Ray traversal is inlined to allow for constant folding and avoid
-spilling.
----
- src/amd/vulkan/radv_pipeline_rt.c |  11 +-
- src/amd/vulkan/radv_rt_shader.c   | 276 ++++++++++++++++++++++++++----
- src/amd/vulkan/radv_shader.h      |   3 +-
- 3 files changed, 248 insertions(+), 42 deletions(-)
-
-diff --git a/src/amd/vulkan/radv_pipeline_rt.c b/src/amd/vulkan/radv_pipeline_rt.c
-index 85afc8cb28e1b..12562c6cf89ba 100644
---- a/src/amd/vulkan/radv_pipeline_rt.c
-+++ b/src/amd/vulkan/radv_pipeline_rt.c
-@@ -356,9 +356,8 @@ move_rt_instructions(nir_shader *shader)
- static VkResult
- radv_rt_nir_to_asm(struct radv_device *device, struct vk_pipeline_cache *cache,
-                    const VkRayTracingPipelineCreateInfoKHR *pCreateInfo, const struct radv_pipeline_key *pipeline_key,
--                   const struct radv_ray_tracing_pipeline *pipeline, struct radv_shader_stage *stage,
--                   uint32_t *stack_size, struct radv_serialized_shader_arena_block *replay_block,
--                   struct radv_shader **out_shader)
-+                   struct radv_ray_tracing_pipeline *pipeline, struct radv_shader_stage *stage, uint32_t *stack_size,
-+                   struct radv_serialized_shader_arena_block *replay_block, struct radv_shader **out_shader)
- {
-    struct radv_shader_binary *binary;
-    bool keep_executable_info = radv_pipeline_capture_shaders(device, pipeline->base.base.create_flags);
-@@ -384,7 +383,8 @@ radv_rt_nir_to_asm(struct radv_device *device, struct vk_pipeline_cache *cache,
-    uint32_t num_resume_shaders = 0;
-    nir_shader **resume_shaders = NULL;
- 
--   if (stage->stage != MESA_SHADER_INTERSECTION) {
-+   bool monolithic_raygen = pipeline_key->rt.monolithic && stage->stage == MESA_SHADER_RAYGEN;
-+   if (stage->stage != MESA_SHADER_INTERSECTION && !monolithic_raygen) {
-       nir_builder b = nir_builder_at(nir_after_cf_list(&nir_shader_get_entrypoint(stage->nir)->body));
-       nir_rt_return_amd(&b);
- 
-@@ -411,7 +411,8 @@ radv_rt_nir_to_asm(struct radv_device *device, struct vk_pipeline_cache *cache,
-    for (uint32_t i = 0; i < num_shaders; i++) {
-       struct radv_shader_stage temp_stage = *stage;
-       temp_stage.nir = shaders[i];
--      radv_nir_lower_rt_abi(temp_stage.nir, pCreateInfo, &temp_stage.args, &stage->info, stack_size, i > 0);
-+      radv_nir_lower_rt_abi(temp_stage.nir, pCreateInfo, &temp_stage.args, &stage->info, stack_size, i > 0, device,
-+                            pipeline, pipeline_key);
-       radv_optimize_nir(temp_stage.nir, pipeline_key->optimisations_disabled);
-       radv_postprocess_nir(device, pipeline_key, &temp_stage);
- 
-diff --git a/src/amd/vulkan/radv_rt_shader.c b/src/amd/vulkan/radv_rt_shader.c
-index 3def324bcccf3..362d918597008 100644
---- a/src/amd/vulkan/radv_rt_shader.c
-+++ b/src/amd/vulkan/radv_rt_shader.c
-@@ -1306,6 +1306,87 @@ handle_candidate_aabb(nir_builder *b, struct radv_leaf_intersection *intersectio
-    nir_pop_if(b, NULL);
- }
- 
-+static void
-+visit_closest_hit_shaders(struct radv_device *device, nir_builder *b, struct radv_ray_tracing_pipeline *pipeline,
-+                          struct rt_variables *vars)
-+{
-+   nir_def *sbt_idx = nir_load_var(b, vars->idx);
-+
-+   if (!(vars->flags & VK_PIPELINE_CREATE_RAY_TRACING_NO_NULL_CLOSEST_HIT_SHADERS_BIT_KHR))
-+      nir_push_if(b, nir_ine_imm(b, sbt_idx, 0));
-+
-+   for (unsigned i = 0; i < pipeline->group_count; ++i) {
-+      struct radv_ray_tracing_group *group = &pipeline->groups[i];
-+
-+      unsigned shader_id = VK_SHADER_UNUSED_KHR;
-+      if (group->type != VK_RAY_TRACING_SHADER_GROUP_TYPE_GENERAL_KHR)
-+         shader_id = group->recursive_shader;
-+
-+      if (shader_id == VK_SHADER_UNUSED_KHR)
-+         continue;
-+
-+      /* Avoid emitting stages with the same shaders/handles multiple times. */
-+      bool is_dup = false;
-+      for (unsigned j = 0; j < i; ++j)
-+         if (pipeline->groups[j].handle.closest_hit_index == pipeline->groups[i].handle.closest_hit_index)
-+            is_dup = true;
-+
-+      if (is_dup)
-+         continue;
-+
-+      nir_shader *nir_stage = radv_pipeline_cache_handle_to_nir(device, pipeline->stages[shader_id].nir);
-+      assert(nir_stage);
-+
-+      insert_rt_case(b, nir_stage, vars, sbt_idx, pipeline->groups[i].handle.closest_hit_index);
-+      ralloc_free(nir_stage);
-+   }
-+
-+   if (!(vars->flags & VK_PIPELINE_CREATE_RAY_TRACING_NO_NULL_CLOSEST_HIT_SHADERS_BIT_KHR))
-+      nir_pop_if(b, NULL);
-+}
-+
-+static void
-+visit_miss_shaders(struct radv_device *device, nir_builder *b, struct radv_ray_tracing_pipeline *pipeline,
-+                   struct rt_variables *vars)
-+{
-+   nir_def *sbt_idx = nir_load_var(b, vars->idx);
-+
-+   if (!(vars->flags & VK_PIPELINE_CREATE_RAY_TRACING_NO_NULL_MISS_SHADERS_BIT_KHR))
-+      nir_push_if(b, nir_ine_imm(b, sbt_idx, 0));
-+
-+   for (unsigned i = 0; i < pipeline->group_count; ++i) {
-+      struct radv_ray_tracing_group *group = &pipeline->groups[i];
-+
-+      unsigned shader_id = VK_SHADER_UNUSED_KHR;
-+      if (group->type == VK_RAY_TRACING_SHADER_GROUP_TYPE_GENERAL_KHR)
-+         shader_id = group->recursive_shader;
-+
-+      if (shader_id == VK_SHADER_UNUSED_KHR)
-+         continue;
-+
-+      if (pipeline->stages[shader_id].stage != MESA_SHADER_MISS)
-+         continue;
-+
-+      /* Avoid emitting stages with the same shaders/handles multiple times. */
-+      bool is_dup = false;
-+      for (unsigned j = 0; j < i; ++j)
-+         if (pipeline->groups[j].handle.general_index == pipeline->groups[i].handle.general_index)
-+            is_dup = true;
-+
-+      if (is_dup)
-+         continue;
-+
-+      nir_shader *nir_stage = radv_pipeline_cache_handle_to_nir(device, pipeline->stages[shader_id].nir);
-+      assert(nir_stage);
-+
-+      insert_rt_case(b, nir_stage, vars, sbt_idx, pipeline->groups[i].handle.general_index);
-+      ralloc_free(nir_stage);
-+   }
-+
-+   if (!(vars->flags & VK_PIPELINE_CREATE_RAY_TRACING_NO_NULL_MISS_SHADERS_BIT_KHR))
-+      nir_pop_if(b, NULL);
-+}
-+
- static void
- store_stack_entry(nir_builder *b, nir_def *index, nir_def *value, const struct radv_ray_traversal_args *args)
- {
-@@ -1414,25 +1495,47 @@ radv_build_traversal(struct radv_device *device, struct radv_ray_tracing_pipelin
-    /* Register storage for hit attributes */
-    nir_variable *hit_attribs[RADV_MAX_HIT_ATTRIB_SIZE / sizeof(uint32_t)];
- 
--   for (uint32_t i = 0; i < ARRAY_SIZE(hit_attribs); i++)
--      hit_attribs[i] = nir_local_variable_create(nir_shader_get_entrypoint(b->shader), glsl_uint_type(), "ahit_attrib");
-+   if (!key->rt.monolithic || b->shader->info.stage != MESA_SHADER_RAYGEN) {
-+      for (uint32_t i = 0; i < ARRAY_SIZE(hit_attribs); i++)
-+         hit_attribs[i] =
-+            nir_local_variable_create(nir_shader_get_entrypoint(b->shader), glsl_uint_type(), "ahit_attrib");
- 
--   lower_hit_attribs(b->shader, hit_attribs, device->physical_device->rt_wave_size);
-+      lower_hit_attribs(b->shader, hit_attribs, device->physical_device->rt_wave_size);
-+   }
- 
-    /* Initialize follow-up shader. */
-    nir_push_if(b, nir_load_var(b, trav_vars.hit));
-    {
--      for (int i = 0; i < ARRAY_SIZE(hit_attribs); ++i)
--         nir_store_hit_attrib_amd(b, nir_load_var(b, hit_attribs[i]), .base = i);
--      nir_execute_closest_hit_amd(b, nir_load_var(b, vars->idx), nir_load_var(b, vars->tmax),
--                                  nir_load_var(b, vars->primitive_id), nir_load_var(b, vars->instance_addr),
--                                  nir_load_var(b, vars->geometry_id_and_flags), nir_load_var(b, vars->hit_kind));
-+      if (key->rt.monolithic && b->shader->info.stage == MESA_SHADER_RAYGEN) {
-+         load_sbt_entry(b, vars, nir_load_var(b, vars->idx), SBT_HIT, SBT_CLOSEST_HIT_IDX);
-+
-+         nir_def *should_return =
-+            nir_test_mask(b, nir_load_var(b, vars->cull_mask_and_flags), SpvRayFlagsSkipClosestHitShaderKHRMask);
-+
-+         /* should_return is set if we had a hit but we won't be calling the closest hit
-+          * shader and hence need to return immediately to the calling shader. */
-+         nir_push_if(b, nir_inot(b, should_return));
-+         visit_closest_hit_shaders(device, b, pipeline, vars);
-+         nir_pop_if(b, NULL);
-+      } else {
-+         for (int i = 0; i < ARRAY_SIZE(hit_attribs); ++i)
-+            nir_store_hit_attrib_amd(b, nir_load_var(b, hit_attribs[i]), .base = i);
-+         nir_execute_closest_hit_amd(b, nir_load_var(b, vars->idx), nir_load_var(b, vars->tmax),
-+                                     nir_load_var(b, vars->primitive_id), nir_load_var(b, vars->instance_addr),
-+                                     nir_load_var(b, vars->geometry_id_and_flags), nir_load_var(b, vars->hit_kind));
-+      }
-    }
-    nir_push_else(b, NULL);
-    {
--      /* Only load the miss shader if we actually miss. It is valid to not specify an SBT pointer
--       * for miss shaders if none of the rays miss. */
--      nir_execute_miss_amd(b, nir_load_var(b, vars->tmax));
-+      if (key->rt.monolithic && b->shader->info.stage == MESA_SHADER_RAYGEN) {
-+         load_sbt_entry(b, vars, nir_load_var(b, vars->miss_index), SBT_MISS, SBT_GENERAL_IDX);
-+
-+         visit_miss_shaders(device, b, pipeline, vars);
-+      } else {
-+         /* Only load the miss shader if we actually miss. It is valid to not specify an SBT pointer
-+          * for miss shaders if none of the rays miss. */
-+         nir_execute_miss_amd(b, nir_load_var(b, vars->tmax));
-+      }
-    }
-    nir_pop_if(b, NULL);
- }
-@@ -1477,6 +1580,98 @@ radv_build_traversal_shader(struct radv_device *device, struct radv_ray_tracing_
-    return b.shader;
- }
- 
-+struct lower_rt_instruction_monolithic_state {
-+   struct radv_device *device;
-+   struct radv_ray_tracing_pipeline *pipeline;
-+   const struct radv_pipeline_key *key;
-+   const VkRayTracingPipelineCreateInfoKHR *pCreateInfo;
-+
-+   struct rt_variables *vars;
-+};
-+
-+static bool
-+lower_rt_instruction_monolithic(nir_builder *b, nir_instr *instr, void *data)
-+{
-+   if (instr->type != nir_instr_type_intrinsic)
-+      return false;
-+
-+   b->cursor = nir_after_instr(instr);
-+
-+   nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
-+
-+   struct lower_rt_instruction_monolithic_state *state = data;
-+   struct rt_variables *vars = state->vars;
-+
-+   switch (intr->intrinsic) {
-+   case nir_intrinsic_execute_callable:
-+      unreachable("nir_intrinsic_execute_callable");
-+   case nir_intrinsic_trace_ray: {
-+      nir_store_var(b, vars->arg, nir_iadd_imm(b, intr->src[10].ssa, -b->shader->scratch_size), 1);
-+
-+      /* Per the SPIR-V extension spec we have to ignore some bits for some arguments. */
-+      nir_store_var(b, vars->accel_struct, intr->src[0].ssa, 0x1);
-+      nir_store_var(b, vars->cull_mask_and_flags, nir_ior(b, nir_ishl_imm(b, intr->src[2].ssa, 24), intr->src[1].ssa),
-+                    0x1);
-+      nir_store_var(b, vars->sbt_offset, nir_iand_imm(b, intr->src[3].ssa, 0xf), 0x1);
-+      nir_store_var(b, vars->sbt_stride, nir_iand_imm(b, intr->src[4].ssa, 0xf), 0x1);
-+      nir_store_var(b, vars->miss_index, nir_iand_imm(b, intr->src[5].ssa, 0xffff), 0x1);
-+      nir_store_var(b, vars->origin, intr->src[6].ssa, 0x7);
-+      nir_store_var(b, vars->tmin, intr->src[7].ssa, 0x1);
-+      nir_store_var(b, vars->direction, intr->src[8].ssa, 0x7);
-+      nir_store_var(b, vars->tmax, intr->src[9].ssa, 0x1);
-+
-+      nir_def *stack_ptr = nir_load_var(b, vars->stack_ptr);
-+      nir_store_var(b, vars->stack_ptr, nir_iadd_imm(b, stack_ptr, b->shader->scratch_size), 0x1);
-+
-+      radv_build_traversal(state->device, state->pipeline, state->pCreateInfo, state->key, b, vars);
-+      b->shader->info.shared_size = MAX2(b->shader->info.shared_size, state->device->physical_device->rt_wave_size *
-+                                                                         MAX_STACK_ENTRY_COUNT * sizeof(uint32_t));
-+
-+      nir_store_var(b, vars->stack_ptr, stack_ptr, 0x1);
-+
-+      nir_instr_remove(instr);
-+      return true;
-+   }
-+   case nir_intrinsic_rt_resume:
-+      unreachable("nir_intrinsic_rt_resume");
-+   case nir_intrinsic_rt_return_amd:
-+      unreachable("nir_intrinsic_rt_return_amd");
-+   case nir_intrinsic_execute_closest_hit_amd:
-+      unreachable("nir_intrinsic_execute_closest_hit_amd");
-+   case nir_intrinsic_execute_miss_amd:
-+      unreachable("nir_intrinsic_execute_miss_amd");
-+   default:
-+      return false;
-+   }
-+}
-+
-+static void
-+lower_rt_instructions_monolithic(nir_shader *shader, struct radv_device *device,
-+                                 struct radv_ray_tracing_pipeline *pipeline, const struct radv_pipeline_key *key,
-+                                 const VkRayTracingPipelineCreateInfoKHR *pCreateInfo, struct rt_variables *vars)
-+{
-+   nir_function_impl *impl = nir_shader_get_entrypoint(shader);
-+
-+   struct lower_rt_instruction_monolithic_state state = {
-+      .device = device,
-+      .pipeline = pipeline,
-+      .key = key,
-+      .pCreateInfo = pCreateInfo,
-+      .vars = vars,
-+   };
-+
-+   nir_shader_instructions_pass(shader, lower_rt_instruction_monolithic, nir_metadata_none, &state);
-+   nir_index_ssa_defs(impl);
-+
-+   /* Register storage for hit attributes */
-+   nir_variable *hit_attribs[RADV_MAX_HIT_ATTRIB_SIZE / sizeof(uint32_t)];
-+
-+   for (uint32_t i = 0; i < ARRAY_SIZE(hit_attribs); i++)
-+      hit_attribs[i] = nir_local_variable_create(impl, glsl_uint_type(), "ahit_attrib");
-+
-+   lower_hit_attribs(shader, hit_attribs, 0);
-+}
-+
- /** Select the next shader based on priorities:
-  *
-  * Detect the priority of the shader stage by the lowest bits in the address (low to high):
-@@ -1517,13 +1712,18 @@ select_next_shader(nir_builder *b, nir_def *shader_addr, unsigned wave_size)
- void
- radv_nir_lower_rt_abi(nir_shader *shader, const VkRayTracingPipelineCreateInfoKHR *pCreateInfo,
-                       const struct radv_shader_args *args, const struct radv_shader_info *info, uint32_t *stack_size,
--                      bool resume_shader)
-+                      bool resume_shader, struct radv_device *device, struct radv_ray_tracing_pipeline *pipeline,
-+                      const struct radv_pipeline_key *key)
- {
-    nir_function_impl *impl = nir_shader_get_entrypoint(shader);
- 
-    const VkPipelineCreateFlagBits2KHR create_flags = radv_get_pipeline_create_flags(pCreateInfo);
- 
-    struct rt_variables vars = create_rt_variables(shader, create_flags);
-+
-+   if (key->rt.monolithic && shader->info.stage == MESA_SHADER_RAYGEN)
-+      lower_rt_instructions_monolithic(shader, device, pipeline, key, pCreateInfo, &vars);
-+
-    lower_rt_instructions(shader, &vars, true);
- 
-    if (stack_size) {
-@@ -1585,32 +1785,36 @@ radv_nir_lower_rt_abi(nir_shader *shader, const VkRayTracingPipelineCreateInfoKH
-    if (shader_guard)
-       nir_pop_if(&b, shader_guard);
- 
--   /* select next shader */
-    b.cursor = nir_after_cf_list(&impl->body);
- 
--   shader_addr = nir_load_var(&b, vars.shader_addr);
--   nir_def *next = select_next_shader(&b, shader_addr, info->wave_size);
--   ac_nir_store_arg(&b, &args->ac, args->ac.rt.uniform_shader_addr, next);
--
--   /* store back all variables to registers */
--   ac_nir_store_arg(&b, &args->ac, args->ac.rt.dynamic_callable_stack_base, nir_load_var(&b, vars.stack_ptr));
--   ac_nir_store_arg(&b, &args->ac, args->ac.rt.shader_addr, shader_addr);
--   ac_nir_store_arg(&b, &args->ac, args->ac.rt.shader_record, nir_load_var(&b, vars.shader_record_ptr));
--   ac_nir_store_arg(&b, &args->ac, args->ac.rt.payload_offset, nir_load_var(&b, vars.arg));
--   ac_nir_store_arg(&b, &args->ac, args->ac.rt.accel_struct, nir_load_var(&b, vars.accel_struct));
--   ac_nir_store_arg(&b, &args->ac, args->ac.rt.cull_mask_and_flags, nir_load_var(&b, vars.cull_mask_and_flags));
--   ac_nir_store_arg(&b, &args->ac, args->ac.rt.sbt_offset, nir_load_var(&b, vars.sbt_offset));
--   ac_nir_store_arg(&b, &args->ac, args->ac.rt.sbt_stride, nir_load_var(&b, vars.sbt_stride));
--   ac_nir_store_arg(&b, &args->ac, args->ac.rt.miss_index, nir_load_var(&b, vars.miss_index));
--   ac_nir_store_arg(&b, &args->ac, args->ac.rt.ray_origin, nir_load_var(&b, vars.origin));
--   ac_nir_store_arg(&b, &args->ac, args->ac.rt.ray_tmin, nir_load_var(&b, vars.tmin));
--   ac_nir_store_arg(&b, &args->ac, args->ac.rt.ray_direction, nir_load_var(&b, vars.direction));
--   ac_nir_store_arg(&b, &args->ac, args->ac.rt.ray_tmax, nir_load_var(&b, vars.tmax));
--
--   ac_nir_store_arg(&b, &args->ac, args->ac.rt.primitive_id, nir_load_var(&b, vars.primitive_id));
--   ac_nir_store_arg(&b, &args->ac, args->ac.rt.instance_addr, nir_load_var(&b, vars.instance_addr));
--   ac_nir_store_arg(&b, &args->ac, args->ac.rt.geometry_id_and_flags, nir_load_var(&b, vars.geometry_id_and_flags));
--   ac_nir_store_arg(&b, &args->ac, args->ac.rt.hit_kind, nir_load_var(&b, vars.hit_kind));
-+   if (key->rt.monolithic && shader->info.stage == MESA_SHADER_RAYGEN) {
-+      nir_terminate(&b);
-+   } else {
-+      /* select next shader */
-+      shader_addr = nir_load_var(&b, vars.shader_addr);
-+      nir_def *next = select_next_shader(&b, shader_addr, info->wave_size);
-+      ac_nir_store_arg(&b, &args->ac, args->ac.rt.uniform_shader_addr, next);
-+
-+      /* store back all variables to registers */
-+      ac_nir_store_arg(&b, &args->ac, args->ac.rt.dynamic_callable_stack_base, nir_load_var(&b, vars.stack_ptr));
-+      ac_nir_store_arg(&b, &args->ac, args->ac.rt.shader_addr, shader_addr);
-+      ac_nir_store_arg(&b, &args->ac, args->ac.rt.shader_record, nir_load_var(&b, vars.shader_record_ptr));
-+      ac_nir_store_arg(&b, &args->ac, args->ac.rt.payload_offset, nir_load_var(&b, vars.arg));
-+      ac_nir_store_arg(&b, &args->ac, args->ac.rt.accel_struct, nir_load_var(&b, vars.accel_struct));
-+      ac_nir_store_arg(&b, &args->ac, args->ac.rt.cull_mask_and_flags, nir_load_var(&b, vars.cull_mask_and_flags));
-+      ac_nir_store_arg(&b, &args->ac, args->ac.rt.sbt_offset, nir_load_var(&b, vars.sbt_offset));
-+      ac_nir_store_arg(&b, &args->ac, args->ac.rt.sbt_stride, nir_load_var(&b, vars.sbt_stride));
-+      ac_nir_store_arg(&b, &args->ac, args->ac.rt.miss_index, nir_load_var(&b, vars.miss_index));
-+      ac_nir_store_arg(&b, &args->ac, args->ac.rt.ray_origin, nir_load_var(&b, vars.origin));
-+      ac_nir_store_arg(&b, &args->ac, args->ac.rt.ray_tmin, nir_load_var(&b, vars.tmin));
-+      ac_nir_store_arg(&b, &args->ac, args->ac.rt.ray_direction, nir_load_var(&b, vars.direction));
-+      ac_nir_store_arg(&b, &args->ac, args->ac.rt.ray_tmax, nir_load_var(&b, vars.tmax));
-+
-+      ac_nir_store_arg(&b, &args->ac, args->ac.rt.primitive_id, nir_load_var(&b, vars.primitive_id));
-+      ac_nir_store_arg(&b, &args->ac, args->ac.rt.instance_addr, nir_load_var(&b, vars.instance_addr));
-+      ac_nir_store_arg(&b, &args->ac, args->ac.rt.geometry_id_and_flags, nir_load_var(&b, vars.geometry_id_and_flags));
-+      ac_nir_store_arg(&b, &args->ac, args->ac.rt.hit_kind, nir_load_var(&b, vars.hit_kind));
-+   }
- 
-    nir_metadata_preserve(impl, nir_metadata_none);
- 
-diff --git a/src/amd/vulkan/radv_shader.h b/src/amd/vulkan/radv_shader.h
-index 6eb95fdd0a097..969f9a56ab7f5 100644
---- a/src/amd/vulkan/radv_shader.h
-+++ b/src/amd/vulkan/radv_shader.h
-@@ -635,7 +635,8 @@ nir_shader *radv_parse_rt_stage(struct radv_device *device, const VkPipelineShad
- 
- void radv_nir_lower_rt_abi(nir_shader *shader, const VkRayTracingPipelineCreateInfoKHR *pCreateInfo,
-                            const struct radv_shader_args *args, const struct radv_shader_info *info,
--                           uint32_t *stack_size, bool resume_shader);
-+                           uint32_t *stack_size, bool resume_shader, struct radv_device *device,
-+                           struct radv_ray_tracing_pipeline *pipeline, const struct radv_pipeline_key *key);
- 
- struct radv_shader_stage;
- 
--- 
-
-
-From 5c9dd4efece8f352d00d1310b556928cccb239c8 Mon Sep 17 00:00:00 2001
-From: Konstantin Seurer <konstantin.seurer@gmail.com>
-Date: Sat, 24 Jun 2023 16:11:16 +0200
-Subject: [PATCH 4/4] radv/rt: Use monolithic pipelines
-
-Only available for non-recursive pipelines that do not have callables.
----
- src/amd/vulkan/radv_pipeline_rt.c | 11 +++++++++++
- 1 file changed, 11 insertions(+)
-
-diff --git a/src/amd/vulkan/radv_pipeline_rt.c b/src/amd/vulkan/radv_pipeline_rt.c
-index 12562c6cf89ba..97449b9cbafac 100644
---- a/src/amd/vulkan/radv_pipeline_rt.c
-+++ b/src/amd/vulkan/radv_pipeline_rt.c
-@@ -103,6 +103,17 @@ radv_generate_rt_pipeline_key(const struct radv_device *device, const struct rad
-       }
-    }
- 
-+   if (!(pCreateInfo->flags & VK_PIPELINE_CREATE_LIBRARY_BIT_KHR)) {
-+      key.rt.monolithic = pCreateInfo->maxPipelineRayRecursionDepth <= 1;
-+
-+      for (uint32_t i = 0; i < pipeline->stage_count; i++) {
-+         if (pipeline->stages[i].stage == MESA_SHADER_CALLABLE) {
-+            key.rt.monolithic = false;
-+            break;
-+         }
-+      }
-+   }
-+
-    return key;
- }
- 
--- 

debian/patches/24720.patch

@@ -1,687 +0,0 @@
-From 42be7a3c53698a165e9612619f6a34a65bbf91ff Mon Sep 17 00:00:00 2001
-From: Konstantin Seurer <konstantin.seurer@gmail.com>
-Date: Wed, 16 Aug 2023 10:37:56 +0200
-Subject: [PATCH 1/3] radv: Remove dead radix_sort_vk_get_memory_requirements
- call
-
----
- src/amd/vulkan/radv_acceleration_structure.c | 4 ----
- 1 file changed, 4 deletions(-)
-
-diff --git a/src/amd/vulkan/radv_acceleration_structure.c b/src/amd/vulkan/radv_acceleration_structure.c
-index ece47b1230c88..5c5eb16e61f9d 100644
---- a/src/amd/vulkan/radv_acceleration_structure.c
-+++ b/src/amd/vulkan/radv_acceleration_structure.c
-@@ -745,10 +745,6 @@ morton_sort(VkCommandBuffer commandBuffer, uint32_t infoCount,
- {
-    RADV_FROM_HANDLE(radv_cmd_buffer, cmd_buffer, commandBuffer);
-    for (uint32_t i = 0; i < infoCount; ++i) {
--      struct radix_sort_vk_memory_requirements requirements;
--      radix_sort_vk_get_memory_requirements(cmd_buffer->device->meta_state.accel_struct_build.radix_sort,
--                                            bvh_states[i].node_count, &requirements);
--
-       struct radix_sort_vk_sort_devaddr_info info = cmd_buffer->device->meta_state.accel_struct_build.radix_sort_info;
-       info.count = bvh_states[i].node_count;
- 
--- 
-GitLab
-
-
-From faa17e5322ea66cd74e37aab48316059a05738d6 Mon Sep 17 00:00:00 2001
-From: Konstantin Seurer <konstantin.seurer@gmail.com>
-Date: Wed, 16 Aug 2023 11:09:25 +0200
-Subject: [PATCH 2/3] radv/radix_sort: Vendor the radix sort dispatch code
-
-This needs to be done so we can optimize it for occpuancy when building
-multiple acceleration structures in parallel. Changes to the original
-code:
-
-- Change // to /* */
-- clang-format
-- Replace vkCmd calls with calls to the driver entrypoints
-- Add a light weight info struct
-- Use radv_fill_buffer directly
----
- src/amd/vulkan/radv_acceleration_structure.c | 218 ++++++++++++++++---
- src/amd/vulkan/radv_private.h                |   1 -
- 2 files changed, 187 insertions(+), 32 deletions(-)
-
-diff --git a/src/amd/vulkan/radv_acceleration_structure.c b/src/amd/vulkan/radv_acceleration_structure.c
-index 5c5eb16e61f9d..9866de2e594a8 100644
---- a/src/amd/vulkan/radv_acceleration_structure.c
-+++ b/src/amd/vulkan/radv_acceleration_structure.c
-@@ -27,7 +27,9 @@
- #include "nir_builder.h"
- #include "radv_cs.h"
- 
-+#include "radix_sort/common/vk/barrier.h"
- #include "radix_sort/radv_radix_sort.h"
-+#include "radix_sort/shaders/push.h"
- 
- #include "bvh/build_interface.h"
- #include "bvh/bvh.h"
-@@ -76,6 +78,7 @@ static const uint32_t header_spv[] = {
- };
- 
- #define KEY_ID_PAIR_SIZE 8
-+#define MORTON_BIT_SIZE  24
- 
- enum internal_build_type {
-    INTERNAL_BUILD_TYPE_LBVH,
-@@ -382,17 +385,6 @@ cleanup:
-    return result;
- }
- 
--static void
--radix_sort_fill_buffer(VkCommandBuffer commandBuffer, radix_sort_vk_buffer_info_t const *buffer_info,
--                       VkDeviceSize offset, VkDeviceSize size, uint32_t data)
--{
--   RADV_FROM_HANDLE(radv_cmd_buffer, cmd_buffer, commandBuffer);
--
--   assert(size != VK_WHOLE_SIZE);
--
--   radv_fill_buffer(cmd_buffer, NULL, NULL, buffer_info->devaddr + buffer_info->offset + offset, size, data);
--}
--
- VkResult
- radv_device_init_null_accel_struct(struct radv_device *device)
- {
-@@ -576,12 +568,6 @@ radv_device_init_accel_struct_build_state(struct radv_device *device)
- 
-    device->meta_state.accel_struct_build.radix_sort =
-       radv_create_radix_sort_u64(radv_device_to_handle(device), &device->meta_state.alloc, device->meta_state.cache);
--
--   struct radix_sort_vk_sort_devaddr_info *radix_sort_info = &device->meta_state.accel_struct_build.radix_sort_info;
--   radix_sort_info->ext = NULL;
--   radix_sort_info->key_bits = 24;
--   radix_sort_info->fill_buffer = radix_sort_fill_buffer;
--
- exit:
-    mtx_unlock(&device->meta_state.mtx);
-    return result;
-@@ -743,28 +729,198 @@ morton_sort(VkCommandBuffer commandBuffer, uint32_t infoCount,
-             const VkAccelerationStructureBuildGeometryInfoKHR *pInfos, struct bvh_state *bvh_states,
-             enum radv_cmd_flush_bits flush_bits)
- {
-+   /* Copyright 2019 The Fuchsia Authors. */
-    RADV_FROM_HANDLE(radv_cmd_buffer, cmd_buffer, commandBuffer);
-+
-+   radix_sort_vk_t *rs = cmd_buffer->device->meta_state.accel_struct_build.radix_sort;
-+
-    for (uint32_t i = 0; i < infoCount; ++i) {
--      struct radix_sort_vk_sort_devaddr_info info = cmd_buffer->device->meta_state.accel_struct_build.radix_sort_info;
--      info.count = bvh_states[i].node_count;
-+      uint32_t count = bvh_states[i].node_count;
-+      uint64_t keyvals_even_addr = pInfos[i].scratchData.deviceAddress + bvh_states[i].scratch.sort_buffer_offset[0];
-+      uint64_t keyvals_odd_addr = pInfos[i].scratchData.deviceAddress + bvh_states[i].scratch.sort_buffer_offset[1];
-+      uint64_t internal_addr = pInfos[i].scratchData.deviceAddress + bvh_states[i].scratch.sort_internal_offset;
-+
-+      /* Anything to do? */
-+      if (!count) {
-+         bvh_states[i].scratch_offset = bvh_states[i].scratch.sort_buffer_offset[0];
-+         continue;
-+      }
-+
-+      /*
-+       * OVERVIEW
-+       *
-+       *   1. Pad the keyvals in `scatter_even`.
-+       *   2. Zero the `histograms` and `partitions`.
-+       *      --- BARRIER ---
-+       *   3. HISTOGRAM is dispatched before PREFIX.
-+       *      --- BARRIER ---
-+       *   4. PREFIX is dispatched before the first SCATTER.
-+       *      --- BARRIER ---
-+       *   5. One or more SCATTER dispatches.
-+       *
-+       * Note that the `partitions` buffer can be zeroed anytime before the first
-+       * scatter.
-+       */
-+
-+      /* How many passes? */
-+      uint32_t keyval_bytes = rs->config.keyval_dwords * (uint32_t)sizeof(uint32_t);
-+      uint32_t keyval_bits = keyval_bytes * 8;
-+      uint32_t key_bits = MIN2(MORTON_BIT_SIZE, keyval_bits);
-+      uint32_t passes = (key_bits + RS_RADIX_LOG2 - 1) / RS_RADIX_LOG2;
-+
-+      bvh_states[i].scratch_offset = bvh_states[i].scratch.sort_buffer_offset[passes & 1];
-+
-+      /*
-+       * PAD KEYVALS AND ZERO HISTOGRAM/PARTITIONS
-+       *
-+       * Pad fractional blocks with max-valued keyvals.
-+       *
-+       * Zero the histograms and partitions buffer.
-+       *
-+       * This assumes the partitions follow the histograms.
-+       */
-+
-+      /* FIXME(allanmac): Consider precomputing some of these values and hang them off `rs`. */
-+
-+      /* How many scatter blocks? */
-+      uint32_t scatter_wg_size = 1 << rs->config.scatter.workgroup_size_log2;
-+      uint32_t scatter_block_kvs = scatter_wg_size * rs->config.scatter.block_rows;
-+      uint32_t scatter_blocks = (count + scatter_block_kvs - 1) / scatter_block_kvs;
-+      uint32_t count_ru_scatter = scatter_blocks * scatter_block_kvs;
-+
-+      /*
-+       * How many histogram blocks?
-+       *
-+       * Note that it's OK to have more max-valued digits counted by the histogram
-+       * than sorted by the scatters because the sort is stable.
-+       */
-+      uint32_t histo_wg_size = 1 << rs->config.histogram.workgroup_size_log2;
-+      uint32_t histo_block_kvs = histo_wg_size * rs->config.histogram.block_rows;
-+      uint32_t histo_blocks = (count_ru_scatter + histo_block_kvs - 1) / histo_block_kvs;
-+      uint32_t count_ru_histo = histo_blocks * histo_block_kvs;
-+
-+      /* Fill with max values */
-+      if (count_ru_histo > count) {
-+         radv_fill_buffer(cmd_buffer, NULL, NULL, keyvals_even_addr + count * keyval_bytes,
-+                          (count_ru_histo - count) * keyval_bytes, 0xFFFFFFFF);
-+      }
-+
-+      /*
-+       * Zero histograms and invalidate partitions.
-+       *
-+       * Note that the partition invalidation only needs to be performed once
-+       * because the even/odd scatter dispatches rely on the the previous pass to
-+       * leave the partitions in an invalid state.
-+       *
-+       * Note that the last workgroup doesn't read/write a partition so it doesn't
-+       * need to be initialized.
-+       */
-+      uint32_t histo_partition_count = passes + scatter_blocks - 1;
-+      uint32_t pass_idx = (keyval_bytes - passes);
-+
-+      uint32_t fill_base = pass_idx * (RS_RADIX_SIZE * sizeof(uint32_t));
-+
-+      radv_fill_buffer(cmd_buffer, NULL, NULL, internal_addr + rs->internal.histograms.offset + fill_base,
-+                       histo_partition_count * (RS_RADIX_SIZE * sizeof(uint32_t)), 0);
-+
-+      /*
-+       * Pipeline: HISTOGRAM
-+       *
-+       * TODO(allanmac): All subgroups should try to process approximately the same
-+       * number of blocks in order to minimize tail effects.  This was implemented
-+       * and reverted but should be reimplemented and benchmarked later.
-+       */
-+      vk_barrier_transfer_w_to_compute_r(commandBuffer);
-+
-+      uint64_t devaddr_histograms = internal_addr + rs->internal.histograms.offset;
-+
-+      /* Dispatch histogram */
-+      struct rs_push_histogram push_histogram = {
-+         .devaddr_histograms = devaddr_histograms,
-+         .devaddr_keyvals = keyvals_even_addr,
-+         .passes = passes,
-+      };
-+
-+      radv_CmdPushConstants(commandBuffer, rs->pipeline_layouts.named.histogram, VK_SHADER_STAGE_COMPUTE_BIT, 0,
-+                            sizeof(push_histogram), &push_histogram);
-+
-+      radv_CmdBindPipeline(commandBuffer, VK_PIPELINE_BIND_POINT_COMPUTE, rs->pipelines.named.histogram);
-+
-+      vk_common_CmdDispatch(commandBuffer, histo_blocks, 1, 1);
- 
--      info.keyvals_even.buffer = VK_NULL_HANDLE;
--      info.keyvals_even.offset = 0;
--      info.keyvals_even.devaddr = pInfos[i].scratchData.deviceAddress + bvh_states[i].scratch.sort_buffer_offset[0];
-+      /*
-+       * Pipeline: PREFIX
-+       *
-+       * Launch one workgroup per pass.
-+       */
-+      vk_barrier_compute_w_to_compute_r(commandBuffer);
-+
-+      struct rs_push_prefix push_prefix = {
-+         .devaddr_histograms = devaddr_histograms,
-+      };
-+
-+      radv_CmdPushConstants(commandBuffer, rs->pipeline_layouts.named.prefix, VK_SHADER_STAGE_COMPUTE_BIT, 0,
-+                            sizeof(push_prefix), &push_prefix);
- 
--      info.keyvals_odd = pInfos[i].scratchData.deviceAddress + bvh_states[i].scratch.sort_buffer_offset[1];
-+      radv_CmdBindPipeline(commandBuffer, VK_PIPELINE_BIND_POINT_COMPUTE, rs->pipelines.named.prefix);
- 
--      info.internal.buffer = VK_NULL_HANDLE;
--      info.internal.offset = 0;
--      info.internal.devaddr = pInfos[i].scratchData.deviceAddress + bvh_states[i].scratch.sort_internal_offset;
-+      vk_common_CmdDispatch(commandBuffer, passes, 1, 1);
- 
--      VkDeviceAddress result_addr;
--      radix_sort_vk_sort_devaddr(cmd_buffer->device->meta_state.accel_struct_build.radix_sort, &info,
--                                 radv_device_to_handle(cmd_buffer->device), commandBuffer, &result_addr);
-+      /* Pipeline: SCATTER */
-+      vk_barrier_compute_w_to_compute_r(commandBuffer);
-+
-+      uint32_t histogram_offset = pass_idx * (RS_RADIX_SIZE * sizeof(uint32_t));
-+      uint64_t devaddr_partitions = internal_addr + rs->internal.partitions.offset;
-+
-+      struct rs_push_scatter push_scatter = {
-+         .devaddr_keyvals_even = keyvals_even_addr,
-+         .devaddr_keyvals_odd = keyvals_odd_addr,
-+         .devaddr_partitions = devaddr_partitions,
-+         .devaddr_histograms = devaddr_histograms + histogram_offset,
-+         .pass_offset = (pass_idx & 3) * RS_RADIX_LOG2,
-+      };
- 
--      assert(result_addr == info.keyvals_even.devaddr || result_addr == info.keyvals_odd);
-+      {
-+         uint32_t pass_dword = pass_idx / 4;
- 
--      bvh_states[i].scratch_offset = (uint32_t)(result_addr - pInfos[i].scratchData.deviceAddress);
-+         radv_CmdPushConstants(commandBuffer, rs->pipeline_layouts.named.scatter[pass_dword].even,
-+                               VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(push_scatter), &push_scatter);
-+
-+         radv_CmdBindPipeline(commandBuffer, VK_PIPELINE_BIND_POINT_COMPUTE,
-+                              rs->pipelines.named.scatter[pass_dword].even);
-+      }
-+
-+      bool is_even = true;
-+
-+      while (true) {
-+         vk_common_CmdDispatch(commandBuffer, scatter_blocks, 1, 1);
-+
-+         /* Continue? */
-+         if (++pass_idx >= keyval_bytes)
-+            break;
-+
-+         vk_barrier_compute_w_to_compute_r(commandBuffer);
-+
-+         is_even ^= true;
-+         push_scatter.devaddr_histograms += (RS_RADIX_SIZE * sizeof(uint32_t));
-+         push_scatter.pass_offset = (pass_idx & 3) * RS_RADIX_LOG2;
-+
-+         uint32_t pass_dword = pass_idx / 4;
-+
-+         /* Update push constants that changed */
-+         VkPipelineLayout pl = is_even ? rs->pipeline_layouts.named.scatter[pass_dword].even
-+                                       : rs->pipeline_layouts.named.scatter[pass_dword].odd;
-+         radv_CmdPushConstants(commandBuffer, pl, VK_SHADER_STAGE_COMPUTE_BIT,
-+                               offsetof(struct rs_push_scatter, devaddr_histograms),
-+                               sizeof(push_scatter.devaddr_histograms) + sizeof(push_scatter.pass_offset),
-+                               &push_scatter.devaddr_histograms);
-+
-+         /* Bind new pipeline */
-+         VkPipeline p =
-+            is_even ? rs->pipelines.named.scatter[pass_dword].even : rs->pipelines.named.scatter[pass_dword].odd;
-+
-+         radv_CmdBindPipeline(commandBuffer, VK_PIPELINE_BIND_POINT_COMPUTE, p);
-+      }
-    }
- 
-    cmd_buffer->state.flush_bits |= flush_bits;
-diff --git a/src/amd/vulkan/radv_private.h b/src/amd/vulkan/radv_private.h
-index 1ea606c2ca111..2b0b9189f33f6 100644
---- a/src/amd/vulkan/radv_private.h
-+++ b/src/amd/vulkan/radv_private.h
-@@ -731,7 +731,6 @@ struct radv_meta_state {
-       VkPipeline copy_pipeline;
- 
-       struct radix_sort_vk *radix_sort;
--      struct radix_sort_vk_sort_devaddr_info radix_sort_info;
- 
-       struct {
-          VkBuffer buffer;
--- 
-GitLab
-
-
-From 04c77145628fe9956ae44a25ba7b1dfe401a9de8 Mon Sep 17 00:00:00 2001
-From: Konstantin Seurer <konstantin.seurer@gmail.com>
-Date: Wed, 16 Aug 2023 11:50:18 +0200
-Subject: [PATCH 3/3] radv: Perform multiple sorts in parallel
-
-This was the last part that didn't scale with multiple infos. Reducing
-the amount of barriers in this case improves DOOM Eternal performance by
-50%. (Running with low resolution)
----
- src/amd/vulkan/radv_acceleration_structure.c | 264 ++++++++++---------
- 1 file changed, 143 insertions(+), 121 deletions(-)
-
-diff --git a/src/amd/vulkan/radv_acceleration_structure.c b/src/amd/vulkan/radv_acceleration_structure.c
-index 9866de2e594a8..85852453fcb29 100644
---- a/src/amd/vulkan/radv_acceleration_structure.c
-+++ b/src/amd/vulkan/radv_acceleration_structure.c
-@@ -598,6 +598,13 @@ struct bvh_state {
-    struct acceleration_structure_layout accel_struct;
-    struct scratch_layout scratch;
-    struct build_config config;
-+
-+   /* Radix sort state */
-+   uint32_t scatter_blocks;
-+   uint32_t count_ru_scatter;
-+   uint32_t histo_blocks;
-+   uint32_t count_ru_histo;
-+   struct rs_push_scatter push_scatter;
- };
- 
- static uint32_t
-@@ -734,75 +741,79 @@ morton_sort(VkCommandBuffer commandBuffer, uint32_t infoCount,
- 
-    radix_sort_vk_t *rs = cmd_buffer->device->meta_state.accel_struct_build.radix_sort;
- 
--   for (uint32_t i = 0; i < infoCount; ++i) {
--      uint32_t count = bvh_states[i].node_count;
--      uint64_t keyvals_even_addr = pInfos[i].scratchData.deviceAddress + bvh_states[i].scratch.sort_buffer_offset[0];
--      uint64_t keyvals_odd_addr = pInfos[i].scratchData.deviceAddress + bvh_states[i].scratch.sort_buffer_offset[1];
--      uint64_t internal_addr = pInfos[i].scratchData.deviceAddress + bvh_states[i].scratch.sort_internal_offset;
-+   /*
-+    * OVERVIEW
-+    *
-+    *   1. Pad the keyvals in `scatter_even`.
-+    *   2. Zero the `histograms` and `partitions`.
-+    *      --- BARRIER ---
-+    *   3. HISTOGRAM is dispatched before PREFIX.
-+    *      --- BARRIER ---
-+    *   4. PREFIX is dispatched before the first SCATTER.
-+    *      --- BARRIER ---
-+    *   5. One or more SCATTER dispatches.
-+    *
-+    * Note that the `partitions` buffer can be zeroed anytime before the first
-+    * scatter.
-+    */
-+
-+   /* How many passes? */
-+   uint32_t keyval_bytes = rs->config.keyval_dwords * (uint32_t)sizeof(uint32_t);
-+   uint32_t keyval_bits = keyval_bytes * 8;
-+   uint32_t key_bits = MIN2(MORTON_BIT_SIZE, keyval_bits);
-+   uint32_t passes = (key_bits + RS_RADIX_LOG2 - 1) / RS_RADIX_LOG2;
- 
--      /* Anything to do? */
--      if (!count) {
-+   for (uint32_t i = 0; i < infoCount; ++i) {
-+      if (bvh_states[i].node_count)
-+         bvh_states[i].scratch_offset = bvh_states[i].scratch.sort_buffer_offset[passes & 1];
-+      else
-          bvh_states[i].scratch_offset = bvh_states[i].scratch.sort_buffer_offset[0];
--         continue;
--      }
--
--      /*
--       * OVERVIEW
--       *
--       *   1. Pad the keyvals in `scatter_even`.
--       *   2. Zero the `histograms` and `partitions`.
--       *      --- BARRIER ---
--       *   3. HISTOGRAM is dispatched before PREFIX.
--       *      --- BARRIER ---
--       *   4. PREFIX is dispatched before the first SCATTER.
--       *      --- BARRIER ---
--       *   5. One or more SCATTER dispatches.
--       *
--       * Note that the `partitions` buffer can be zeroed anytime before the first
--       * scatter.
--       */
--
--      /* How many passes? */
--      uint32_t keyval_bytes = rs->config.keyval_dwords * (uint32_t)sizeof(uint32_t);
--      uint32_t keyval_bits = keyval_bytes * 8;
--      uint32_t key_bits = MIN2(MORTON_BIT_SIZE, keyval_bits);
--      uint32_t passes = (key_bits + RS_RADIX_LOG2 - 1) / RS_RADIX_LOG2;
-+   }
- 
--      bvh_states[i].scratch_offset = bvh_states[i].scratch.sort_buffer_offset[passes & 1];
-+   /*
-+    * PAD KEYVALS AND ZERO HISTOGRAM/PARTITIONS
-+    *
-+    * Pad fractional blocks with max-valued keyvals.
-+    *
-+    * Zero the histograms and partitions buffer.
-+    *
-+    * This assumes the partitions follow the histograms.
-+    */
-+
-+   /* FIXME(allanmac): Consider precomputing some of these values and hang them off `rs`. */
-+
-+   /* How many scatter blocks? */
-+   uint32_t scatter_wg_size = 1 << rs->config.scatter.workgroup_size_log2;
-+   uint32_t scatter_block_kvs = scatter_wg_size * rs->config.scatter.block_rows;
-+
-+   /*
-+    * How many histogram blocks?
-+    *
-+    * Note that it's OK to have more max-valued digits counted by the histogram
-+    * than sorted by the scatters because the sort is stable.
-+    */
-+   uint32_t histo_wg_size = 1 << rs->config.histogram.workgroup_size_log2;
-+   uint32_t histo_block_kvs = histo_wg_size * rs->config.histogram.block_rows;
-+
-+   uint32_t pass_idx = (keyval_bytes - passes);
- 
--      /*
--       * PAD KEYVALS AND ZERO HISTOGRAM/PARTITIONS
--       *
--       * Pad fractional blocks with max-valued keyvals.
--       *
--       * Zero the histograms and partitions buffer.
--       *
--       * This assumes the partitions follow the histograms.
--       */
-+   for (uint32_t i = 0; i < infoCount; ++i) {
-+      if (!bvh_states[i].node_count)
-+         continue;
- 
--      /* FIXME(allanmac): Consider precomputing some of these values and hang them off `rs`. */
-+      uint64_t keyvals_even_addr = pInfos[i].scratchData.deviceAddress + bvh_states[i].scratch.sort_buffer_offset[0];
-+      uint64_t internal_addr = pInfos[i].scratchData.deviceAddress + bvh_states[i].scratch.sort_internal_offset;
- 
--      /* How many scatter blocks? */
--      uint32_t scatter_wg_size = 1 << rs->config.scatter.workgroup_size_log2;
--      uint32_t scatter_block_kvs = scatter_wg_size * rs->config.scatter.block_rows;
--      uint32_t scatter_blocks = (count + scatter_block_kvs - 1) / scatter_block_kvs;
--      uint32_t count_ru_scatter = scatter_blocks * scatter_block_kvs;
-+      bvh_states[i].scatter_blocks = (bvh_states[i].node_count + scatter_block_kvs - 1) / scatter_block_kvs;
-+      bvh_states[i].count_ru_scatter = bvh_states[i].scatter_blocks * scatter_block_kvs;
- 
--      /*
--       * How many histogram blocks?
--       *
--       * Note that it's OK to have more max-valued digits counted by the histogram
--       * than sorted by the scatters because the sort is stable.
--       */
--      uint32_t histo_wg_size = 1 << rs->config.histogram.workgroup_size_log2;
--      uint32_t histo_block_kvs = histo_wg_size * rs->config.histogram.block_rows;
--      uint32_t histo_blocks = (count_ru_scatter + histo_block_kvs - 1) / histo_block_kvs;
--      uint32_t count_ru_histo = histo_blocks * histo_block_kvs;
-+      bvh_states[i].histo_blocks = (bvh_states[i].count_ru_scatter + histo_block_kvs - 1) / histo_block_kvs;
-+      bvh_states[i].count_ru_histo = bvh_states[i].histo_blocks * histo_block_kvs;
- 
-       /* Fill with max values */
--      if (count_ru_histo > count) {
--         radv_fill_buffer(cmd_buffer, NULL, NULL, keyvals_even_addr + count * keyval_bytes,
--                          (count_ru_histo - count) * keyval_bytes, 0xFFFFFFFF);
-+      if (bvh_states[i].count_ru_histo > bvh_states[i].node_count) {
-+         radv_fill_buffer(cmd_buffer, NULL, NULL, keyvals_even_addr + bvh_states[i].node_count * keyval_bytes,
-+                          (bvh_states[i].count_ru_histo - bvh_states[i].node_count) * keyval_bytes, 0xFFFFFFFF);
-       }
- 
-       /*
-@@ -815,28 +826,35 @@ morton_sort(VkCommandBuffer commandBuffer, uint32_t infoCount,
-        * Note that the last workgroup doesn't read/write a partition so it doesn't
-        * need to be initialized.
-        */
--      uint32_t histo_partition_count = passes + scatter_blocks - 1;
--      uint32_t pass_idx = (keyval_bytes - passes);
-+      uint32_t histo_partition_count = passes + bvh_states[i].scatter_blocks - 1;
- 
-       uint32_t fill_base = pass_idx * (RS_RADIX_SIZE * sizeof(uint32_t));
- 
-       radv_fill_buffer(cmd_buffer, NULL, NULL, internal_addr + rs->internal.histograms.offset + fill_base,
-                        histo_partition_count * (RS_RADIX_SIZE * sizeof(uint32_t)), 0);
-+   }
- 
--      /*
--       * Pipeline: HISTOGRAM
--       *
--       * TODO(allanmac): All subgroups should try to process approximately the same
--       * number of blocks in order to minimize tail effects.  This was implemented
--       * and reverted but should be reimplemented and benchmarked later.
--       */
--      vk_barrier_transfer_w_to_compute_r(commandBuffer);
-+   /*
-+    * Pipeline: HISTOGRAM
-+    *
-+    * TODO(allanmac): All subgroups should try to process approximately the same
-+    * number of blocks in order to minimize tail effects.  This was implemented
-+    * and reverted but should be reimplemented and benchmarked later.
-+    */
-+   vk_barrier_transfer_w_to_compute_r(commandBuffer);
-+
-+   radv_CmdBindPipeline(commandBuffer, VK_PIPELINE_BIND_POINT_COMPUTE, rs->pipelines.named.histogram);
-+
-+   for (uint32_t i = 0; i < infoCount; ++i) {
-+      if (!bvh_states[i].node_count)
-+         continue;
- 
--      uint64_t devaddr_histograms = internal_addr + rs->internal.histograms.offset;
-+      uint64_t keyvals_even_addr = pInfos[i].scratchData.deviceAddress + bvh_states[i].scratch.sort_buffer_offset[0];
-+      uint64_t internal_addr = pInfos[i].scratchData.deviceAddress + bvh_states[i].scratch.sort_internal_offset;
- 
-       /* Dispatch histogram */
-       struct rs_push_histogram push_histogram = {
--         .devaddr_histograms = devaddr_histograms,
-+         .devaddr_histograms = internal_addr + rs->internal.histograms.offset,
-          .devaddr_keyvals = keyvals_even_addr,
-          .passes = passes,
-       };
-@@ -844,83 +862,87 @@ morton_sort(VkCommandBuffer commandBuffer, uint32_t infoCount,
-       radv_CmdPushConstants(commandBuffer, rs->pipeline_layouts.named.histogram, VK_SHADER_STAGE_COMPUTE_BIT, 0,
-                             sizeof(push_histogram), &push_histogram);
- 
--      radv_CmdBindPipeline(commandBuffer, VK_PIPELINE_BIND_POINT_COMPUTE, rs->pipelines.named.histogram);
-+      vk_common_CmdDispatch(commandBuffer, bvh_states[i].histo_blocks, 1, 1);
-+   }
- 
--      vk_common_CmdDispatch(commandBuffer, histo_blocks, 1, 1);
-+   /*
-+    * Pipeline: PREFIX
-+    *
-+    * Launch one workgroup per pass.
-+    */
-+   vk_barrier_compute_w_to_compute_r(commandBuffer);
- 
--      /*
--       * Pipeline: PREFIX
--       *
--       * Launch one workgroup per pass.
--       */
--      vk_barrier_compute_w_to_compute_r(commandBuffer);
-+   radv_CmdBindPipeline(commandBuffer, VK_PIPELINE_BIND_POINT_COMPUTE, rs->pipelines.named.prefix);
-+
-+   for (uint32_t i = 0; i < infoCount; ++i) {
-+      if (!bvh_states[i].node_count)
-+         continue;
-+
-+      uint64_t internal_addr = pInfos[i].scratchData.deviceAddress + bvh_states[i].scratch.sort_internal_offset;
- 
-       struct rs_push_prefix push_prefix = {
--         .devaddr_histograms = devaddr_histograms,
-+         .devaddr_histograms = internal_addr + rs->internal.histograms.offset,
-       };
- 
-       radv_CmdPushConstants(commandBuffer, rs->pipeline_layouts.named.prefix, VK_SHADER_STAGE_COMPUTE_BIT, 0,
-                             sizeof(push_prefix), &push_prefix);
- 
--      radv_CmdBindPipeline(commandBuffer, VK_PIPELINE_BIND_POINT_COMPUTE, rs->pipelines.named.prefix);
--
-       vk_common_CmdDispatch(commandBuffer, passes, 1, 1);
-+   }
- 
--      /* Pipeline: SCATTER */
--      vk_barrier_compute_w_to_compute_r(commandBuffer);
-+   /* Pipeline: SCATTER */
-+   vk_barrier_compute_w_to_compute_r(commandBuffer);
- 
--      uint32_t histogram_offset = pass_idx * (RS_RADIX_SIZE * sizeof(uint32_t));
--      uint64_t devaddr_partitions = internal_addr + rs->internal.partitions.offset;
-+   uint32_t histogram_offset = pass_idx * (RS_RADIX_SIZE * sizeof(uint32_t));
- 
--      struct rs_push_scatter push_scatter = {
-+   for (uint32_t i = 0; i < infoCount; i++) {
-+      uint64_t keyvals_even_addr = pInfos[i].scratchData.deviceAddress + bvh_states[i].scratch.sort_buffer_offset[0];
-+      uint64_t keyvals_odd_addr = pInfos[i].scratchData.deviceAddress + bvh_states[i].scratch.sort_buffer_offset[1];
-+      uint64_t internal_addr = pInfos[i].scratchData.deviceAddress + bvh_states[i].scratch.sort_internal_offset;
-+
-+      bvh_states[i].push_scatter = (struct rs_push_scatter){
-          .devaddr_keyvals_even = keyvals_even_addr,
-          .devaddr_keyvals_odd = keyvals_odd_addr,
--         .devaddr_partitions = devaddr_partitions,
--         .devaddr_histograms = devaddr_histograms + histogram_offset,
--         .pass_offset = (pass_idx & 3) * RS_RADIX_LOG2,
-+         .devaddr_partitions = internal_addr + rs->internal.partitions.offset,
-+         .devaddr_histograms = internal_addr + rs->internal.histograms.offset + histogram_offset,
-       };
-+   }
- 
--      {
--         uint32_t pass_dword = pass_idx / 4;
-+   bool is_even = true;
- 
--         radv_CmdPushConstants(commandBuffer, rs->pipeline_layouts.named.scatter[pass_dword].even,
--                               VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(push_scatter), &push_scatter);
-+   while (true) {
-+      uint32_t pass_dword = pass_idx / 4;
- 
--         radv_CmdBindPipeline(commandBuffer, VK_PIPELINE_BIND_POINT_COMPUTE,
--                              rs->pipelines.named.scatter[pass_dword].even);
--      }
-+      /* Bind new pipeline */
-+      VkPipeline p =
-+         is_even ? rs->pipelines.named.scatter[pass_dword].even : rs->pipelines.named.scatter[pass_dword].odd;
-+      radv_CmdBindPipeline(commandBuffer, VK_PIPELINE_BIND_POINT_COMPUTE, p);
- 
--      bool is_even = true;
-+      /* Update push constants that changed */
-+      VkPipelineLayout pl = is_even ? rs->pipeline_layouts.named.scatter[pass_dword].even //
-+                                    : rs->pipeline_layouts.named.scatter[pass_dword].odd;
- 
--      while (true) {
--         vk_common_CmdDispatch(commandBuffer, scatter_blocks, 1, 1);
-+      for (uint32_t i = 0; i < infoCount; i++) {
-+         if (!bvh_states[i].node_count)
-+            continue;
- 
--         /* Continue? */
--         if (++pass_idx >= keyval_bytes)
--            break;
-+         bvh_states[i].push_scatter.pass_offset = (pass_idx & 3) * RS_RADIX_LOG2;
- 
--         vk_barrier_compute_w_to_compute_r(commandBuffer);
-+         radv_CmdPushConstants(commandBuffer, pl, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(struct rs_push_scatter),
-+                               &bvh_states[i].push_scatter);
- 
--         is_even ^= true;
--         push_scatter.devaddr_histograms += (RS_RADIX_SIZE * sizeof(uint32_t));
--         push_scatter.pass_offset = (pass_idx & 3) * RS_RADIX_LOG2;
-+         vk_common_CmdDispatch(commandBuffer, bvh_states[i].scatter_blocks, 1, 1);
- 
--         uint32_t pass_dword = pass_idx / 4;
-+         bvh_states[i].push_scatter.devaddr_histograms += (RS_RADIX_SIZE * sizeof(uint32_t));
-+      }
- 
--         /* Update push constants that changed */
--         VkPipelineLayout pl = is_even ? rs->pipeline_layouts.named.scatter[pass_dword].even
--                                       : rs->pipeline_layouts.named.scatter[pass_dword].odd;
--         radv_CmdPushConstants(commandBuffer, pl, VK_SHADER_STAGE_COMPUTE_BIT,
--                               offsetof(struct rs_push_scatter, devaddr_histograms),
--                               sizeof(push_scatter.devaddr_histograms) + sizeof(push_scatter.pass_offset),
--                               &push_scatter.devaddr_histograms);
-+      /* Continue? */
-+      if (++pass_idx >= keyval_bytes)
-+         break;
- 
--         /* Bind new pipeline */
--         VkPipeline p =
--            is_even ? rs->pipelines.named.scatter[pass_dword].even : rs->pipelines.named.scatter[pass_dword].odd;
-+      vk_barrier_compute_w_to_compute_r(commandBuffer);
- 
--         radv_CmdBindPipeline(commandBuffer, VK_PIPELINE_BIND_POINT_COMPUTE, p);
--      }
-+      is_even ^= true;
-    }
- 
-    cmd_buffer->state.flush_bits |= flush_bits;
--- 
-GitLab

debian/patches/24789.patch

@@ -1,34 +0,0 @@
-From 87f95fa7f24415f51391f128adf7f048358be226 Mon Sep 17 00:00:00 2001
-From: Friedrich Vock <friedrich.vock@gmx.de>
-Date: Sat, 19 Aug 2023 11:00:45 +0200
-Subject: [PATCH] nir/load_store_vectorize: Handle intrinsics with constant
- base
-
-This includes nir_load_stack and nir_store_stack, which are vectorized
-in nir_lower_shader_calls. If not adjusted, we end up loading from
-the wrong base.
-
-Closes: https://gitlab.freedesktop.org/mesa/mesa/-/issues/9596
-Closes: https://gitlab.freedesktop.org/mesa/mesa/-/issues/9587
-Cc: mesa-stable
-Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/24789>
----
- src/compiler/nir/nir_opt_load_store_vectorize.c | 2 ++
- 1 file changed, 2 insertions(+)
-
-diff --git a/src/compiler/nir/nir_opt_load_store_vectorize.c b/src/compiler/nir/nir_opt_load_store_vectorize.c
-index 73e6ff6d8878..4bea8fbea6ff 100644
---- a/src/compiler/nir/nir_opt_load_store_vectorize.c
-+++ b/src/compiler/nir/nir_opt_load_store_vectorize.c
-@@ -756,6 +756,8 @@ vectorize_loads(nir_builder *b, struct vectorize_ctx *ctx,
- 
-       nir_intrinsic_set_range_base(first->intrin, low_base);
-       nir_intrinsic_set_range(first->intrin, MAX2(low_end, high_end) - low_base);
-+   } else if (nir_intrinsic_has_base(first->intrin) && info->base_src == -1 && info->deref_src == -1) {
-+      nir_intrinsic_set_base(first->intrin, nir_intrinsic_base(low->intrin));
-    }
- 
-    first->key = low->key;
--- 
-GitLab
-

debian/patches/revert-af1ee8e01044.diff

@@ -1,153 +0,0 @@
-From 0c3587a2f8e1b6cfadf9a4bbb6ae4b2c3e14a651 Mon Sep 17 00:00:00 2001
-From: Leandro Ribeiro <leandro.ribeiro@collabora.com>
-Date: Sun, 10 Apr 2022 22:54:36 -0300
-Subject: [PATCH] Revert "egl/wayland: deprecate drm_handle_format() and
- drm_handle_capabilities()"
-
-Commit af1ee8e010441f8f2ed8c77065b159652a4ac9fe dropped support to
-wl_drm, as we thought that most compositors from active projects were
-already supporting zwp_linux_dmabuf_v1.
-
-But that's not true, so revert this commit in order to give these
-projects a longer transition period.
-
-Note that we didn't add back the support to GEM name API, and that was
-on purpose.
-
-Signed-off-by: Leandro Ribeiro <leandro.ribeiro@collabora.com>
----
- src/egl/drivers/dri2/egl_dri2.h         |  1 +
- src/egl/drivers/dri2/platform_wayland.c | 59 +++++++++++++++++++------
- 2 files changed, 47 insertions(+), 13 deletions(-)
-
-diff --git a/src/egl/drivers/dri2/egl_dri2.h b/src/egl/drivers/dri2/egl_dri2.h
-index 89158993efdd..1c840a966b3c 100644
---- a/src/egl/drivers/dri2/egl_dri2.h
-+++ b/src/egl/drivers/dri2/egl_dri2.h
-@@ -284,6 +284,7 @@ struct dri2_egl_display
-    struct zwp_linux_dmabuf_feedback_v1 *wl_dmabuf_feedback;
-    struct dmabuf_feedback_format_table format_table;
-    bool authenticated;
-+   uint32_t capabilities;
-    char *device_name;
- #endif
- 
-diff --git a/src/egl/drivers/dri2/platform_wayland.c b/src/egl/drivers/dri2/platform_wayland.c
-index e9ecf6d1e716..19fad8bfa08e 100644
---- a/src/egl/drivers/dri2/platform_wayland.c
-+++ b/src/egl/drivers/dri2/platform_wayland.c
-@@ -1344,7 +1344,7 @@ create_wl_buffer(struct dri2_egl_display *dri2_dpy,
-                  struct dri2_egl_surface *dri2_surf,
-                  __DRIimage *image)
- {
--   struct wl_buffer *ret;
-+   struct wl_buffer *ret = NULL;
-    EGLBoolean query;
-    int width, height, fourcc, num_planes;
-    uint64_t modifier = DRM_FORMAT_MOD_INVALID;
-@@ -1448,11 +1448,28 @@ create_wl_buffer(struct dri2_egl_display *dri2_dpy,
-       ret = zwp_linux_buffer_params_v1_create_immed(params, width, height,
-                                                     fourcc, 0);
-       zwp_linux_buffer_params_v1_destroy(params);
-+   } else {
-+      struct wl_drm *wl_drm =
-+         dri2_surf ? dri2_surf->wl_drm_wrapper : dri2_dpy->wl_drm;
-+      int fd, stride;
-+
-+      if (num_planes > 1)
-+         return NULL;
-+
-+      query = dri2_dpy->image->queryImage(image, __DRI_IMAGE_ATTRIB_FD, &fd);
-+      query &= dri2_dpy->image->queryImage(image, __DRI_IMAGE_ATTRIB_STRIDE, &stride);
-+      if (!query) {
-+         if (fd >= 0)
-+            close(fd);
-+         return NULL;
-+      }
- 
--      return ret;
-+      ret = wl_drm_create_prime_buffer(wl_drm, fd, width, height, fourcc, 0,
-+                                       stride, 0, 0, 0, 0);
-+      close(fd);
-    }
- 
--   return NULL;
-+   return ret;
- }
- 
- static EGLBoolean
-@@ -1699,16 +1716,21 @@ drm_handle_device(void *data, struct wl_drm *drm, const char *device)
- static void
- drm_handle_format(void *data, struct wl_drm *drm, uint32_t format)
- {
--   /* deprecated, as compositors already support the dma-buf protocol extension
--    * and so we can rely on dmabuf_handle_modifier() to receive formats and
--    * modifiers */
-+   struct dri2_egl_display *dri2_dpy = data;
-+   int visual_idx = dri2_wl_visual_idx_from_fourcc(format);
-+
-+   if (visual_idx == -1)
-+      return;
-+
-+   BITSET_SET(dri2_dpy->formats.formats_bitmap, visual_idx);
- }
- 
- static void
- drm_handle_capabilities(void *data, struct wl_drm *drm, uint32_t value)
- {
--   /* deprecated, as compositors already support the dma-buf protocol extension
--    * and so we can rely on it to create wl_buffer's */
-+   struct dri2_egl_display *dri2_dpy = data;
-+
-+   dri2_dpy->capabilities = value;
- }
- 
- static void
-@@ -2077,13 +2099,12 @@ dri2_initialize_wayland_drm(_EGLDisplay *disp)
-    wl_registry_add_listener(dri2_dpy->wl_registry,
-                             &registry_listener_drm, dri2_dpy);
- 
--   /* The compositor must expose the dma-buf interface. */
--   if (roundtrip(dri2_dpy) < 0 || dri2_dpy->wl_dmabuf == NULL)
-+   if (roundtrip(dri2_dpy) < 0)
-       goto cleanup;
- 
-    /* Get default dma-buf feedback */
--   if (zwp_linux_dmabuf_v1_get_version(dri2_dpy->wl_dmabuf) >=
--       ZWP_LINUX_DMABUF_V1_GET_DEFAULT_FEEDBACK_SINCE_VERSION) {
-+   if (dri2_dpy->wl_dmabuf && zwp_linux_dmabuf_v1_get_version(dri2_dpy->wl_dmabuf) >=
-+                              ZWP_LINUX_DMABUF_V1_GET_DEFAULT_FEEDBACK_SINCE_VERSION) {
-       dmabuf_feedback_format_table_init(&dri2_dpy->format_table);
-       dri2_dpy->wl_dmabuf_feedback =
-          zwp_linux_dmabuf_v1_get_default_feedback(dri2_dpy->wl_dmabuf);
-@@ -2091,7 +2112,6 @@ dri2_initialize_wayland_drm(_EGLDisplay *disp)
-                                                 &dmabuf_feedback_listener, dri2_dpy);
-    }
- 
--   /* Receive events from the interfaces */
-    if (roundtrip(dri2_dpy) < 0)
-       goto cleanup;
- 
-@@ -2178,6 +2198,19 @@ dri2_initialize_wayland_drm(_EGLDisplay *disp)
- 
-    dri2_wl_setup_swap_interval(disp);
- 
-+   if (dri2_dpy->wl_drm) {
-+      /* To use Prime, we must have _DRI_IMAGE v7 at least. createImageFromFds
-+       * support indicates that Prime export/import is supported by the driver.
-+       * We deprecated the support to GEM names API, so we bail out if the
-+       * driver does not suport Prime. */
-+      if (!(dri2_dpy->capabilities & WL_DRM_CAPABILITY_PRIME) ||
-+          (dri2_dpy->image->base.version < 7) ||
-+          (dri2_dpy->image->createImageFromFds == NULL)) {
-+         _eglLog(_EGL_WARNING, "wayland-egl: display does not support prime");
-+         goto cleanup;
-+      }
-+   }
-+
-    if (dri2_dpy->is_different_gpu &&
-        (dri2_dpy->image->base.version < 9 ||
-         dri2_dpy->image->blitImage == NULL)) {
--- 
-GitLab
-

debian/patches/series

@@ -1,6 +0,0 @@
-#07_gallium-fix-build-failure-on-powerpcspe.diff
-#path_max.diff
-#src_glx_dri_common.h.diff
-#21929.patch
-#24720.patch
-#24789.patch

debian/rules

@@ -13,138 +13,121 @@ DEB_HOST_ARCH_OS   ?= $(shell dpkg-architecture -qDEB_HOST_ARCH_OS)
 DEB_HOST_ARCH_CPU  ?= $(shell dpkg-architecture -qDEB_HOST_ARCH_CPU)
 
 # for finding the correct llvm-config when meson doesn't know about it yet
-LLVM_VERSION = 15
+LLVM_VERSION = 18
+RUST_VERSION = 1.80
 export PATH:=/usr/lib/llvm-$(LLVM_VERSION)/bin/:$(PATH)
+export PATH:=/usr/lib/rust-$(RUST_VERSION)/bin/:$(PATH)
 
-export DEB_BUILD_MAINT_OPTIONS=optimize=-lto
-# enable LTO everywhere:
-#confflags += -Db_lto=true
-
-ifeq (,$(filter $(DEB_HOST_ARCH), armhf ppc64el sh3 sh4))
-buildflags = \
-	$(shell DEB_CFLAGS_MAINT_APPEND=-Wall DEB_CXXFLAGS_MAINT_APPEND=-Wall dpkg-buildflags --export=configure)
+ifneq (,$(filter $(DEB_HOST_ARCH), amd64))
+	export DEB_BUILD_MAINT_OPTIONS=optimize=+lto
 else
-  ifneq (,$(filter $(DEB_HOST_ARCH), armhf))
-  # Workaround for a variant of LP: #725126
-  buildflags = \
-	$(shell DEB_CFLAGS_MAINT_APPEND="-Wall -fno-optimize-sibling-calls" DEB_CXXFLAGS_MAINT_APPEND="-Wall -fno-optimize-sibling-calls" dpkg-buildflags --export=configure)
-  else
-    ifneq (,$(filter $(DEB_HOST_ARCH), ppc64el))
-    # Workaround for https://gitlab.freedesktop.org/mesa/mesa/-/issues/5315
-    buildflags = \
-	$(shell DEB_CFLAGS_MAINT_APPEND="-Wall -O2" DEB_CXXFLAGS_MAINT_APPEND="-Wall -O2" dpkg-buildflags --export=configure)
-    else
-    # Workaround for https://gcc.gnu.org/bugzilla/show_bug.cgi?id=83143
-    buildflags = \
-	$(shell DEB_CFLAGS_MAINT_APPEND="-Wall -O1" DEB_CXXFLAGS_MAINT_APPEND="-Wall -O1" dpkg-buildflags --export=configure)
-    endif
-  endif
+	export DEB_BUILD_MAINT_OPTIONS=optimize=-lto
 endif
+export MESON_PACKAGE_CACHE_DIR=/usr/share/cargo/registry/
+
+DEB_CFLAGS_MAINT_APPEND := -Wall
+DEB_CXXFLAGS_MAINT_APPEND := -Wall
+ifneq (,$(filter $(DEB_HOST_ARCH), armhf))
+  # Workaround for a variant of LP: #725126
+  DEB_CFLAGS_MAINT_APPEND += -fno-optimize-sibling-calls
+  DEB_CXXFLAGS_MAINT_APPEND += -fno-optimize-sibling-calls
+else ifneq (,$(filter $(DEB_HOST_ARCH), m68k))
+  # This library has huge jump tables: Debian #1067207
+  DEB_CFLAGS_MAINT_APPEND += -mlong-jump-table-offsets
+  DEB_CXXFLAGS_MAINT_APPEND += -mlong-jump-table-offsets
+else ifneq (,$(filter $(DEB_HOST_ARCH), sh4))
+  # Workaround for https://gcc.gnu.org/bugzilla/show_bug.cgi?id=83143
+  DEB_CFLAGS_MAINT_APPEND += -freorder-blocks-algorithm=simple
+  DEB_CXXFLAGS_MAINT_APPEND += -freorder-blocks-algorithm=simple
+endif
+buildflags = $(shell \
+	DEB_CFLAGS_MAINT_APPEND='$(DEB_CFLAGS_MAINT_APPEND)' \
+	DEB_CXXFLAGS_MAINT_APPEND='$(DEB_CXXFLAGS_MAINT_APPEND)' \
+	dpkg-buildflags --export=configure)
 
 EGL_PLATFORMS = x11
-GALLIUM_DRIVERS = swrast
+GALLIUM_DRIVERS =
 VULKAN_DRIVERS =
 VULKAN_LAYERS =
 
-confflags_DRI3 = -Ddri3=disabled
-confflags_OSMESA =  -Dosmesa=true
+confflags_OSMESA = -Dosmesa=true
 confflags_SSE2 = -Dsse2=true
+confflags_TEFLON = -Dteflon=false
 
-LLVM_ARCHS     = amd64 arm64 armel armhf i386 mips64el mipsel powerpc ppc64 ppc64el riscv64 s390x sparc64 x32
-RUSTICL_ARCHS  = amd64 arm64 armel armhf mips64el mipsel ppc64el s390x
-#ifeq ($(DEB_DISTRIBUTION), jammy)
-#       RUSTICL_ARCHS  = arm64
-#else
-VALGRIND_ARCHS = amd64 arm64 armhf i386 mips64el mipsel powerpc ppc64 ppc64el s390x
-WINE_ARCHS     = amd64 arm64 armel armhf i386 powerpc
+LLVM_ARCHS     = amd64 arm64 armel armhf i386 loong64 mips64el powerpc ppc64 ppc64el riscv64 s390x sparc64 x32
+RUSTICL_ARCHS  = amd64 arm64 armel armhf loong64 mips64el powerpc ppc64 ppc64el riscv64 s390x x32
+NVK_ARCHS      = amd64 arm64 armhf ppc64 riscv64 x32
+VALGRIND_ARCHS = amd64 arm64 armhf i386 mips64el powerpc ppc64 ppc64el s390x
+WINE_ARCHS     = amd64 arm64 armel armhf i386
+WSL_ARCHS      = amd64 arm64
 
 # hurd doesn't do direct rendering
 ifeq ($(DEB_HOST_ARCH_OS), hurd)
-	confflags_DIRECT_RENDERING = -Dglx-direct=false
-	confflags_GBM = -Dgbm=disabled
+  GALLIUM_DRIVERS += softpipe
+  confflags_DIRECT_RENDERING = -Dglx-direct=false
+  confflags_GBM = -Dgbm=disabled
 else
-	GALLIUM_DRIVERS += r300 r600
+  # Non-Linux ports lack epoll, so wayland isn't ready yet:
+  # https://gitlab.freedesktop.org/wayland/wayland/-/issues/72
+  # hurd also lacks *_CLOEXEC
+  EGL_PLATFORMS += wayland
 
-	confflags_DIRECT_RENDERING = -Dglx-direct=true
-	confflags_GBM = -Dgbm=enabled
-	confflags_GALLIUM += -Dgallium-extra-hud=true
-	confflags_GALLIUM += -Dgallium-vdpau=enabled
-	confflags_GALLIUM += -Dlmsensors=enabled
+  GALLIUM_DRIVERS += nouveau r300 r600 virgl
+  confflags_DIRECT_RENDERING = -Dglx-direct=true
+  confflags_GBM = -Dgbm=enabled
+  confflags_GALLIUM += -Dgallium-extra-hud=true
+  confflags_GALLIUM += -Dgallium-vdpau=enabled
+  confflags_GALLIUM += -Dlmsensors=enabled
 
-  # radv/lavapipe needs LLVM and the Vulkan loader, so only build on the subset of
-  # arches where we have LLVM enabled and where the Vulkan loader is built.
-  ifneq (,$(filter $(DEB_HOST_ARCH), amd64 arm64 armel armhf i386 mips64el mipsel powerpc ppc64 ppc64el s390x sparc64))
-	VULKAN_DRIVERS += amd swrast virtio nouveau-experimental
+  # Freedreno requires arm in addition
+  ifneq (,$(filter arm arm64,$(DEB_HOST_ARCH_CPU)))
+	GALLIUM_DRIVERS += freedreno
+	# XXX - broken
+	#GALLIUM_DRIVERS += asahi
   endif
 
-  # Only enable amd on riscv64, swrast needs CPU JIT support which doesn't work properly yet
-  ifneq (,$(filter $(DEB_HOST_ARCH), riscv64))
-	VULKAN_DRIVERS += amd
-	confflags_GALLIUM += -Ddraw-use-llvm=false
+  # etnaviv, tegra, vc4 and v3d kernel support are only available on armhf and arm64
+  ifneq (,$(filter $(DEB_HOST_ARCH), armhf arm64))
+	GALLIUM_DRIVERS += etnaviv panfrost svga tegra vc4 v3d
+	VULKAN_DRIVERS += broadcom freedreno panfrost
   endif
 
-  ifeq ($(DEB_HOST_ARCH_OS), linux)
-	confflags_DRI3 = -Ddri3=enabled
-	# Gallium drivers which require kernel support, not yet ported to non-Linux
-	GALLIUM_DRIVERS += nouveau virgl
-
-	# Freedreno requires arm in addition
-	ifneq (,$(filter arm arm64,$(DEB_HOST_ARCH_CPU)))
-		GALLIUM_DRIVERS += freedreno asahi
-	endif
-
-	# etnaviv, tegra, vc4 and v3d kernel support are only available on armhf and arm64
-	ifneq (,$(filter $(DEB_HOST_ARCH), armhf arm64))
-		GALLIUM_DRIVERS += etnaviv panfrost svga tegra vc4 v3d
-		VULKAN_DRIVERS += broadcom freedreno panfrost
-	endif
-
-	ifneq (,$(filter $(DEB_HOST_ARCH), armhf arm64 riscv64))
-		GALLIUM_DRIVERS += lima
-	endif
-
-	ifneq (,$(filter $(DEB_HOST_ARCH), amd64 i386 x32))
-		GALLIUM_DRIVERS += svga
-		# svga needs xa state tracker
-		confflags_GALLIUM += -Dgallium-xa=enabled
-		VULKAN_DRIVERS += intel intel_hasvk
-	endif
-	ifneq (,$(filter $(DEB_HOST_ARCH), amd64))
-		confflags_GALLIUM += -Dintel-clc=enabled
-	endif
-
-	# Non-Linux ports lack epoll, so wayland isn't ready yet:
-	# https://gitlab.freedesktop.org/wayland/wayland/-/issues/72
-	# hurd also lacks *_CLOEXEC
-	EGL_PLATFORMS += ,wayland
-
-	ifneq (,$(filter $(DEB_HOST_ARCH), $(VALGRIND_ARCHS)))
-		confflags_VALGRIND += -Dvalgrind=enabled
-	endif
-
-	# WSL supports only amd64 and arm64
-	ifneq (,$(filter amd64 arm64,$(DEB_HOST_ARCH)))
-		GALLIUM_DRIVERS += d3d12
-#		VULKAN_DRIVERS += microsoft-experimental
-	endif
+  ifneq (,$(filter $(DEB_HOST_ARCH), armhf arm64 loong64 riscv64))
+	GALLIUM_DRIVERS += lima
   endif
 
   # Build intel drivers on archs where libdrm-intel is installed
-  ifneq (,$(filter $(DEB_HOST_ARCH),amd64 i386 kfreebsd-amd64 kfreebsd-i386 x32))
-	GALLIUM_DRIVERS += iris crocus i915
+  ifneq (,$(filter $(DEB_HOST_ARCH), amd64 i386 x32))
+	GALLIUM_DRIVERS += crocus i915 iris svga
+	VULKAN_DRIVERS += intel intel_hasvk
+	# svga needs xa state tracker
+	confflags_GALLIUM += -Dgallium-xa=enabled
+  endif
+
+  ifneq (,$(filter $(DEB_HOST_ARCH), amd64))
+	confflags_GALLIUM += -Dintel-clc=enabled
+  endif
+
+  ifneq (,$(filter $(DEB_HOST_ARCH), $(VALGRIND_ARCHS)))
+	confflags_VALGRIND += -Dvalgrind=enabled
+  endif
+
+  # WSL supports only amd64 and arm64
+  ifneq (,$(filter $(DEB_HOST_ARCH), $(WSL_ARCHS)))
+	GALLIUM_DRIVERS += d3d12
+	VULKAN_DRIVERS += microsoft-experimental
   endif
 
   # LLVM is required for building r300g, radeonsi and llvmpipe drivers.
   # It's also required for building OpenCL support.
   ifneq (,$(filter $(DEB_HOST_ARCH), $(LLVM_ARCHS)))
-	GALLIUM_DRIVERS += radeonsi zink
+	GALLIUM_DRIVERS += radeonsi zink llvmpipe
 
-	# Only enable amd on riscv64, swrast needs CPU JIT support which doesn't work properly yet
-	ifneq (,$(filter $(DEB_HOST_ARCH), riscv64))
-		VULKAN_DRIVERS += amd
-		confflags_GALLIUM += -Ddraw-use-llvm=false
-	else
+	# drop virtio from armel, it doesn't build
+	ifneq (,$(filter $(DEB_HOST_ARCH), armel))
 		VULKAN_DRIVERS += amd swrast
+	else
+		VULKAN_DRIVERS += amd swrast virtio
 	endif
 
 	VULKAN_LAYERS += device-select intel-nullhw overlay
@@ -153,7 +136,15 @@ else
 
 	# Build rusticl for archs where rustc is available
 	ifneq (,$(filter $(DEB_HOST_ARCH), $(RUSTICL_ARCHS)))
-		confflags_GALLIUM += -Dgallium-rusticl=true
+		ifneq ($(DEB_DISTRIBUTION), xxx)
+			confflags_GALLIUM += -Dgallium-rusticl=true
+		else
+			confflags_GALLIUM += -Dgallium-rusticl=false
+		endif
+	endif
+
+	ifneq (,$(filter $(DEB_HOST_ARCH), $(NVK_ARCHS)))
+		VULKAN_DRIVERS += nouveau
 	endif
 
 	# nine makes sense only on archs that build wine
@@ -161,22 +152,30 @@ else
 		confflags_GALLIUM += -Dgallium-nine=true
 	endif
   else
+	GALLIUM_DRIVERS += softpipe
 	confflags_GALLIUM += -Dllvm=disabled
   endif
 
   ifeq (,$(filter pkg.mesa.nolibva,$(DEB_BUILD_PROFILES)))
 	confflags_GALLIUM += -Dgallium-va=enabled
-	confflags_GALLIUM += -Dvideo-codecs="vc1dec, h264dec, h264enc, h265dec, h265enc"
+	confflags_GALLIUM += -Dvideo-codecs="all"
+  endif
+
+  # Teflon only supports arm64
+  ifneq (,$(filter $(DEB_HOST_ARCH), arm64))
+	confflags_TEFLON = -Dteflon=true
   endif
 endif
 
-ifeq ($(DEB_HOST_ARCH), i386)
-	confflags_SSE2 = -Dsse2=false
-endif
+#ifeq ($(DEB_HOST_ARCH), i386)
+#	confflags_SSE2 = -Dsse2=false
+#endif
 
 empty:=
 space := $(empty) $(empty)
 comma := ,
+EGL_PLATFORMS := $(patsubst %,'%',$(EGL_PLATFORMS))
+EGL_PLATFORMS_LIST := $(subst $(space),$(comma),$(EGL_PLATFORMS))
 GALLIUM_DRIVERS := $(patsubst %,'%',$(GALLIUM_DRIVERS))
 GALLIUM_DRIVERS_LIST := $(subst $(space),$(comma),$(GALLIUM_DRIVERS))
 VULKAN_DRIVERS := $(patsubst %,'%',$(VULKAN_DRIVERS))
@@ -184,44 +183,58 @@ VULKAN_DRIVERS_LIST := $(subst $(space),$(comma),$(VULKAN_DRIVERS))
 VULKAN_LAYERS := $(patsubst %,'%',$(VULKAN_LAYERS))
 VULKAN_LAYERS_LIST := $(subst $(space),$(comma),$(VULKAN_LAYERS))
 
-confflags_EGL = -Dplatforms="$(EGL_PLATFORMS)"
 confflags_GLES = -Dgles1=disabled -Dgles2=enabled
 confflags_GALLIUM += -Dgallium-drivers="[$(GALLIUM_DRIVERS_LIST)]"
 
 confflags += \
 	-Ddri-drivers-path=/usr/lib/$(DEB_HOST_MULTIARCH)/dri \
 	-Ddri-search-path='/usr/lib/$(DEB_HOST_MULTIARCH)/dri:\$$$${ORIGIN}/dri:/usr/lib/dri' \
+	-Dplatforms="[$(EGL_PLATFORMS_LIST)]" \
 	-Dvulkan-drivers="[$(VULKAN_DRIVERS_LIST)]" \
 	-Dvulkan-layers="[$(VULKAN_LAYERS_LIST)]" \
-	-Dglvnd=true \
+	-Dglvnd=enabled \
 	-Dshared-glapi=enabled \
-	-Dgallium-omx=disabled \
 	-Db_ndebug=true \
 	-Dbuild-tests=true \
 	-Dtools=drm-shim \
 	$(confflags_DIRECT_RENDERING) \
 	$(confflags_GBM) \
-	$(confflags_DRI3) \
-	$(confflags_EGL) \
 	$(confflags_GALLIUM) \
 	$(confflags_GLES) \
 	$(confflags_OSMESA) \
 	$(confflags_SSE2) \
-	$(confflags_VALGRIND)
+	$(confflags_VALGRIND) \
+	$(confflags_TEFLON)
+
+
+rewrite_wrap_files:
+	cp -r subprojects subprojects-save
+	for crate in paste proc-macro2 quote syn unicode-ident; \
+	do \
+		export crate_namever=`basename $$MESON_PACKAGE_CACHE_DIR/$$crate-*`; \
+		sed -e"/source.*/d" -e"s,$${crate}-.*,$${crate_namever}," -i subprojects/$${crate}.wrap; \
+	done
+	touch subprojects/rewrite
 
 override_dh_clean: regen_control
-	sed -i 's/-.*//' VERSION
+	dpkg-parsechangelog | awk '/^Version:/ {print $$2}' | sed 's/-.*//;s/~/-/' > VERSION
 	rm -rf .pc
 	rm -rf build
 	rm -rf configure bin/config.guess bin/config.sub config.h.in
 	rm -rf $$(find -name Makefile.in)
 	rm -rf bin/install-sh bin/ltmain.sh
+	rm -rf src/intel/vulkan/grl/parser.out
+	rm -rf src/intel/vulkan/grl/parsetab.py
 	for file in debian/*.links.in; do rm -f $${file%%.in}; done
 	find -name '*.pyc' -delete
 	find -name '__pycache__' -delete
+	if [ -f subprojects/rewrite ]; then \
+		rm -rf subprojects; \
+		mv subprojects-save subprojects; \
+	fi
 	dh_clean
 
-override_dh_auto_configure:
+override_dh_auto_configure: rewrite_wrap_files
 	dpkg-parsechangelog | awk '/^Version:/ {print $$2}' > VERSION
 	$(buildflags) dh_auto_configure -- \
 		$(confflags)
@@ -248,13 +261,16 @@ override_dh_installchangelogs:
 override_dh_install:
 	# purge .la files
 	find debian/tmp/ -name '*.la' -exec rm '{}' ';'
-	find debian/tmp/usr/bin -name 'mme_*_sim_hw_test' -exec rm '{}' ';'
 
 	# Get rid of some files which aren't installed. Do not
 	# use -f to ensure we notice disappearing files:
 	rm debian/tmp/usr/lib/*/libglapi.so
 	rm debian/tmp/usr/lib/*/libEGL_mesa.so
 	rm debian/tmp/usr/lib/*/libGLX_mesa.so
+  ifneq (,$(filter $(DEB_HOST_ARCH), $(NVK_ARCHS)))
+	rm debian/tmp/usr/bin/mme_fermi_sim_hw_test
+	rm debian/tmp/usr/bin/mme_tu104_sim_hw_test
+  endif
 	# use -f here though
 	rm -f debian/tmp/usr/lib/*/libgrl.a
 
@@ -277,14 +293,24 @@ override_dh_install:
     endif
   endif
 
-  ifneq (,$(filter $(DEB_HOST_ARCH), $(RUSTICL_ARCHS)))
+ifneq (,$(filter $(DEB_HOST_ARCH), $(RUSTICL_ARCHS)))
 	install -m755 -d debian/mesa-opencl-icd-git/etc/OpenCL/vendors
 	mv debian/tmp/etc/OpenCL/vendors/rusticl.icd \
 	   debian/mesa-opencl-icd-git/etc/OpenCL/vendors
 	install -m755 -d debian/mesa-opencl-icd-git/usr/lib/${DEB_HOST_MULTIARCH}/
 	mv debian/tmp/usr/lib/${DEB_HOST_MULTIARCH}/libRusticlOpenCL* \
 	   debian/mesa-opencl-icd-git/usr/lib/${DEB_HOST_MULTIARCH}/
-  endif
+endif
+
+# Copy spirv2dxil and its dependencies on architectures that support WSL/Dozen (dzn)
+ifneq (,$(filter $(DEB_HOST_ARCH), $(WSL_ARCHS)))
+	install -m755 -d debian/mesa-vulkan-drivers-git/usr/bin
+	mv debian/tmp/usr/bin/spirv2dxil \
+		debian/mesa-vulkan-drivers-git/usr/bin
+	install -m755 -d debian/mesa-vulkan-drivers-git/usr/lib/${DEB_HOST_MULTIARCH}/
+	mv debian/tmp/usr/lib/${DEB_HOST_MULTIARCH}/libspirv_to_dxil.* \
+		debian/mesa-vulkan-drivers-git/usr/lib/${DEB_HOST_MULTIARCH}/
+endif
 
 	dh_install -a
 
@@ -295,9 +321,6 @@ override_dh_install:
 		install -m 755 debian/local/script  debian/$$p/usr/share/bug/$$p; \
 	done
 
-override_dh_missing:
-	dh_missing --fail-missing
-
 override_dh_makeshlibs:
 	dh_makeshlibs -a -- -c0
 
@@ -313,8 +336,10 @@ gentarball:
 
 regen_control:
 	sed -e"s,@LLVM_VERSION@,$(LLVM_VERSION),g" \
+	    -e"s,@RUST_VERSION@,$(RUST_VERSION),g" \
 	    -e"s,@LLVM_ARCHS@,$(LLVM_ARCHS),g" \
 	    -e"s,@RUSTICL_ARCHS@,$(RUSTICL_ARCHS),g" \
+	    -e"s,@NVK_ARCHS@,$(NVK_ARCHS),g" \
 	    -e"s,@VALGRIND_ARCHS@,$(VALGRIND_ARCHS),g" \
 	    -e"s,@WINE_ARCHS@,$(WINE_ARCHS),g" \
 	    debian/control.in > debian/control

main.sh

@@ -1,20 +1,36 @@
 #! /bin/bash
+
 set -e
 
+source ./pika-build-config.sh
+
+echo "$PIKA_BUILD_ARCH" > pika-build-arch
+
+VERSION="24.4"
+
+# TEMP I386 SPECFIC FIX
+if [[ "$PIKA_BUILD_ARCH" == "i386" ]]
+then
+    wget http://ftp.us.debian.org/debian/pool/main/s/spirv-llvm-translator-18/libllvmspirvlib18.1_18.1.4-1_i386.deb -O ./32bit-spirv.deb
+    wget http://ftp.us.debian.org/debian/pool/main/s/spirv-llvm-translator-18/libllvmspirvlib-18-dev_18.1.4-1_i386.deb -O ./32bit-spirv-dev.deb
+    apt install -y ./32bit-spirv-dev.deb ./32bit-spirv.deb
+fi
+#
+
 # Clone Upstream
-git clone https://gitlab.freedesktop.org/mesa/mesa ./mesa-git
+git clone https://gitlab.freedesktop.org/mesa/mesa mesa-git
 cp -rvf ./debian ./mesa-git/
 cd ./mesa-git
-git submodule update --init
+for i in $(cat ../patches/series) ; do echo "Applying Patch: $i" && patch -Np1 -i ../patches/$i || bash -c "echo "Applying Patch $i Failed!" && exit 2"; done
 sed -i ' 1 s/.*/& - PikaOS YellowBirb Mesa Git /' ./VERSION
 touch debian/changelog
-echo -e "mesa-git (23.3-99pika"$(date '+%Y%m%d')".git.1."$(git rev-parse --short HEAD)") lunar; urgency=medium\n\n  * New GIT Release\n\n -- Ward Nakchbandi <hotrod.master@hotmail.com> Sat, 01 Oct 2022 14:50:00 +0200" > debian/changelog
+echo -e "mesa-git ($VERSION."$(date '+%Y%m%d')".git."$(git rev-parse --short HEAD)") pika; urgency=medium\n\n  * New GIT Release\n\n -- Ward Nakchbandi <hotrod.master@hotmail.com> Sat, 01 Oct 2022 14:50:00 +0200" > debian/changelog
 
 # Get build deps
 apt-get build-dep ./ -y
 
 # Build package
-
+#LOGNAME=root dh_make --createorig -y -l -p mesa-git_"$VERSION" || echo "dh-make: Ignoring Last Error"
 dpkg-buildpackage --no-sign
 
 # Move the debs to output

main32.sh

@@ -1,23 +0,0 @@
-#! /bin/bash
-set -e
-
-# Clone Upstream
-git clone https://gitlab.freedesktop.org/mesa/mesa ./mesa-git
-cp -rvf ./debian ./mesa-git/
-cd ./mesa-git
-git submodule update --init
-sed -i ' 1 s/.*/& - PikaOS YellowBirb Mesa Git /' ./VERSION
-touch debian/changelog
-echo -e "mesa-git (23.3-99pika"$(date '+%Y%m%d')".git.1."$(git rev-parse --short HEAD)") lunar; urgency=medium\n\n  * New GIT Release\n\n -- Ward Nakchbandi <hotrod.master@hotmail.com> Sat, 01 Oct 2022 14:50:00 +0200" > debian/changelog
-
-# Get build deps
-apt-get build-dep ./ -y
-
-# Build package
-
-dpkg-buildpackage --no-sign
-
-# Move the debs to output
-cd ../
-mkdir -p ./output
-mv ./*.deb ./output/

patches/series

@@ -0,0 +1,2 @@
+#path_max.diff
+#src_glx_dri_common.h.diff

pika-build-config/amd64-v3.sh

@@ -0,0 +1,10 @@
+#! /bin/bash
+export PIKA_BUILD_ARCH="amd64-v3"
+export DEBIAN_FRONTEND="noninteractive"
+export DEB_BUILD_MAINT_OPTIONS="optimize=+lto -march=x86-64-v3 -O3 -flto -fuse-linker-plugin -falign-functions=32"
+export DEB_CFLAGS_MAINT_APPEND="-march=x86-64-v3 -O3 -flto -fuse-linker-plugin -falign-functions=32"
+export DEB_CPPFLAGS_MAINT_APPEND="-march=x86-64-v3 -O3 -flto -fuse-linker-plugin -falign-functions=32"
+export DEB_CXXFLAGS_MAINT_APPEND="-march=x86-64-v3 -O3 -flto -fuse-linker-plugin -falign-functions=32"
+export DEB_LDFLAGS_MAINT_APPEND="-march=x86-64-v3 -O3 -flto -fuse-linker-plugin -falign-functions=32"
+export DEB_BUILD_OPTIONS="nocheck notest terse"
+export DPKG_GENSYMBOLS_CHECK_LEVEL=0

pika-build-config/i386.sh

@@ -0,0 +1,5 @@
+#! /bin/bash
+export PIKA_BUILD_ARCH="i386"
+export DEBIAN_FRONTEND="noninteractive"
+export DEB_BUILD_OPTIONS="nocheck notest terse"
+export DPKG_GENSYMBOLS_CHECK_LEVEL=0

release.sh

@@ -1,8 +1,2 @@
 # send debs to server
-rsync -azP --include './' --include '*.deb' --exclude '*' ./output/ ferreo@direct.pika-os.com:/srv/www/incoming/
-
-# add debs to repo
-ssh ferreo@direct.pika-os.com 'aptly repo add -force-replace -remove-files pika-main /srv/www/incoming/'
-
-# publish the repo
-ssh ferreo@direct.pika-os.com 'aptly publish update -batch -skip-contents -force-overwrite lunar filesystem:pikarepo:'
+rsync -azP --include './' --include '*.deb' --exclude '*' ./output/ ferreo@direct.pika-os.com:/srv/www/cockatiel-incoming/