Compare commits

...

59 Commits
lunar ... main

Author SHA1 Message Date
Ward from fusion-voyager-3
2ad132ee7e fix i386 maybe attempt #1
All checks were successful
PikaOS Package Build & Release (i386) / build (push) Successful in 6m12s
2024-09-17 19:36:06 +03:00
Ward from fusion-voyager-3
ff4e963f21 Try fixing mesa-git again
Some checks failed
PikaOS Package Build & Release (amd64-v3 on debian) / build (push) Successful in 12m44s
PikaOS Package Build & Release (i386) / build (push) Failing after 17s
2024-09-17 19:21:31 +03:00
c2188c0003 Update .github/release-nest-i386
Some checks failed
PikaOS Package Build & Release (i386) / build (push) Failing after 2m9s
2024-09-16 06:30:29 +02:00
3d21425253 Update .github/release-debian-v3
All checks were successful
PikaOS Package Build & Release (amd64-v3 on debian) / build (push) Successful in 11m56s
2024-09-16 06:30:17 +02:00
cd94c5a12d Update debian/rules 2024-09-16 06:30:05 +02:00
ce334a02e3 Update .github/release-nest-i386
Some checks failed
PikaOS Package Build & Release (i386) / build (push) Failing after 2m8s
2024-09-15 21:30:10 +02:00
ce56e8e6d8 Update .github/release-nest-v3
Some checks failed
PikaOS Package Build & Release (amd64-v3) / build (push) Has been cancelled
2024-09-15 21:29:35 +02:00
efdb2cec08 Update debian/rules 2024-09-15 21:29:13 +02:00
Ward from fusion-voyager-3
ec3420e422 fu
Some checks failed
PikaOS Package Build & Release (amd64-v3) / build (push) Failing after 3m23s
2024-09-15 22:12:43 +03:00
Ward from fusion-voyager-3
e4907b303c
Some checks failed
PikaOS Package Build & Release (i386) / build (push) Failing after 2m37s
2024-09-15 22:03:51 +03:00
Ward from fusion-voyager-3
3f92c21e93
Some checks failed
PikaOS Package Build & Release (i386) / build (push) Failing after 1m36s
2024-09-15 21:50:48 +03:00
Ward from fusion-voyager-3
798a01fe84 f
All checks were successful
PikaOS Package Build & Release (amd64-v3 on debian) / build (push) Successful in 12m11s
2024-09-15 21:37:31 +03:00
Ward from fusion-voyager-3
e232b4a5e0 shti
Some checks failed
PikaOS Package Build & Release (amd64-v3) / build (push) Failing after 3m29s
2024-09-15 21:32:41 +03:00
Ward from fusion-voyager-3
ae391fb96d fix build dep
Some checks failed
PikaOS Package Build & Release (amd64-v3) / build (push) Failing after 1m34s
2024-09-15 21:27:16 +03:00
Ward from fusion-voyager-3
6448fcfd7c update rules
Some checks failed
PikaOS Package Build & Release (amd64-v3) / build (push) Failing after 2m11s
2024-09-15 21:20:51 +03:00
1e6c42e4f2 Update .github/release-nest-v3
Some checks failed
PikaOS Package Build & Release (amd64-v3) / build (push) Failing after 2m2s
2024-09-15 20:09:50 +02:00
f879753ee0 Update .github/release-nest-i386
All checks were successful
PikaOS Package Build & Release (i386) / build (push) Successful in 5m39s
2024-08-11 22:23:37 +02:00
80d57d65fd Update .github/release-nest-v3
All checks were successful
PikaOS Package Build & Release (amd64-v3) / build (push) Successful in 14m26s
2024-08-11 22:23:29 +02:00
6a77a70c41 Update debian/source/format 2024-08-11 22:23:18 +02:00
0cb3553e44 Update main.sh 2024-08-11 22:23:02 +02:00
409c37a911 Update .github/release-nest-v3
Some checks failed
PikaOS Package Build & Release (amd64-v3) / build (push) Failing after 3m10s
2024-08-11 02:43:37 +02:00
2560c665fb Update .github/release-nest-i386
Some checks failed
PikaOS Package Build & Release (i386) / build (push) Failing after 2m47s
2024-08-11 02:43:25 +02:00
e20c20d5b0 Update debian/source/format 2024-08-11 02:43:08 +02:00
3cbb30be2b Update .github/release-nest-i386
Some checks failed
PikaOS Package Build & Release (i386) / build (push) Has been cancelled
2024-08-11 02:42:54 +02:00
fb54a47121 Update main.sh 2024-08-11 02:42:35 +02:00
57f8b4a63d Update .github/release-nest-v3
Some checks failed
PikaOS Package Build & Release (amd64-v3) / build (push) Failing after 1m59s
2024-08-11 02:39:33 +02:00
92c66d1762 Update .github/release-nest-i386
Some checks failed
PikaOS Package Build & Release (i386) / build (push) Failing after 2m0s
2024-08-11 02:39:24 +02:00
8f6a044740 Update main.sh 2024-08-11 02:39:08 +02:00
7d5b6fa68d Update .github/release-nest-i386
Some checks failed
PikaOS Package Build & Release (i386) / build (push) Failing after 1m40s
2024-08-11 02:32:11 +02:00
bc897780a9 Update .github/release-nest-v3
Some checks failed
PikaOS Package Build & Release (amd64-v3) / build (push) Failing after 1m43s
2024-08-11 02:32:05 +02:00
943bada916 Update main.sh 2024-08-11 02:31:47 +02:00
bd63bccb56 Update .github/release-nest-i386
Some checks failed
PikaOS Package Build & Release (i386) / build (push) Has been cancelled
2024-08-11 02:04:14 +02:00
b2878bc88d Update .github/release-nest-v3
All checks were successful
PikaOS Package Build & Release (amd64-v3) / build (push) Successful in 14m26s
2024-08-11 02:04:00 +02:00
cec18fe57f Update .github/release-canary-i386
All checks were successful
PikaOS Package Build & Release (Canary) (i386) / build (push) Successful in 5m37s
2024-08-09 20:23:25 +02:00
Ward from fusion-voyager-3
68a0e48ddc fix docs
All checks were successful
PikaOS Package Build & Release (Canary) (amd64-v3) / build (push) Successful in 11m30s
2024-07-29 23:12:30 +03:00
Ward from fusion-voyager-3
211dea2b75 fix weird distro check
Some checks failed
PikaOS Package Build & Release (Canary) (amd64-v3) / build (push) Failing after 10m40s
2024-07-29 22:45:04 +03:00
Ward from fusion-voyager-3
d0889bfa96 port to pika os 4
Some checks failed
PikaOS Package Build & Release (Canary) (amd64-v3) / build (push) Failing after 9m24s
2024-07-29 18:47:23 +03:00
Ward from fusion-voyager-3
4f1e3321c7 port to pika os 4
Some checks failed
PikaOS Package Build Only (Canary) (i386) / build (push) Failing after 1s
PikaOS Package Build Only (Canary) (amd64-v3) / build (push) Has been cancelled
PikaOS Package Build Only (i386) / build (push) Failing after 1s
PikaOS Package Build Only (amd64-v3) / build (push) Failing after 1s
PikaOS Package Build & Release (Canary) (i386) / build (push) Failing after 1s
PikaOS Package Build & Release (Canary) (amd64-v3) / build (push) Failing after 2m5s
PikaOS Package Build & Release (i386) / build (push) Failing after 0s
PikaOS Package Build & Release (amd64-v3) / build (push) Failing after 0s
2024-07-29 18:38:10 +03:00
f6f9040ac7
Update main32.sh 2024-04-09 17:43:32 +01:00
4bcd0ce687
Update main.sh 2024-04-09 17:43:18 +01:00
b0a52830f6
Update main32.sh 2024-02-27 15:04:51 +00:00
45ea9846e4
Update main.sh 2024-02-27 15:04:38 +00:00
1ef2e09a3f
Update main.sh 2024-02-27 15:02:39 +00:00
4921128572
Update control 2024-02-27 14:51:05 +00:00
Ward Nakchbandi (Cosmic Fusion)
6da9ce8095
Update main32.sh 2024-01-23 07:52:55 +03:00
Ward Nakchbandi (Cosmic Fusion)
f4f874c035
Update main.sh 2024-01-23 07:52:30 +03:00
Ward Nakchbandi (Cosmic Fusion)
68f06c7e03
Update rules 2023-11-27 21:17:35 +03:00
Ward Nakchbandi (Cosmic Fusion)
f958f8c846
Update main32.sh 2023-11-27 21:16:21 +03:00
Ward Nakchbandi (Cosmic Fusion)
64482dcdda
Update main.sh 2023-11-27 21:16:06 +03:00
Ward Nakchbandi (Cosmic Fusion)
f1dc2a4b02
Update rules 2023-11-27 20:49:34 +03:00
Ward Nakchbandi (Cosmic Fusion)
aee983e8f0
Update mesa-common-git-dev.docs 2023-11-27 20:06:14 +03:00
Ward Nakchbandi (Cosmic Fusion)
879d836f52
Update rules 2023-11-27 19:39:22 +03:00
Ward Nakchbandi (Cosmic Fusion)
41bf4917a2
Update control.in 2023-11-22 19:48:20 +03:00
Ward Nakchbandi (Cosmic Fusion)
5e2bd3b231
Update control 2023-11-22 19:47:21 +03:00
Ward from fusion-voyager-3
39daf5aa26 24.0 2023-11-22 19:22:20 +03:00
Ward Nakchbandi (Cosmic Fusion)
854c3b1ff8
2023-10-06 17:29:24 +03:00
Ward Nakchbandi (Cosmic Fusion)
50b2c25787
2023-10-06 17:13:23 +03:00
Ward Nakchbandi (Cosmic Fusion)
9d266840d2
2023-10-05 22:52:53 +03:00
Ward Nakchbandi (Cosmic Fusion)
1f12539671
2023-10-05 22:22:50 +03:00
45 changed files with 637 additions and 3271 deletions

1
.github/build-canary-i386 vendored Normal file
View File

@ -0,0 +1 @@
1

1
.github/build-canary-v3 vendored Normal file
View File

@ -0,0 +1 @@
1

1
.github/build-nest-i386 vendored Normal file
View File

@ -0,0 +1 @@
1

1
.github/build-nest-v3 vendored Normal file
View File

@ -0,0 +1 @@
1

1
.github/release-canary-i386 vendored Normal file
View File

@ -0,0 +1 @@
2

1
.github/release-canary-v3 vendored Normal file
View File

@ -0,0 +1 @@
4

1
.github/release-debian-v3 vendored Normal file
View File

@ -0,0 +1 @@
13

1
.github/release-nest-i386 vendored Normal file
View File

@ -0,0 +1 @@
14

1
.github/release-nest-v3 vendored Normal file
View File

@ -0,0 +1 @@
14

37
.github/workflows/build-canaryi386.yml vendored Normal file
View File

@ -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

37
.github/workflows/build-canaryv3.yml vendored Normal file
View File

@ -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

37
.github/workflows/build-nesti386.yml vendored Normal file
View File

@ -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

37
.github/workflows/build-nestv3.yml vendored Normal file
View File

@ -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

View File

@ -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

40
.github/workflows/release-canaryv3.yml vendored Normal file
View File

@ -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

40
.github/workflows/release-debian-v3.yml vendored Executable file
View File

@ -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

40
.github/workflows/release-nesti386.yml vendored Normal file
View File

@ -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

40
.github/workflows/release-nestv3.yml vendored Normal file
View File

@ -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

View File

@ -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 }}

View File

@ -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 }}

24
LICENSE.md Normal file
View File

@ -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.

71
debian/control vendored
View File

@ -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},

38
debian/control.in vendored
View File

@ -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/

View File

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

View File

@ -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

View File

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

File diff suppressed because it is too large Load Diff

View File

@ -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

View File

@ -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

View File

@ -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.*

View File

@ -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

View File

@ -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;
}
--

View File

@ -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

View File

@ -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

View File

@ -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

View File

@ -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

207
debian/rules vendored
View File

@ -13,84 +13,77 @@ 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_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)
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
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
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
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
GALLIUM_DRIVERS += freedreno
# XXX - broken
#GALLIUM_DRIVERS += asahi
endif
# etnaviv, tegra, vc4 and v3d kernel support are only available on armhf and arm64
@ -99,52 +92,42 @@ else
VULKAN_DRIVERS += broadcom freedreno panfrost
endif
ifneq (,$(filter $(DEB_HOST_ARCH), armhf arm64 riscv64))
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 x32))
GALLIUM_DRIVERS += svga
GALLIUM_DRIVERS += crocus i915 iris svga
VULKAN_DRIVERS += intel intel_hasvk
# 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)))
ifneq (,$(filter $(DEB_HOST_ARCH), $(WSL_ARCHS)))
GALLIUM_DRIVERS += d3d12
# VULKAN_DRIVERS += microsoft-experimental
endif
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
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)))
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

24
main.sh
View File

@ -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

View File

@ -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/

2
patches/series Normal file
View File

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

10
pika-build-config/amd64-v3.sh Executable file
View File

@ -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

5
pika-build-config/i386.sh Executable file
View File

@ -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

View File

@ -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/