diff --git a/.github/workflows/build-linux-turnip.yml b/.github/workflows/build-linux-turnip.yml new file mode 100644 index 0000000..3d65229 --- /dev/null +++ b/.github/workflows/build-linux-turnip.yml @@ -0,0 +1,142 @@ +name: Build WN Linux Turnip Drivers + +on: + workflow_call: + inputs: + source_ref: + type: string + default: feature/linux-drivers + publish: + type: boolean + default: false + schedule: + - cron: '0 12 * * 3' + push: + branches: [feature/linux-drivers] + paths: ['linux/**', 'patches/**', 'build_linux_turnip.sh', '.github/workflows/build-linux-turnip.yml'] + pull_request: + paths: ['linux/**', 'patches/**', 'build_linux_turnip.sh', '.github/workflows/build-linux-turnip.yml'] + workflow_dispatch: + inputs: + publish: + description: Publish the Linux driver release + type: boolean + default: false + +concurrency: + group: linux-turnip-${{ github.event_name == 'pull_request' && github.event.pull_request.number || 'release' }} + cancel-in-progress: false + +permissions: + contents: read + +jobs: + build: + runs-on: ubuntu-24.04-arm + timeout-minutes: 45 + outputs: + version: ${{ steps.version.outputs.version }} + commit: ${{ steps.mesa.outputs.commit }} + mesa_version: ${{ steps.mesa.outputs.version }} + recipe_commit: ${{ steps.version.outputs.recipe_commit }} + steps: + - uses: actions/checkout@v4 + with: + ref: ${{ inputs.source_ref || github.sha }} + persist-credentials: false + - name: Version + id: version + env: + GH_TOKEN: ${{ github.token }} + run: | + set -euo pipefail + python3 -m unittest discover -s linux -p 'test_*.py' + gh api --paginate --slurp "repos/${GITHUB_REPOSITORY}/releases?per_page=100" > releases.json + echo "version=$(python3 linux/release_version.py < releases.json)" >> "$GITHUB_OUTPUT" + echo "recipe_commit=$(git rev-parse HEAD)" >> "$GITHUB_OUTPUT" + - name: Dependencies + run: | + sudo apt-get update + sudo apt-get install -y --no-install-recommends build-essential git python3-pip ninja-build pkg-config \ + flex bison zip libdrm-dev libwayland-dev wayland-protocols libx11-xcb-dev libxcb-dri3-dev \ + libxcb-present-dev libxcb-randr0-dev libxcb-sync-dev libxcb-xfixes0-dev libxshmfence-dev \ + libxrandr-dev libxcb-keysyms1-dev libxcb-shm0-dev libzstd-dev zlib1g-dev libexpat1-dev glslang-tools + python3 -m pip install --break-system-packages 'meson>=1.5' mako pyyaml packaging + - name: Build balanced and performance + env: + BUILD_VERSION: ${{ steps.version.outputs.version }} + BUILD_JOBS: '4' + run: | + set -euo pipefail + ./build_linux_turnip.sh 2>&1 | tee linux-build.log + - name: Mesa provenance + id: mesa + run: | + echo "commit=$(cat linux_work/mesa_commit.txt)" >> "$GITHUB_OUTPUT" + echo "version=$(cat linux_work/mesa_version.txt)" >> "$GITHUB_OUTPUT" + - uses: actions/upload-artifact@v4 + with: + name: WN-Linux-Turnip-${{ steps.version.outputs.version }}-b_Axxx + path: WN-Linux-Turnip-${{ steps.version.outputs.version }}-b_Axxx.zip + if-no-files-found: error + - uses: actions/upload-artifact@v4 + with: + name: WN-Linux-Turnip-${{ steps.version.outputs.version }}-p_Axxx + path: WN-Linux-Turnip-${{ steps.version.outputs.version }}-p_Axxx.zip + if-no-files-found: error + - uses: actions/upload-artifact@v4 + if: always() + with: + name: linux-build-provenance-${{ steps.version.outputs.version }} + path: | + linux-build.log + linux_work/patch-*.log + linux_work/mesa*.txt + linux_work/mesa-*.patch + linux_work/package-*/meta.json + if-no-files-found: warn + + release: + needs: build + if: github.event_name == 'schedule' || github.event_name == 'workflow_dispatch' + runs-on: ubuntu-latest + permissions: + contents: write + steps: + - uses: actions/download-artifact@v4 + with: + pattern: WN-Linux-Turnip-${{ needs.build.outputs.version }}-*_Axxx + merge-multiple: true + - name: Release + env: + GH_TOKEN: ${{ github.token }} + VERSION: ${{ needs.build.outputs.version }} + MESA_COMMIT: ${{ needs.build.outputs.commit }} + MESA_VERSION: ${{ needs.build.outputs.mesa_version }} + RECIPE_COMMIT: ${{ needs.build.outputs.recipe_commit }} + PUBLISH: ${{ github.event_name == 'schedule' || inputs.publish }} + run: | + set -euo pipefail + cat > notes.md < releases.json + existing=$(jq -r --arg tag "linux-v${VERSION}" '[.[][] | select(.tag_name == $tag)] | if length == 0 then "absent" elif .[0].draft then "draft" else "published" end' releases.json) + if [[ "$existing" == published ]]; then + echo "Refusing to replace published linux-v${VERSION}" >&2 + exit 1 + fi + if [[ "$existing" == draft ]]; then + gh release delete "linux-v${VERSION}" --repo "$GITHUB_REPOSITORY" --yes + fi + gh release create "linux-v${VERSION}" WN-Linux-Turnip-*.zip --repo "$GITHUB_REPOSITORY" \ + --target "$RECIPE_COMMIT" --title "WN Linux Turnip ${VERSION}" --notes-file notes.md "${options[@]}" diff --git a/.gitignore b/.gitignore index a24dc06..daec851 100644 --- a/.gitignore +++ b/.gitignore @@ -27,3 +27,5 @@ build_full_*.log # mesa reference checkout mesa_ref/ mesa_version.txt + +linux_work/ diff --git a/README.md b/README.md index 3a3f258..97a20ff 100644 --- a/README.md +++ b/README.md @@ -198,3 +198,7 @@ implemented and device-confirmed in [`Leb-Sun/Drivers`](https://github.com/Leb-S ## License MIT — see `LICENSE`. + +## Linux GameScope/Wayland drivers + +The `feature/linux-drivers` branch adds glibc ARM64 Turnip packages alongside the existing Android builds. See [linux/README.md](linux/README.md) for the build, package format, validation and scheduling requirements. Android build scripts and releases keep their existing behavior. diff --git a/build_linux_turnip.sh b/build_linux_turnip.sh new file mode 100755 index 0000000..bb8f090 --- /dev/null +++ b/build_linux_turnip.sh @@ -0,0 +1,86 @@ +#!/usr/bin/env bash +set -euo pipefail +here=$(cd "$(dirname "$0")" && pwd) +cd "$here" +version=${BUILD_VERSION:-0.1.0} +[[ "$version" =~ ^[a-zA-Z0-9._-]+$ ]] +work="$here/linux_work" +mkdir -p "$work" +if [[ -z "${MESA_LOCAL_SRC:-}" ]]; then + rm -rf "$work/upstream" + git clone --depth=1 --branch main https://gitlab.freedesktop.org/mesa/mesa.git "$work/upstream" + MESA_LOCAL_SRC="$work/upstream" +fi +source_dir=$(cd "$MESA_LOCAL_SRC" && pwd) +commit=$(git -C "$source_dir" rev-parse HEAD) +mesa_version=$(tr -d '[:space:]' < "$source_dir/VERSION") +printf '%s\n' "$commit" > "$work/mesa_commit.txt" +printf '%s\n' "$mesa_version" > "$work/mesa_version.txt" +read -ra variants <<< "${BUILD_VARIANTS:-b p}" +for variant in "${variants[@]}"; do + [[ "$variant" == b || "$variant" == p ]] + src="$work/mesa-$variant" + rm -rf "$src" "$work/build-$variant" "$work/package-$variant" + git clone --no-hardlinks --shared "$source_dir" "$src" + git -C "$src" checkout --detach "$commit" + ( + cd "$src" + for script in fix_a8xx_dev_info apply_a8xx_gpus apply_a7xx_gen1_quirks apply_a7xx_gen2_ubwc_hint; do + python3 "$here/patches/$script.py" + done + if [[ "$variant" == b ]]; then + python3 "$here/patches/apply_balance_variant.py" + else + BUILD_VARIANT=p python3 "$here/patches/apply_perf_variant.py" + fi + python3 "$here/linux/patch_mesa.py" + for patch in "$here"/linux/patches/*.patch; do + git apply "$patch" + echo "Applied $(basename "$patch")" + done + ) 2>&1 | tee "$work/patch-$variant.log" + python3 "$here/linux/verify_patches.py" "$src" "$variant" "$work/patch-$variant.log" + git -C "$src" diff --binary > "$work/mesa-$variant.patch" + options=() + if [[ -n "${LINUX_SYSROOT:-}" ]]; then + rootfs=$(cd "$LINUX_SYSROOT" && pwd) + cat > "$work/cross.ini" <=1.5, Python Mako/PyYAML, Ninja, C/C++ compilers, binutils, pkg-config, zip, flex, bison, glslang, and the Wayland/XCB/libdrm development packages are required. Mesa's checked-in wrap downloads verified Wayland protocols if the distro version is too old. + +For a cross build using the existing WinNative Linux rootfs: + +```sh +LINUX_SYSROOT=/absolute/path/to/rootfs ./build_linux_turnip.sh +``` + +The cross build additionally needs `aarch64-linux-gnu-gcc/g++`, ARM64 binutils and a host `wayland-scanner`. Rootfs development headers and pkg-config files must be available. `MESA_LOCAL_SRC` optionally supplies a Git checkout; both variants always use its HEAD. `BUILD_VERSION`, `BUILD_VARIANTS="b p"` and `BUILD_JOBS` control version, variants and parallelism. Build output is under `linux_work/`; distributable ZIPs are in the repository root. + +Balanced (`-b`) uses the WN GPU fixes and GMEM bandwidth multiplier 10. Performance (`-p`) also requests KGSL PWR_MAX, sets the context/submission power-constraint flags, and refreshes the request every 1000 submissions. The upstream draw-call threshold was restructured; that retired tweak is deliberately skipped. Android gralloc/IMapper/AHB patches are excluded because this build uses glibc and Linux WSI. + +`patch_mesa.py` ports the two WinNative `tools/linuxfs/turnip` fixes: KGSL dma-buf feedback device reporting, and disabling unsupported calibrated timestamps/present timing. It fails when its source anchors change. `verify_patches.py` also rejects warnings/missing patch anchors, aside from the retired draw-call threshold. Update and re-review patches when upstream changes; do not suppress failures to produce a release. + +`linux/patches/` holds Mesa commits applied with `git apply` after the scripts, in name order. They are needed by DX12 Ultimate games such as FINAL FANTASY VII REBIRTH: + +- `0001` implements `VK_EXT_mesh_shader` by running task and mesh shaders as compute into a device ring that a generated vertex shader draws. Render passes with mesh draws use sysmem rendering. +- `0002` supports a required subgroup size of 32 on A8XX's 64-wide waves, so vkd3d-proton accepts `WaveSize(32)` shaders. Each half of a wave acts as one subgroup. +- `0003` sanitizes cube map sampling directions on A8XX. A NaN or zero direction can hang the GPU. +- `0004` invalidates bindless descriptors through A8XX's dedicated registers, because A8XX's `SP_UPDATE_CNTL` has no bindless bits. +- `0005` fetches A8XX command streams through a KGSL virtual BO alias that is unbound before the memory is freed. Freeing command stream memory directly leads to GPU hangs. + +A patch that no longer applies fails the build. Rebase it in a Mesa checkout, regenerate it with `git format-patch`, and re-run the `dEQP-VK.mesh_shader.ext.*`, `dEQP-VK.subgroups.*` and `dEQP-VK.texture.*cube*` groups on an Adreno device. Changes to `0005` also need a long game session on A8XX, since the hangs it prevents take minutes to appear. + +Displayed names start at **WN Linux Turnip 0.1.0-b** and **WN Linux Turnip 0.1.0-p**, packaged as `WN-Linux-Turnip-0.1.0-b_Axxx.zip` and `WN-Linux-Turnip-0.1.0-p_Axxx.zip`. Linux has its own semantic version series. CI increments the patch component after the highest published stable Linux release; previews and drafts reuse the next unpublished version. A repeated draft build replaces that draft, while published releases cannot be replaced. Local builds default to `0.1.0`; set `BUILD_VERSION` explicitly for another version. + +## ZIP contract + +Each `WN-Linux-Turnip--{b,p}_Axxx.zip` contains a flat `meta.json` and `libvulkan_freedreno.so`. Metadata identifies `platform=linux`, `architecture=aarch64`, `libc=glibc`, `variant`, Mesa repository/SHA/version, build-recipe commit and SHA-256 of the library. The app creates its own ICD manifest after installing, so no device-specific absolute paths are distributed. + +The `WN-Linux-` prefix reserves a separate catalog namespace. Linux releases use `linux-v...` tags and do not change the repository's latest Android release. Neither the Android mirror nor Android release notes are modified. + +The app detects actual library ABI through ELF dependencies and corrects an Android/Linux destination mistake. Packages with contradictory platform metadata or an incorrect digest are rejected. New installs select that Linux driver; Settings can select another or return to bundled Mesa. The old Components/Linux Client driver remains separately available. + +## Validation and provenance + +The script verifies that all variants use one Mesa SHA, their binaries differ, they are ARM64/glibc libraries with Wayland and X11 entry points, and only `-p` contains the power-constraint code. CI uploads source diffs, patch logs, metadata and build logs along with separate variant artifacts. The app's package installation and UI tests run on AVD; rendering and KGSL power behavior require physical Adreno hardware. + +Initial local validation built Mesa `5ff61a7646d29b54c324af0a60aa3bfb5cdd24d1`, version `26.3.0-devel`, against the existing Arch Linux ARM sysroot. Both variants passed package validation, and the Linux loader resolved the performance build's dependencies against that rootfs. Existing bundled Mesa 26.2.2 is preserved as fallback in the app. + +## Release workflow + +`build-linux-turnip.yml` builds both variants on Ubuntu 24.04 ARM64. Branch pushes build preview artifacts. Manual runs create a draft unless `publish` is true; scheduled runs publish. The cron is the same as the Android workflow: Wednesday 12:00 UTC (`0 12 * * 3`). + +GitHub executes scheduled workflows only from the default branch. The feature branch's cron is not active by itself. To activate weekly builds, merge this workflow into the default branch or add a default-branch entrypoint that checks out the Linux branch. Preserve the Android workflow and its existing schedule. After the push registers the workflow, manual dispatch with `--ref feature/linux-drivers` works; this was verified through the GitHub API. Push builds work on this branch immediately. + +A ready default-branch entrypoint is provided at `linux/scheduler/schedule-linux-turnip.yml`. Install that file as `.github/workflows/schedule-linux-turnip.yml` on `main` to call the reusable build from `feature/linux-drivers`. Its manual trigger defaults to a draft, while weekly runs publish. The reusable workflow checks out the Linux branch and records its actual recipe commit for release tags, even when the caller runs from main. Do not install both this scheduler and a separately scheduled copy of the full workflow on main. diff --git a/linux/package.py b/linux/package.py new file mode 100644 index 0000000..5090730 --- /dev/null +++ b/linux/package.py @@ -0,0 +1,27 @@ +#!/usr/bin/env python3 +import hashlib +import json +from pathlib import Path +import subprocess +import sys + +output, version, variant, commit, mesa_version = sys.argv[1:] +path = Path(output) +meta = { + 'schemaVersion': 1, + 'name': f'WN Linux Turnip {version}-{variant}', + 'description': 'Turnip for the WinNative glibc ARM64 GameScope/Wayland runtime', + 'author': 'WinNative', + 'platform': 'linux', + 'architecture': 'aarch64', + 'libc': 'glibc', + 'driverVersion': f'{version}-{variant}', + 'libraryName': 'libvulkan_freedreno.so', + 'variant': variant, + 'mesaVersion': mesa_version, + 'mesaCommit': commit, + 'mesaSource': 'https://gitlab.freedesktop.org/mesa/mesa', + 'buildCommit': subprocess.check_output(['git', 'rev-parse', 'HEAD'], text=True).strip(), + 'librarySha256': hashlib.sha256((path / 'libvulkan_freedreno.so').read_bytes()).hexdigest(), +} +(path / 'meta.json').write_text(json.dumps(meta, indent=2) + '\n') diff --git a/linux/patch_mesa.py b/linux/patch_mesa.py new file mode 100644 index 0000000..9528726 --- /dev/null +++ b/linux/patch_mesa.py @@ -0,0 +1,40 @@ +#!/usr/bin/env python3 +from pathlib import Path + + +def replace(path, old, new): + file = Path(path) + text = file.read_text() + if new in text: + return + if text.count(old) != 1: + raise SystemExit(f"Unsupported Mesa source: {path}: {old!r}") + file.write_text(text.replace(old, new, 1)) + + +kgsl = 'src/freedreno/vulkan/tu_knl_kgsl.cc' +device = 'src/freedreno/vulkan/tu_device.cc' +replace(kgsl, '#include ', '#include \n#include \n#include ') +replace(kgsl, ' result = tu_physical_device_init(device, instance);', ''' struct stat kgsl_st; + if (fstat(fd, &kgsl_st) == 0 && S_ISCHR(kgsl_st.st_mode)) { + device->has_local = device->has_master = true; + device->local_major = device->master_major = major(kgsl_st.st_rdev); + device->local_minor = device->master_minor = minor(kgsl_st.st_rdev); + } + + result = tu_physical_device_init(device, instance);''') +replace(device, '.EXT_physical_device_drm = !is_kgsl(device->instance),', '.EXT_physical_device_drm = !is_kgsl(device->instance) || device->has_local,') +text = Path(kgsl).read_text() +function = 'static int\nkgsl_device_get_gpu_timestamp(struct tu_device *dev, uint64_t *ts)\n{\n UNREACHABLE("");\n return 0;\n}\n\n' +if function not in text: + raise SystemExit('KGSL timestamp implementation changed; review before building') +text = text.replace(function, '').replace(' .device_get_gpu_timestamp = kgsl_device_get_gpu_timestamp,\n', '') +Path(kgsl).write_text(text) +replace('src/freedreno/vulkan/tu_knl.cc', ' return dev->instance->knl->device_get_gpu_timestamp(dev, ts);', ''' if (dev->instance->knl->device_get_gpu_timestamp == NULL) + return -1; + return dev->instance->knl->device_get_gpu_timestamp(dev, ts);''') +for extension in ('KHR_calibrated_timestamps', 'EXT_calibrated_timestamps', 'EXT_present_timing'): + replace(device, f'.{extension} = device->info->props.has_persistent_counter,', f'.{extension} = device->info->props.has_persistent_counter && device->instance->knl->device_get_gpu_timestamp != NULL,') +for feature in ('presentTiming', 'presentAtRelativeTime', 'presentAtAbsoluteTime'): + replace(device, f'features->{feature} = true;', f'features->{feature} = pdevice->vk.supported_extensions.EXT_present_timing;') +print('Linux KGSL Wayland and timestamp fixes verified') diff --git a/linux/patches/0001-tu-Emulate-VK_EXT_mesh_shader-with-compute.patch b/linux/patches/0001-tu-Emulate-VK_EXT_mesh_shader-with-compute.patch new file mode 100644 index 0000000..0881724 --- /dev/null +++ b/linux/patches/0001-tu-Emulate-VK_EXT_mesh_shader-with-compute.patch @@ -0,0 +1,2733 @@ +From 0000000000000000000000000000000000000000 Mon Sep 17 00:00:00 2001 +From: MaxsTechReview +Date: Sun, 27 Sep 2026 02:49:01 -0400 +Subject: [PATCH 1/5] tu: Emulate VK_EXT_mesh_shader with compute + +Task and mesh shaders are compiled as compute shaders that write their +outputs into a device-wide ring. A generated vertex shader then pulls one +vertex per primitive corner out of the ring in a non-indexed draw, so render +passes containing mesh draws run in sysmem mode. Indirect and indirect count +draws build their workgroup tables on the GPU and are split into chunks +predicated on those tables. +--- + src/compiler/nir/nir_divergence_analysis.c | 1 + + src/compiler/nir/nir_intrinsics.py | 3 + + src/freedreno/vulkan/meson.build | 1 + + src/freedreno/vulkan/tu_cmd_buffer.cc | 675 ++++++++++++-- + src/freedreno/vulkan/tu_cmd_buffer.h | 9 +- + src/freedreno/vulkan/tu_common.h | 2 + + src/freedreno/vulkan/tu_device.cc | 68 +- + src/freedreno/vulkan/tu_device.h | 6 + + src/freedreno/vulkan/tu_mesh.cc | 992 +++++++++++++++++++++ + src/freedreno/vulkan/tu_mesh.h | 127 +++ + src/freedreno/vulkan/tu_pipeline.cc | 33 +- + src/freedreno/vulkan/tu_pipeline.h | 12 +- + src/freedreno/vulkan/tu_shader.cc | 188 +++- + src/freedreno/vulkan/tu_shader.h | 13 +- + 14 files changed, 2034 insertions(+), 96 deletions(-) + create mode 100644 src/freedreno/vulkan/tu_mesh.cc + create mode 100644 src/freedreno/vulkan/tu_mesh.h + +diff --git a/src/compiler/nir/nir_divergence_analysis.c b/src/compiler/nir/nir_divergence_analysis.c +index d45de5b..63ce20f 100644 +--- a/src/compiler/nir/nir_divergence_analysis.c ++++ b/src/compiler/nir/nir_divergence_analysis.c +@@ -353,6 +353,7 @@ visit_intrinsic(nir_intrinsic_instr *instr, struct divergence_state *state) + case nir_intrinsic_load_hs_patch_stride_ir3: + case nir_intrinsic_load_tess_factor_base_ir3: + case nir_intrinsic_load_tess_param_base_ir3: ++ case nir_intrinsic_load_mesh_ring_ir3: + case nir_intrinsic_load_primitive_location_ir3: + case nir_intrinsic_preamble_start_ir3: + case nir_intrinsic_optimization_barrier_sgpr_amd: +diff --git a/src/compiler/nir/nir_intrinsics.py b/src/compiler/nir/nir_intrinsics.py +index c80a232..463e0c2 100644 +--- a/src/compiler/nir/nir_intrinsics.py ++++ b/src/compiler/nir/nir_intrinsics.py +@@ -1638,6 +1638,9 @@ system_value("tess_param_base_ir3", 2) + system_value("tcs_header_ir3", 1) + system_value("rel_patch_id_ir3", 1) + ++# Base address of turnip's mesh shading ring. ++system_value("mesh_ring_ir3", 2) ++ + # System values for freedreno compute shaders. + system_value("subgroup_id_shift_ir3", 1) + +diff --git a/src/freedreno/vulkan/meson.build b/src/freedreno/vulkan/meson.build +index b019710..4647b23 100644 +--- a/src/freedreno/vulkan/meson.build ++++ b/src/freedreno/vulkan/meson.build +@@ -49,6 +49,7 @@ libtu_files = files( + 'tu_image.cc', + 'tu_knl.cc', + 'tu_lrz.cc', ++ 'tu_mesh.cc', + 'tu_nir_lower_demote_samples.cc', + 'tu_nir_lower_multiview.cc', + 'tu_nir_lower_ray_query.cc', +diff --git a/src/freedreno/vulkan/tu_cmd_buffer.cc b/src/freedreno/vulkan/tu_cmd_buffer.cc +index 01f0aa6..f87b236 100644 +--- a/src/freedreno/vulkan/tu_cmd_buffer.cc ++++ b/src/freedreno/vulkan/tu_cmd_buffer.cc +@@ -1346,12 +1346,21 @@ tu6_update_msaa(struct tu_cmd_buffer *cmd) + tu6_emit_msaa(&cmd->draw_cs, samples, cmd->state.msaa_disable); + } + ++static VkPrimitiveTopology ++tu_primitive_topology(const struct tu_cmd_buffer *cmd) ++{ ++ const struct tu_shader *ms = cmd->state.shaders[MESA_SHADER_MESH]; ++ if (ms) ++ return (VkPrimitiveTopology) ms->mesh.topology; ++ return (VkPrimitiveTopology) ++ cmd->vk.dynamic_graphics_state.ia.primitive_topology; ++} ++ + template + static void + tu6_update_msaa_disable(struct tu_cmd_buffer *cmd) + { +- VkPrimitiveTopology topology = +- (VkPrimitiveTopology)cmd->vk.dynamic_graphics_state.ia.primitive_topology; ++ VkPrimitiveTopology topology = tu_primitive_topology(cmd); + bool is_line = + topology == VK_PRIMITIVE_TOPOLOGY_LINE_LIST || + topology == VK_PRIMITIVE_TOPOLOGY_LINE_LIST_WITH_ADJACENCY || +@@ -1453,6 +1462,11 @@ use_sysmem_rendering(struct tu_cmd_buffer *cmd, + return true; + } + ++ if (cmd->state.rp.has_mesh) { ++ cmd->state.rp.force_render_mode_reason = "Uses mesh shaders"; ++ return true; ++ } ++ + if (cmd->state.rp.disable_gmem) { + /* force_render_mode_reason is set where disable_gmem is set. */ + return true; +@@ -4767,6 +4781,49 @@ tu_CmdBindIndexBuffer2KHR(VkCommandBuffer commandBuffer, + } + TU_GENX(tu_CmdBindIndexBuffer2KHR); + ++/* Points the bindless base registers of the graphics or compute stages at ++ * the given descriptor sets. ++ */ ++template ++static void ++tu6_emit_bindless_bases(struct tu_cmd_buffer *cmd, struct tu_cs *cs, ++ const struct tu_descriptor_state *descriptors_state, ++ bool compute) ++{ ++ uint32_t sp_bindless_base_reg = compute ? ++ __SP_CS_BINDLESS_BASE_DESCRIPTOR(0, {}).reg : ++ __SP_GFX_BINDLESS_BASE_DESCRIPTOR(0, {}).reg; ++ uint32_t hlsq_bindless_base_reg = compute ? ++ REG_A6XX_HLSQ_CS_BINDLESS_BASE(0) : REG_A6XX_HLSQ_BINDLESS_BASE(0); ++ ++ if (descriptors_state->max_sets_bound > 0) { ++ tu_cs_emit_pkt4(cs, sp_bindless_base_reg, 2 * descriptors_state->max_sets_bound); ++ tu_cs_emit_array(cs, (const uint32_t*)descriptors_state->set_iova, 2 * descriptors_state->max_sets_bound); ++ if (CHIP == A6XX) { ++ tu_cs_emit_pkt4(cs, hlsq_bindless_base_reg, 2 * descriptors_state->max_sets_bound); ++ tu_cs_emit_array(cs, (const uint32_t*)descriptors_state->set_iova, 2 * descriptors_state->max_sets_bound); ++ } ++ } ++ ++ /* Dynamic descriptors get the reserved descriptor set. */ ++ if (descriptors_state->max_dynamic_offset_size) { ++ int reserved_set_idx = cmd->device->physical_device->reserved_set_idx; ++ assert(reserved_set_idx >= 0); /* reserved set must be bound */ ++ ++ tu_cs_emit_pkt4(cs, sp_bindless_base_reg + reserved_set_idx * 2, 2); ++ tu_cs_emit_qw(cs, descriptors_state->set_iova[reserved_set_idx]); ++ if (CHIP == A6XX) { ++ tu_cs_emit_pkt4(cs, hlsq_bindless_base_reg + reserved_set_idx * 2, 2); ++ tu_cs_emit_qw(cs, descriptors_state->set_iova[reserved_set_idx]); ++ } ++ } ++ ++ tu_cs_emit_regs(cs, SP_UPDATE_CNTL(CHIP, ++ .cs_bindless = compute ? CHIP == A6XX ? 0x1f : 0xff : 0, ++ .gfx_bindless = !compute ? CHIP == A6XX ? 0x1f : 0xff : 0, ++ )); ++} ++ + template + static void + tu6_emit_descriptor_sets(struct tu_cmd_buffer *cmd, +@@ -4774,13 +4831,9 @@ tu6_emit_descriptor_sets(struct tu_cmd_buffer *cmd, + { + struct tu_descriptor_state *descriptors_state = + tu_get_descriptors_state(cmd, bind_point); +- uint32_t sp_bindless_base_reg, hlsq_bindless_base_reg; + struct tu_cs *cs, state_cs; + + if (bind_point == VK_PIPELINE_BIND_POINT_GRAPHICS) { +- sp_bindless_base_reg = __SP_GFX_BINDLESS_BASE_DESCRIPTOR(0, {}).reg; +- hlsq_bindless_base_reg = REG_A6XX_HLSQ_BINDLESS_BASE(0); +- + unsigned bindless_pkt_size = descriptors_state->max_sets_bound ? + 1 + 2 * descriptors_state->max_sets_bound : + 0; +@@ -4799,39 +4852,11 @@ tu6_emit_descriptor_sets(struct tu_cmd_buffer *cmd, + cs = &state_cs; + } else { + assert(bind_point == VK_PIPELINE_BIND_POINT_COMPUTE); +- +- sp_bindless_base_reg = __SP_CS_BINDLESS_BASE_DESCRIPTOR(0, {}).reg; +- hlsq_bindless_base_reg = REG_A6XX_HLSQ_CS_BINDLESS_BASE(0); +- + cs = &cmd->cs; + } + +- if (descriptors_state->max_sets_bound > 0) { +- tu_cs_emit_pkt4(cs, sp_bindless_base_reg, 2 * descriptors_state->max_sets_bound); +- tu_cs_emit_array(cs, (const uint32_t*)descriptors_state->set_iova, 2 * descriptors_state->max_sets_bound); +- if (CHIP == A6XX) { +- tu_cs_emit_pkt4(cs, hlsq_bindless_base_reg, 2 * descriptors_state->max_sets_bound); +- tu_cs_emit_array(cs, (const uint32_t*)descriptors_state->set_iova, 2 * descriptors_state->max_sets_bound); +- } +- } +- +- /* Dynamic descriptors get the reserved descriptor set. */ +- if (descriptors_state->max_dynamic_offset_size) { +- int reserved_set_idx = cmd->device->physical_device->reserved_set_idx; +- assert(reserved_set_idx >= 0); /* reserved set must be bound */ +- +- tu_cs_emit_pkt4(cs, sp_bindless_base_reg + reserved_set_idx * 2, 2); +- tu_cs_emit_qw(cs, descriptors_state->set_iova[reserved_set_idx]); +- if (CHIP == A6XX) { +- tu_cs_emit_pkt4(cs, hlsq_bindless_base_reg + reserved_set_idx * 2, 2); +- tu_cs_emit_qw(cs, descriptors_state->set_iova[reserved_set_idx]); +- } +- } +- +- tu_cs_emit_regs(cs, SP_UPDATE_CNTL(CHIP, +- .cs_bindless = bind_point == VK_PIPELINE_BIND_POINT_COMPUTE ? CHIP == A6XX ? 0x1f : 0xff : 0, +- .gfx_bindless = bind_point == VK_PIPELINE_BIND_POINT_GRAPHICS ? CHIP == A6XX ? 0x1f : 0xff : 0, +- )); ++ tu6_emit_bindless_bases(cmd, cs, descriptors_state, ++ bind_point == VK_PIPELINE_BIND_POINT_COMPUTE); + + if (bind_point == VK_PIPELINE_BIND_POINT_GRAPHICS) { + assert(cs->cur == cs->end); /* validate draw state size */ +@@ -5633,6 +5658,8 @@ tu_CmdBindPipeline(VkCommandBuffer commandBuffer, + tu_bind_tes(cmd, pipeline->shaders[MESA_SHADER_TESS_EVAL]); + tu_bind_gs(cmd, pipeline->shaders[MESA_SHADER_GEOMETRY]); + tu_bind_fs(cmd, pipeline->shaders[MESA_SHADER_FRAGMENT]); ++ cmd->state.shaders[MESA_SHADER_TASK] = pipeline->shaders[MESA_SHADER_TASK]; ++ cmd->state.shaders[MESA_SHADER_MESH] = pipeline->shaders[MESA_SHADER_MESH]; + + /* We precompile static state and count it as dynamic, so we have to + * manually clear bitset that tells which dynamic state is set, in order to +@@ -6313,6 +6340,7 @@ tu_render_pass_state_merge(struct tu_render_pass_state *dst, + { + dst->xfb_used |= src->xfb_used; + dst->has_tess |= src->has_tess; ++ dst->has_mesh |= src->has_mesh; + dst->has_prim_generated_query_in_rp |= src->has_prim_generated_query_in_rp; + dst->has_vtx_stats_query_in_rp |= src->has_vtx_stats_query_in_rp; + dst->has_zpass_done_sample_count_write_in_rp |= src->has_zpass_done_sample_count_write_in_rp; +@@ -7787,6 +7815,15 @@ tu6_emit_per_stage_push_consts(struct tu_cs *cs, + } + } + ++static uint64_t ++tu_inline_ubo_iova(struct tu_cs *cs, const struct tu_inline_ubo *ubo, ++ const struct tu_descriptor_state *descriptors) ++{ ++ if (ubo->mesh_ring) ++ return cs->device->mesh_ring->iova; ++ return (descriptors->set_iova[ubo->base] & ~0x3f) + ubo->offset; ++} ++ + static void + tu6_emit_inline_ubo(struct tu_cs *cs, + const struct tu_const_state *const_state, +@@ -7805,7 +7842,7 @@ tu6_emit_inline_ubo(struct tu_cs *cs, + if (constlen <= ubo->const_offset_vec4) + continue; + +- uint64_t va = descriptors->set_iova[ubo->base] & ~0x3f; ++ uint64_t va = tu_inline_ubo_iova(cs, ubo, descriptors); + + tu_cs_emit_pkt7(cs, tu6_stage2opcode(type), ubo->push_address ? 7 : 3); + tu_cs_emit(cs, CP_LOAD_STATE6_0_DST_OFF(ubo->const_offset_vec4) | +@@ -7816,11 +7853,11 @@ tu6_emit_inline_ubo(struct tu_cs *cs, + if (ubo->push_address) { + tu_cs_emit(cs, 0); + tu_cs_emit(cs, 0); +- tu_cs_emit_qw(cs, va + ubo->offset); ++ tu_cs_emit_qw(cs, va); + tu_cs_emit(cs, 0); + tu_cs_emit(cs, 0); + } else { +- tu_cs_emit_qw(cs, va + ubo->offset); ++ tu_cs_emit_qw(cs, va); + } + } + } +@@ -7833,18 +7870,14 @@ tu7_emit_inline_ubo(struct tu_cs *cs, + mesa_shader_stage type, + struct tu_descriptor_state *descriptors) + { +- uint64_t addresses[7] = {0}; ++ uint64_t addresses[ARRAY_SIZE(const_state->ubos)] = {0}; + unsigned offset = const_state->inline_uniforms_ubo.idx; + + if (offset == -1) + return; + +- for (unsigned i = 0; i < const_state->num_inline_ubos; i++) { +- const struct tu_inline_ubo *ubo = &const_state->ubos[i]; +- +- uint64_t va = descriptors->set_iova[ubo->base] & ~0x3f; +- addresses[i] = va + ubo->offset; +- } ++ for (unsigned i = 0; i < const_state->num_inline_ubos; i++) ++ addresses[i] = tu_inline_ubo_iova(cs, &const_state->ubos[i], descriptors); + + /* A7XX TODO: Emit data via sub_cs instead of NOP */ + uint64_t iova = tu_cs_emit_data_nop(cs, (uint32_t *)addresses, const_state->num_inline_ubos * 2, 4); +@@ -8044,6 +8077,38 @@ tu_emit_bindless_base_addresses(struct tu_cs *cs, + } + } + ++template ++static void ++tu_emit_shared_consts(struct tu_cmd_buffer *cmd, struct tu_cs *cs, ++ const struct tu_push_constant_range *shared_consts, ++ bool compute) ++{ ++ if (shared_consts->type == IR3_PUSH_CONSTS_SHARED) { ++ tu6_emit_shared_consts(cs, shared_consts, cmd->push_constants, compute); ++ } else if (shared_consts->type == IR3_PUSH_CONSTS_SHARED_PREAMBLE) { ++ tu7_emit_shared_preamble_consts(cs, shared_consts, cmd->push_constants); ++ } ++} ++ ++/* Emits the constants of a compute shader reading the given descriptors. */ ++template ++static void ++tu_emit_cs_consts(struct tu_cmd_buffer *cmd, struct tu_cs *cs, ++ const struct tu_shader *shader, ++ struct tu_descriptor_state *descriptors) ++{ ++ tu_emit_shared_consts(cmd, cs, &shader->const_state.push_consts, true); ++ tu6_emit_per_stage_push_consts(cs, &shader->const_state, ++ shader->variant->const_state, ++ MESA_SHADER_COMPUTE, cmd->push_constants); ++ tu_emit_inline_ubo(cs, &shader->const_state, shader->variant->const_state, ++ shader->variant->constlen, MESA_SHADER_COMPUTE, ++ descriptors); ++ tu_emit_bindless_base_addresses(cs, &shader->const_state, ++ shader->variant->const_state, ++ MESA_SHADER_COMPUTE, descriptors); ++} ++ + template + static struct tu_draw_state + tu_emit_consts(struct tu_cmd_buffer *cmd, bool compute) +@@ -8061,29 +8126,13 @@ tu_emit_consts(struct tu_cmd_buffer *cmd, bool compute) + struct tu_cs cs; + tu_cs_begin_sub_stream(&cmd->sub_cs, dwords, &cs); + +- if (shared_consts->type == IR3_PUSH_CONSTS_SHARED) { +- tu6_emit_shared_consts(&cs, shared_consts, cmd->push_constants, compute); +- } else if (shared_consts->type == IR3_PUSH_CONSTS_SHARED_PREAMBLE) { +- tu7_emit_shared_preamble_consts(&cs, shared_consts, cmd->push_constants); +- } +- + if (compute) { +- tu6_emit_per_stage_push_consts( +- &cs, &cmd->state.shaders[MESA_SHADER_COMPUTE]->const_state, +- cmd->state.shaders[MESA_SHADER_COMPUTE]->variant->const_state, +- MESA_SHADER_COMPUTE, cmd->push_constants); +- tu_emit_inline_ubo( +- &cs, &cmd->state.shaders[MESA_SHADER_COMPUTE]->const_state, +- cmd->state.shaders[MESA_SHADER_COMPUTE]->variant->const_state, +- cmd->state.shaders[MESA_SHADER_COMPUTE]->variant->constlen, +- MESA_SHADER_COMPUTE, +- tu_get_descriptors_state(cmd, VK_PIPELINE_BIND_POINT_COMPUTE)); +- tu_emit_bindless_base_addresses( +- &cs, &cmd->state.shaders[MESA_SHADER_COMPUTE]->const_state, +- cmd->state.shaders[MESA_SHADER_COMPUTE]->variant->const_state, +- MESA_SHADER_COMPUTE, ++ tu_emit_cs_consts( ++ cmd, &cs, cmd->state.shaders[MESA_SHADER_COMPUTE], + tu_get_descriptors_state(cmd, VK_PIPELINE_BIND_POINT_COMPUTE)); + } else { ++ tu_emit_shared_consts(cmd, &cs, shared_consts, false); ++ + struct tu_descriptor_state *descriptors = + tu_get_descriptors_state(cmd, VK_PIPELINE_BIND_POINT_GRAPHICS); + for (uint32_t type = MESA_SHADER_VERTEX; type <= MESA_SHADER_FRAGMENT; type++) { +@@ -8807,7 +8856,7 @@ tu6_draw_common(struct tu_cmd_buffer *cmd, + MESA_VK_DYNAMIC_IA_PRIMITIVE_TOPOLOGY) || + BITSET_TEST(cmd->vk.dynamic_graphics_state.dirty, + MESA_VK_DYNAMIC_RS_LINE_MODE) || +- (cmd->state.dirty & TU_CMD_DIRTY_TES) || ++ (cmd->state.dirty & (TU_CMD_DIRTY_TES | TU_CMD_DIRTY_PROGRAM)) || + (cmd->state.dirty & TU_CMD_DIRTY_DRAW_STATE)) { + tu6_update_msaa_disable(cmd); + } +@@ -8914,8 +8963,7 @@ tu6_draw_common(struct tu_cmd_buffer *cmd, + static uint32_t + tu_draw_initiator(struct tu_cmd_buffer *cmd, enum pc_di_src_sel src_sel) + { +- enum pc_di_primtype primtype = +- tu6_primtype((VkPrimitiveTopology)cmd->vk.dynamic_graphics_state.ia.primitive_topology); ++ enum pc_di_primtype primtype = tu6_primtype(tu_primitive_topology(cmd)); + + if (primtype == DI_PT_PATCHES0) + primtype = (enum pc_di_primtype) (primtype + +@@ -9481,10 +9529,10 @@ template + static void + tu_emit_compute_driver_params(struct tu_cmd_buffer *cmd, + struct tu_cs *cs, ++ const struct tu_shader *shader, + const struct tu_dispatch_info *info) + { + mesa_shader_stage type = MESA_SHADER_COMPUTE; +- const struct tu_shader *shader = cmd->state.shaders[MESA_SHADER_COMPUTE]; + const struct ir3_shader_variant *variant = shader->variant; + const struct ir3_const_state *const_state = variant->const_state; + unsigned subgroup_size = variant->info.subgroup_size; +@@ -9678,11 +9726,14 @@ tu_dispatch(struct tu_cmd_buffer *cmd, + cmd->device->physical_device->info->props.instr_cache_size; + + /* We don't use draw states for dispatches, so the bound pipeline +- * could be overwritten by reg stomping in a renderpass or blit. ++ * could be overwritten by reg stomping in a renderpass or blit, or by the ++ * compute shaders of a mesh draw. + */ +- if (cmd->device->dbg_renderpass_stomp_cs) { ++ if (cmd->device->dbg_renderpass_stomp_cs || ++ cmd->state.compute_program_stale) { + tu_cs_emit_state_ib(&cmd->cs, shader->state); + cmd->state.dirty |= TU_CMD_DIRTY_COMPUTE_DESC_SETS; ++ cmd->state.compute_program_stale = false; + } + + /* There appears to be a HW bug where in some rare circumstances it appears +@@ -9715,7 +9766,7 @@ tu_dispatch(struct tu_cmd_buffer *cmd, + /* note: no reason to have this in a separate IB */ + tu_cs_emit_state_ib(cs, tu_emit_consts(cmd, true)); + +- tu_emit_compute_driver_params(cmd, cs, info); ++ tu_emit_compute_driver_params(cmd, cs, shader, info); + + if (cmd->state.dirty & TU_CMD_DIRTY_COMPUTE_DESC_SETS) { + tu6_emit_descriptor_sets(cmd, VK_PIPELINE_BIND_POINT_COMPUTE); +@@ -9996,6 +10047,484 @@ tu_dispatch_unaligned_indirect(VkCommandBuffer commandBuffer, + TU_CALLX(cmd_buffer->device, tu_dispatch)(cmd_buffer, &info); + } + ++/* Mesh draws run the task and mesh shaders as compute dispatches inside the ++ * render pass. A generated vertex shader then draws the primitives the mesh ++ * shader wrote to the mesh ring, one chunk of mesh workgroups at a time. ++ * Direct draws know their workgroup counts up front. Indirect and task draws ++ * build the workgroup table with the setup shader and predicate every chunk ++ * on the arguments it writes. ++ */ ++ ++struct tu_mesh_draw { ++ uint32_t groups[3]; ++ uint64_t indirect; ++ uint64_t count; ++ uint32_t draw_count; ++ uint32_t stride; ++}; ++ ++struct tu_mesh_setup { ++ uint32_t table; ++ enum tu_mesh_source source; ++ uint64_t src; ++ uint32_t stride; ++ uint32_t max_count; ++ uint32_t count; ++ uint64_t count_iova; ++ uint32_t first; ++ uint32_t chunk; ++ uint32_t vertices; ++ uint32_t chunks; ++}; ++ ++static uint64_t ++tu_mesh_ring(const struct tu_cmd_buffer *cmd, uint32_t offset) ++{ ++ return cmd->device->mesh_ring->iova + offset; ++} ++ ++static uint64_t ++tu_mesh_chunk_args(const struct tu_cmd_buffer *cmd, uint32_t table, ++ uint32_t chunk) ++{ ++ return tu_mesh_ring(cmd, table + TU_MESH_TABLE_ARGS + ++ chunk * TU_MESH_ARGS_SIZE); ++} ++ ++/* Makes compute shader writes visible to later shaders and to the CP. */ ++template ++static void ++tu_mesh_emit_shader_barrier(struct tu_cmd_buffer *cmd, struct tu_cs *cs) ++{ ++ tu_cs_emit_wfi(cs); ++ tu_emit_event_write(cmd, cs, FD_CACHE_CLEAN); ++ tu_emit_event_write(cmd, cs, FD_CACHE_INVALIDATE); ++ tu_cs_emit_wfi(cs); ++ tu_cs_emit_pkt7(cs, CP_WAIT_FOR_ME, 0); ++} ++ ++/* Makes CP memory writes visible to shaders. */ ++template ++static void ++tu_mesh_emit_cp_barrier(struct tu_cmd_buffer *cmd, struct tu_cs *cs) ++{ ++ tu_cs_emit_pkt7(cs, CP_WAIT_MEM_WRITES, 0); ++ tu_emit_event_write(cmd, cs, FD_CACHE_INVALIDATE); ++ tu_cs_emit_wfi(cs); ++} ++ ++/* Loads one of the compute shaders of a mesh draw. Dispatches can't use draw ++ * states and the draw CS can't call another IB, so the shader state is copied ++ * inline. ++ */ ++template ++static void ++tu_mesh_emit_program(struct tu_cmd_buffer *cmd, struct tu_cs *cs, ++ const struct tu_shader *shader) ++{ ++ const uint32_t *state = (const uint32_t *) ++ ((const char *) shader->bo.bo->map + ++ (shader->state.iova - shader->bo.bo->iova)); ++ tu_cs_reserve(cs, shader->state.size); ++ tu_cs_emit_array(cs, state, shader->state.size); ++ ++ struct tu_descriptor_state *descriptors = ++ tu_get_descriptors_state(cmd, VK_PIPELINE_BIND_POINT_GRAPHICS); ++ tu_emit_cs_consts(cmd, cs, shader, descriptors); ++ tu6_emit_dynamic_offset(cs, shader->variant, shader, &cmd->state.program); ++ tu6_emit_bindless_bases(cmd, cs, descriptors, true); ++} ++ ++/* Dispatches groups workgroups along x numbered from base, or the workgroups ++ * given by the arguments at indirect. ++ */ ++template ++static void ++tu_mesh_emit_dispatch(struct tu_cmd_buffer *cmd, struct tu_cs *cs, ++ const struct tu_shader *shader, uint32_t base, ++ uint32_t groups, uint64_t indirect) ++{ ++ const struct ir3_shader_variant *v = shader->variant; ++ const uint16_t *local_size = v->local_size; ++ ++ struct tu_dispatch_info info = {}; ++ info.blocks[0] = groups; ++ info.blocks[1] = 1; ++ info.blocks[2] = 1; ++ info.offsets[0] = base; ++ tu_emit_compute_driver_params(cmd, cs, shader, &info); ++ ++ /* See tu_dispatch(). */ ++ bool emit_instrlen_workaround = ++ v->instrlen > cmd->device->physical_device->info->props.instr_cache_size; ++ if (emit_instrlen_workaround) { ++ tu_cs_emit_regs(cs, A6XX_SP_PS_INSTR_SIZE(v->instrlen)); ++ tu_emit_event_write(cmd, cs, FD_LABEL); ++ } ++ ++ tu_set_render_mode(cs, {RM6_COMPUTE}); ++ ++ tu_cs_emit_regs(cs, ++ SP_CS_NDRANGE_0(CHIP, .kerneldim = 3, ++ .localsizex = local_size[0] - 1, ++ .localsizey = local_size[1] - 1, ++ .localsizez = local_size[2] - 1), ++ SP_CS_NDRANGE_1(CHIP, .globalsize_x = local_size[0] * groups), ++ SP_CS_NDRANGE_2(CHIP, .globaloff_x = 0), ++ SP_CS_NDRANGE_3(CHIP, .globalsize_y = local_size[1]), ++ SP_CS_NDRANGE_4(CHIP, .globaloff_y = 0), ++ SP_CS_NDRANGE_5(CHIP, .globalsize_z = local_size[2]), ++ SP_CS_NDRANGE_6(CHIP, .globaloff_z = 0)); ++ if (CHIP >= A7XX) { ++ tu_cs_emit_regs(cs, ++ SP_CS_NDRANGE_7(CHIP, .localsizex = local_size[0] - 1, ++ .localsizey = local_size[1] - 1, ++ .localsizez = local_size[2] - 1)); ++ } ++ ++ if (cmd->device->physical_device->info->props.has_rt_workaround && ++ v->info.uses_ray_intersection) { ++ tu_set_render_mode(cs, { .shader_uses_rt = true }); ++ } ++ ++ if (indirect) { ++ tu_cs_emit_pkt7(cs, CP_EXEC_CS_INDIRECT, 4); ++ tu_cs_emit(cs, 0x00000000); ++ tu_cs_emit_qw(cs, indirect); ++ tu_cs_emit(cs, ++ A5XX_CP_EXEC_CS_INDIRECT_3_LOCALSIZEX(local_size[0] - 1) | ++ A5XX_CP_EXEC_CS_INDIRECT_3_LOCALSIZEY(local_size[1] - 1) | ++ A5XX_CP_EXEC_CS_INDIRECT_3_LOCALSIZEZ(local_size[2] - 1)); ++ } else { ++ tu_cs_emit_pkt7(cs, CP_EXEC_CS, 4); ++ tu_cs_emit(cs, 0x00000000); ++ tu_cs_emit(cs, CP_EXEC_CS_1_NGROUPS_X(groups)); ++ tu_cs_emit(cs, CP_EXEC_CS_2_NGROUPS_Y(1)); ++ tu_cs_emit(cs, CP_EXEC_CS_3_NGROUPS_Z(1)); ++ } ++ ++ if (emit_instrlen_workaround) ++ tu_emit_event_write(cmd, cs, FD_LABEL); ++ ++ tu_set_render_mode(cs, {RM6_DIRECT_RENDER}); ++} ++ ++/* Executes the following packets only if the chunk arguments at args are ++ * enabled. ++ */ ++template ++static void ++tu_mesh_begin_chunk_cond(struct tu_cs *cs, uint64_t args, ++ enum tu_predicate_bit bit) ++{ ++ struct fd_reg_pair scratch = tu_scratch_reg(0); ++ cs->mem_to_reg(scratch, args + TU_MESH_ARG_ENABLE); ++ ++ tu_cs_emit_pkt7(cs, CP_REG_TEST, 1); ++ tu_cs_emit(cs, A6XX_CP_REG_TEST_0_REG(scratch.reg) | ++ A6XX_CP_REG_TEST_0_BIT(0) | ++ A6XX_CP_REG_TEST_0_PRED_BIT(bit)); ++ ++ tu_cond_exec_start(cs, CP_COND_REG_EXEC_0_MODE(PRED_TEST) | ++ CP_COND_REG_EXEC_0_PRED_BIT(bit)); ++} ++ ++/* Writes a workgroup table holding a single direct launch. */ ++template ++static void ++tu_mesh_emit_direct_table(struct tu_cmd_buffer *cmd, struct tu_cs *cs, ++ uint32_t table, const uint32_t groups[3]) ++{ ++ tu_cs_emit_pkt7(cs, CP_MEM_WRITE, 3); ++ tu_cs_emit_qw(cs, tu_mesh_ring(cmd, table)); ++ tu_cs_emit(cs, 1); ++ ++ tu_cs_emit_pkt7(cs, CP_MEM_WRITE, 2 + TU_MESH_TABLE_ENTRY_SIZE / 4); ++ tu_cs_emit_qw(cs, tu_mesh_ring(cmd, table + TU_MESH_TABLE_ENTRIES)); ++ tu_cs_emit(cs, 0); ++ tu_cs_emit(cs, groups[0]); ++ tu_cs_emit(cs, groups[1]); ++ tu_cs_emit(cs, groups[2]); ++ for (unsigned i = 4; i < TU_MESH_TABLE_ENTRY_SIZE / 4; i++) ++ tu_cs_emit(cs, 0); ++ ++ tu_mesh_emit_cp_barrier(cmd, cs); ++} ++ ++/* Builds a workgroup table and its chunk arguments with the setup shader. */ ++template ++static void ++tu_mesh_emit_setup(struct tu_cmd_buffer *cmd, struct tu_cs *cs, ++ const struct tu_mesh_setup *setup) ++{ ++ uint32_t params[TU_MESH_PARAM_NUM]; ++ params[TU_MESH_PARAM_TABLE] = setup->table; ++ params[TU_MESH_PARAM_SOURCE] = setup->source; ++ params[TU_MESH_PARAM_SRC_LO] = setup->src; ++ params[TU_MESH_PARAM_SRC_HI] = setup->src >> 32; ++ params[TU_MESH_PARAM_STRIDE] = setup->stride; ++ params[TU_MESH_PARAM_MAX_COUNT] = setup->max_count; ++ params[TU_MESH_PARAM_COUNT] = setup->count; ++ params[TU_MESH_PARAM_FIRST] = setup->first; ++ params[TU_MESH_PARAM_CHUNK] = setup->chunk; ++ params[TU_MESH_PARAM_VERTICES] = setup->vertices; ++ params[TU_MESH_PARAM_CHUNKS] = setup->chunks; ++ ++ tu_cs_emit_pkt7(cs, CP_MEM_WRITE, 2 + TU_MESH_PARAM_NUM); ++ tu_cs_emit_qw(cs, tu_mesh_ring(cmd, TU_MESH_PARAMS_OFFSET)); ++ tu_cs_emit_array(cs, params, TU_MESH_PARAM_NUM); ++ ++ if (setup->count_iova) { ++ tu_cs_emit_pkt7(cs, CP_WAIT_MEM_WRITES, 0); ++ tu_cs_emit_pkt7(cs, CP_MEM_TO_MEM, 5); ++ tu_cs_emit(cs, 0); ++ tu_cs_emit_qw(cs, tu_mesh_ring(cmd, TU_MESH_PARAMS_OFFSET + ++ TU_MESH_PARAM_COUNT * 4)); ++ tu_cs_emit_qw(cs, setup->count_iova); ++ } ++ ++ tu_mesh_emit_cp_barrier(cmd, cs); ++ ++ tu_mesh_emit_program(cmd, cs, cmd->device->mesh_setup); ++ tu_mesh_emit_dispatch(cmd, cs, cmd->device->mesh_setup, 0, 1, 0); ++ tu_mesh_emit_shader_barrier(cmd, cs); ++} ++ ++/* Runs the mesh shader over one chunk of the mesh table and draws its ++ * primitives. ++ */ ++template ++static void ++tu_mesh_emit_chunk(struct tu_cmd_buffer *cmd, struct tu_cs *cs, ++ const struct tu_shader *ms, uint32_t base, ++ uint32_t groups, uint64_t args) ++{ ++ tu_mesh_emit_dispatch(cmd, cs, ms, base, groups, args); ++ tu_mesh_emit_shader_barrier(cmd, cs); ++ ++ /* The compute state overwrote part of the fragment shader state. */ ++ tu_cs_emit_pkt7(cs, CP_SET_DRAW_STATE, 3); ++ tu_cs_emit_draw_state(cs, TU_DRAW_STATE_FS, cmd->state.program.fs_state); ++ ++ if (args) { ++ tu_cs_emit_pkt7(cs, CP_DRAW_INDIRECT_MULTI, 6); ++ tu_cs_emit(cs, tu_draw_initiator(cmd, DI_SRC_SEL_AUTO_INDEX)); ++ tu_cs_emit(cs, A6XX_CP_DRAW_INDIRECT_MULTI_1_OPCODE(INDIRECT_OP_NORMAL) | ++ A6XX_CP_DRAW_INDIRECT_MULTI_1_DST_OFF(vs_params_offset(cmd))); ++ tu_cs_emit(cs, 1); ++ tu_cs_emit_qw(cs, args + TU_MESH_ARG_DRAW); ++ tu_cs_emit(cs, 0); ++ } else { ++ tu_cs_emit_pkt7(cs, CP_DRAW_INDX_OFFSET, 3); ++ tu_cs_emit(cs, tu_draw_initiator(cmd, DI_SRC_SEL_AUTO_INDEX)); ++ tu_cs_emit(cs, 1); ++ tu_cs_emit(cs, groups * ms->mesh.max_primitives * ms->mesh.verts_per_prim); ++ } ++ ++ /* The next chunk reuses the records this draw reads. */ ++ tu_cs_emit_wfi(cs); ++} ++ ++/* Runs the mesh shader over the mesh table: count workgroups when known on ++ * the CPU, otherwise the chunks enabled by the setup shader. ++ */ ++template ++static void ++tu_mesh_emit_ms_chunks(struct tu_cmd_buffer *cmd, struct tu_cs *cs, ++ bool gpu_driven, uint32_t count) ++{ ++ const struct tu_shader *ms = cmd->state.shaders[MESA_SHADER_MESH]; ++ uint32_t chunk = ms->mesh.chunk_workgroups; ++ ++ tu_mesh_emit_program(cmd, cs, ms); ++ ++ if (!gpu_driven) { ++ for (uint32_t base = 0; base < count; base += chunk) { ++ tu_mesh_emit_chunk(cmd, cs, ms, base, MIN2(chunk, count - base), ++ 0); ++ } ++ return; ++ } ++ ++ for (uint32_t i = 0; i < TU_MESH_MAX_CHUNKS; i++) { ++ uint64_t args = tu_mesh_chunk_args(cmd, TU_MESH_MS_TABLE_OFFSET, i); ++ tu_mesh_begin_chunk_cond(cs, args, TU_PREDICATE_MESH); ++ tu_mesh_emit_chunk(cmd, cs, ms, i * chunk, 0, args); ++ tu_cond_exec_end(cs); ++ } ++} ++ ++/* Runs one chunk of the task table, then the mesh workgroups it launched. */ ++template ++static void ++tu_mesh_emit_task_chunk(struct tu_cmd_buffer *cmd, struct tu_cs *cs, ++ uint32_t base, uint32_t groups, uint64_t args) ++{ ++ const struct tu_shader *ts = cmd->state.shaders[MESA_SHADER_TASK]; ++ const struct tu_shader *ms = cmd->state.shaders[MESA_SHADER_MESH]; ++ ++ tu_mesh_emit_program(cmd, cs, ts); ++ tu_mesh_emit_dispatch(cmd, cs, ts, base, groups, args); ++ tu_mesh_emit_shader_barrier(cmd, cs); ++ ++ struct tu_mesh_setup setup = { ++ .table = TU_MESH_MS_TABLE_OFFSET, ++ .source = TU_MESH_SOURCE_TASK, ++ .stride = ts->mesh.task_payload_stride, ++ .max_count = ts->mesh.chunk_workgroups, ++ .count = groups, ++ .count_iova = args, ++ .chunk = ms->mesh.chunk_workgroups, ++ .vertices = (uint32_t) ms->mesh.max_primitives * ms->mesh.verts_per_prim, ++ .chunks = TU_MESH_MAX_CHUNKS, ++ }; ++ tu_mesh_emit_setup(cmd, cs, &setup); ++ tu_mesh_emit_ms_chunks(cmd, cs, true, 0); ++} ++ ++template ++static void ++tu_mesh_draw(struct tu_cmd_buffer *cmd, const struct tu_mesh_draw *draw) ++{ ++ struct tu_cs *cs = &cmd->draw_cs; ++ const struct tu_shader *ts = cmd->state.shaders[MESA_SHADER_TASK]; ++ const struct tu_shader *ms = cmd->state.shaders[MESA_SHADER_MESH]; ++ ++ if (draw->indirect) ++ tu6_emit_empty_vs_params(cmd); ++ else ++ tu6_emit_vs_params(cmd, 0, 0, 0); ++ ++ tu6_draw_common(cmd, cs, false, 0); ++ ++ cmd->state.rp.has_mesh = true; ++ cmd->state.compute_program_stale = true; ++ cmd->state.dirty |= TU_CMD_DIRTY_COMPUTE_DESC_SETS; ++ ++ if (!draw->indirect) { ++ uint32_t count = draw->groups[0] * draw->groups[1] * draw->groups[2]; ++ tu_mesh_emit_direct_table( ++ cmd, cs, ts ? TU_MESH_TS_TABLE_OFFSET : TU_MESH_MS_TABLE_OFFSET, ++ draw->groups); ++ ++ if (ts) { ++ uint32_t chunk = ts->mesh.chunk_workgroups; ++ for (uint32_t base = 0; base < count; base += chunk) { ++ tu_mesh_emit_task_chunk(cmd, cs, base, ++ MIN2(chunk, count - base), 0); ++ } ++ } else { ++ tu_mesh_emit_ms_chunks(cmd, cs, false, count); ++ } ++ } else { ++ for (uint32_t first = 0; first < draw->draw_count; ++ first += TU_MESH_TABLE_MAX_ENTRIES) { ++ struct tu_mesh_setup setup = { ++ .table = ts ? TU_MESH_TS_TABLE_OFFSET : TU_MESH_MS_TABLE_OFFSET, ++ .source = TU_MESH_SOURCE_INDIRECT, ++ .src = draw->indirect, ++ .stride = draw->stride, ++ .max_count = MIN2(draw->draw_count - first, ++ TU_MESH_TABLE_MAX_ENTRIES), ++ .count = draw->draw_count, ++ .count_iova = draw->count, ++ .first = first, ++ .chunk = ts ? ts->mesh.chunk_workgroups ++ : ms->mesh.chunk_workgroups, ++ .vertices = ts ? 0 : (uint32_t) ms->mesh.max_primitives * ++ ms->mesh.verts_per_prim, ++ .chunks = ts ? TU_MESH_MAX_TASK_CHUNKS : TU_MESH_MAX_CHUNKS, ++ }; ++ tu_mesh_emit_setup(cmd, cs, &setup); ++ ++ if (!ts) { ++ tu_mesh_emit_ms_chunks(cmd, cs, true, 0); ++ continue; ++ } ++ ++ for (uint32_t i = 0; i < TU_MESH_MAX_TASK_CHUNKS; i++) { ++ uint64_t args = tu_mesh_chunk_args(cmd, TU_MESH_TS_TABLE_OFFSET, i); ++ tu_mesh_begin_chunk_cond(cs, args, TU_PREDICATE_MESH_TASK); ++ tu_mesh_emit_task_chunk(cmd, cs, ++ i * ts->mesh.chunk_workgroups, 0, ++ args); ++ tu_cond_exec_end(cs); ++ } ++ } ++ } ++ ++ trace_end_draw(&cmd->rp_trace, cs); ++} ++ ++template ++VKAPI_ATTR void VKAPI_CALL ++tu_CmdDrawMeshTasksEXT(VkCommandBuffer commandBuffer, ++ uint32_t groupCountX, ++ uint32_t groupCountY, ++ uint32_t groupCountZ) ++{ ++ VK_FROM_HANDLE(tu_cmd_buffer, cmd, commandBuffer); ++ ++ if (!groupCountX || !groupCountY || !groupCountZ) ++ return; ++ ++ struct tu_mesh_draw draw = { ++ .groups = { groupCountX, groupCountY, groupCountZ }, ++ }; ++ tu_mesh_draw(cmd, &draw); ++} ++TU_GENX(tu_CmdDrawMeshTasksEXT); ++ ++template ++VKAPI_ATTR void VKAPI_CALL ++tu_CmdDrawMeshTasksIndirectEXT(VkCommandBuffer commandBuffer, ++ VkBuffer _buffer, ++ VkDeviceSize offset, ++ uint32_t drawCount, ++ uint32_t stride) ++{ ++ VK_FROM_HANDLE(tu_cmd_buffer, cmd, commandBuffer); ++ VK_FROM_HANDLE(tu_buffer, buf, _buffer); ++ ++ if (!drawCount) ++ return; ++ ++ struct tu_mesh_draw draw = { ++ .indirect = vk_buffer_address(&buf->vk, offset), ++ .draw_count = drawCount, ++ .stride = stride, ++ }; ++ tu_mesh_draw(cmd, &draw); ++} ++TU_GENX(tu_CmdDrawMeshTasksIndirectEXT); ++ ++template ++VKAPI_ATTR void VKAPI_CALL ++tu_CmdDrawMeshTasksIndirectCountEXT(VkCommandBuffer commandBuffer, ++ VkBuffer _buffer, ++ VkDeviceSize offset, ++ VkBuffer countBuffer, ++ VkDeviceSize countBufferOffset, ++ uint32_t maxDrawCount, ++ uint32_t stride) ++{ ++ VK_FROM_HANDLE(tu_cmd_buffer, cmd, commandBuffer); ++ VK_FROM_HANDLE(tu_buffer, buf, _buffer); ++ VK_FROM_HANDLE(tu_buffer, count_buf, countBuffer); ++ ++ if (!maxDrawCount) ++ return; ++ ++ struct tu_mesh_draw draw = { ++ .indirect = vk_buffer_address(&buf->vk, offset), ++ .count = vk_buffer_address(&count_buf->vk, countBufferOffset), ++ .draw_count = maxDrawCount, ++ .stride = stride, ++ }; ++ tu_mesh_draw(cmd, &draw); ++} ++TU_GENX(tu_CmdDrawMeshTasksIndirectCountEXT); ++ + VKAPI_ATTR void VKAPI_CALL + tu_CmdEndRenderPass2(VkCommandBuffer commandBuffer, + const VkSubpassEndInfo *pSubpassEndInfo) +diff --git a/src/freedreno/vulkan/tu_cmd_buffer.h b/src/freedreno/vulkan/tu_cmd_buffer.h +index 1a78b9e..3372fb3 100644 +--- a/src/freedreno/vulkan/tu_cmd_buffer.h ++++ b/src/freedreno/vulkan/tu_cmd_buffer.h +@@ -312,6 +312,7 @@ struct tu_render_pass_state + { + bool xfb_used; + bool has_tess; ++ bool has_mesh; + bool has_prim_generated_query_in_rp; + bool has_vtx_stats_query_in_rp; + bool has_zpass_done_sample_count_write_in_rp; +@@ -475,7 +476,7 @@ struct tu_cmd_state + { + uint32_t dirty; + +- struct tu_shader *shaders[MESA_SHADER_STAGES]; ++ struct tu_shader *shaders[MESA_SHADER_MESH_STAGES]; + + struct tu_program_state program; + +@@ -596,6 +597,12 @@ struct tu_cmd_state + bool fdm_custom_resolve_subsampled; + + bool tessfactor_addr_set; ++ ++ /* Mesh draws run their own compute shaders, so the bound compute pipeline ++ * state has to be emitted again before the next dispatch. ++ */ ++ bool compute_program_stale; ++ + bool predication_active; + bool msaa_disable; + tu_lrz_blend_status lrz_blend_status; +diff --git a/src/freedreno/vulkan/tu_common.h b/src/freedreno/vulkan/tu_common.h +index dfa46d8..589daf6 100644 +--- a/src/freedreno/vulkan/tu_common.h ++++ b/src/freedreno/vulkan/tu_common.h +@@ -160,6 +160,8 @@ enum tu_predicate_bit { + TU_PREDICATE_SUBSAMPLED_NO_FAST_STORE = 7, + TU_PREDICATE_NON_SUBSAMPLED_FAST_STORE = 8, + TU_PREDICATE_NON_SUBSAMPLED_NO_FAST_STORE = 9, ++ TU_PREDICATE_MESH_TASK = 10, ++ TU_PREDICATE_MESH = 11, + }; + + /* Onchip timestamp register layout. */ +diff --git a/src/freedreno/vulkan/tu_device.cc b/src/freedreno/vulkan/tu_device.cc +index 180591a..ff079f7 100644 +--- a/src/freedreno/vulkan/tu_device.cc ++++ b/src/freedreno/vulkan/tu_device.cc +@@ -205,6 +205,16 @@ static bool tu_is_vk_1_1(const struct tu_physical_device *device) + return tu_has_multiview(device); + } + ++/* Mesh and task shaders are emulated with compute dispatches inside ++ * sysmem render passes. ++ */ ++static bool ++tu_has_mesh_shader(const struct tu_physical_device *device) ++{ ++ return device->info->chip >= 7 && ++ device->info->cs_shared_mem_size >= 32 * 1024; ++} ++ + static uint32_t + tu_subgroup_size(const struct tu_physical_device *device) + { +@@ -378,6 +388,7 @@ get_device_extensions(const struct tu_physical_device *device, + .EXT_load_store_op_none = true, + .EXT_map_memory_placed = true, + .EXT_memory_budget = true, ++ .EXT_mesh_shader = tu_has_mesh_shader(device), + .EXT_multi_draw = true, + .EXT_multisampled_render_to_single_sampled = true, + .EXT_mutable_descriptor_type = true, +@@ -829,6 +840,13 @@ tu_get_features(struct tu_physical_device *pdevice, + features->memoryMapRangePlaced = false; + features->memoryUnmapReserve = true; + ++ /* VK_EXT_mesh_shader */ ++ features->taskShader = tu_has_mesh_shader(pdevice); ++ features->meshShader = tu_has_mesh_shader(pdevice); ++ features->multiviewMeshShader = false; ++ features->primitiveFragmentShadingRateMeshShader = false; ++ features->meshShaderQueries = false; ++ + /* VK_EXT_multi_draw */ + features->multiDraw = true; + +@@ -993,6 +1011,10 @@ tu_get_physical_device_properties_1_1(struct tu_physical_device *pdevice, + VK_SUBGROUP_FEATURE_ROTATE_CLUSTERED_BIT_KHR | + VK_SUBGROUP_FEATURE_CLUSTERED_BIT | + VK_SUBGROUP_FEATURE_ARITHMETIC_BIT; ++ if (tu_has_mesh_shader(pdevice)) { ++ p->subgroupSupportedStages |= ++ VK_SHADER_STAGE_TASK_BIT_EXT | VK_SHADER_STAGE_MESH_BIT_EXT; ++ } + if (pdevice->info->props.has_getfiberid) { + p->subgroupSupportedStages |= VK_SHADER_STAGE_ALL_GRAPHICS; + p->subgroupSupportedOperations |= VK_SUBGROUP_FEATURE_QUAD_BIT; +@@ -1130,7 +1152,8 @@ tu_get_physical_device_properties_1_3(struct tu_physical_device *pdevice, + p->requiredSubgroupSizeStages = + p->minSubgroupSize == p->maxSubgroupSize + ? VK_SHADER_STAGE_ALL +- : (VK_SHADER_STAGE_COMPUTE_BIT | VK_SHADER_STAGE_FRAGMENT_BIT); ++ : (VK_SHADER_STAGE_COMPUTE_BIT | VK_SHADER_STAGE_FRAGMENT_BIT | ++ VK_SHADER_STAGE_TASK_BIT_EXT | VK_SHADER_STAGE_MESH_BIT_EXT); + + p->maxInlineUniformBlockSize = MAX_INLINE_UBO_RANGE; + p->maxPerStageDescriptorInlineUniformBlocks = MAX_INLINE_UBOS; +@@ -1459,6 +1482,44 @@ tu_get_properties(struct tu_physical_device *pdevice, + os_get_page_size(&os_page_size); + props->minPlacedMemoryMapAlignment = os_page_size; + ++ /* VK_EXT_mesh_shader */ ++ props->maxTaskWorkGroupTotalCount = 1u << 22; ++ props->maxTaskWorkGroupCount[0] = 65535; ++ props->maxTaskWorkGroupCount[1] = 65535; ++ props->maxTaskWorkGroupCount[2] = 65535; ++ props->maxTaskWorkGroupInvocations = 128; ++ props->maxTaskWorkGroupSize[0] = 128; ++ props->maxTaskWorkGroupSize[1] = 128; ++ props->maxTaskWorkGroupSize[2] = 128; ++ props->maxTaskPayloadSize = 16384; ++ props->maxTaskSharedMemorySize = 32768; ++ props->maxTaskPayloadAndSharedMemorySize = 32768; ++ props->maxMeshWorkGroupTotalCount = 1u << 22; ++ props->maxMeshWorkGroupCount[0] = 65535; ++ props->maxMeshWorkGroupCount[1] = 65535; ++ props->maxMeshWorkGroupCount[2] = 65535; ++ props->maxMeshWorkGroupInvocations = 128; ++ props->maxMeshWorkGroupSize[0] = 128; ++ props->maxMeshWorkGroupSize[1] = 128; ++ props->maxMeshWorkGroupSize[2] = 128; ++ props->maxMeshSharedMemorySize = 28672; ++ props->maxMeshPayloadAndSharedMemorySize = 16384 + 28672; ++ props->maxMeshOutputMemorySize = 32768; ++ props->maxMeshPayloadAndOutputMemorySize = 16384 + 32768; ++ props->maxMeshOutputComponents = 128; ++ props->maxMeshOutputVertices = 256; ++ props->maxMeshOutputPrimitives = 256; ++ props->maxMeshOutputLayers = 8; ++ props->maxMeshMultiviewViewCount = 1; ++ props->meshOutputPerVertexGranularity = 1; ++ props->meshOutputPerPrimitiveGranularity = 1; ++ props->maxPreferredTaskWorkGroupInvocations = 64; ++ props->maxPreferredMeshWorkGroupInvocations = 128; ++ props->prefersLocalInvocationVertexOutput = true; ++ props->prefersLocalInvocationPrimitiveOutput = true; ++ props->prefersCompactVertexOutput = true; ++ props->prefersCompactPrimitiveOutput = true; ++ + /* VK_EXT_multi_draw */ + props->maxMultiDrawCount = 2048; + +@@ -3371,6 +3432,11 @@ tu_DestroyDevice(VkDevice _device, const VkAllocationCallbacks *pAllocator) + + tu_destroy_empty_shaders(device); + ++ if (device->mesh_setup) ++ vk_pipeline_cache_object_unref(&device->vk, &device->mesh_setup->base); ++ if (device->mesh_ring) ++ tu_bo_finish(device, device->mesh_ring); ++ + tu_destroy_dynamic_rendering(device); + + vk_meta_device_finish(&device->vk, &device->meta); +diff --git a/src/freedreno/vulkan/tu_device.h b/src/freedreno/vulkan/tu_device.h +index a0295de..2b15ae7 100644 +--- a/src/freedreno/vulkan/tu_device.h ++++ b/src/freedreno/vulkan/tu_device.h +@@ -374,6 +374,12 @@ struct tu_device + /* Lazily allocated, protected by the device mutex. */ + struct tu_bo *tess_bo; + ++ /* Mesh shading emulation state, created with the first mesh pipeline and ++ * protected by the device mutex. ++ */ ++ struct tu_bo *mesh_ring; ++ struct tu_shader *mesh_setup; ++ + struct ir3_shader_variant *global_shader_variants[GLOBAL_SH_COUNT]; + struct ir3_shader *global_shaders[GLOBAL_SH_COUNT]; + uint64_t global_shader_va[GLOBAL_SH_COUNT]; +diff --git a/src/freedreno/vulkan/tu_mesh.cc b/src/freedreno/vulkan/tu_mesh.cc +new file mode 100644 +index 0000000..f6770d7 +--- /dev/null ++++ b/src/freedreno/vulkan/tu_mesh.cc +@@ -0,0 +1,992 @@ ++/* ++ * Copyright © 2026 MaxsTechReview ++ * SPDX-License-Identifier: MIT ++ */ ++ ++#include "tu_mesh.h" ++ ++#include "nir/nir_builder.h" ++#include "util/u_math.h" ++ ++static bool ++is_indices_or_cull(const nir_variable *var) ++{ ++ return var->data.location == VARYING_SLOT_PRIMITIVE_INDICES || ++ var->data.location == VARYING_SLOT_CULL_PRIMITIVE; ++} ++ ++/* Number of vec4 record slots a per-vertex or per-primitive element uses. */ ++static unsigned ++io_slots(const nir_variable *var, const glsl_type *elem) ++{ ++ if (var->data.compact) ++ return DIV_ROUND_UP(var->data.location_frac + glsl_get_length(elem), 4); ++ return glsl_count_vec4_slots(elem, false, true); ++} ++ ++/* Explicit layout of an output element inside a record: every slot takes 16 ++ * bytes and compact arrays pack four scalars per slot, matching how varyings ++ * are assigned to locations and components. ++ */ ++static const glsl_type * ++io_explicit_type(const glsl_type *type, bool compact) ++{ ++ if (glsl_type_is_array(type)) { ++ const glsl_type *elem = glsl_get_array_element(type); ++ unsigned stride = compact ? glsl_get_bit_size(elem) / 8 : ++ glsl_count_vec4_slots(elem, false, true) * 16; ++ return glsl_array_type(io_explicit_type(elem, compact), ++ glsl_get_length(type), stride); ++ } ++ ++ if (glsl_type_is_matrix(type)) ++ return glsl_explicit_matrix_type(type, 16, false); ++ ++ if (glsl_type_is_struct(type)) { ++ unsigned num_fields = glsl_get_length(type); ++ glsl_struct_field *fields = ++ (glsl_struct_field *) calloc(num_fields, sizeof(*fields)); ++ unsigned slot = 0; ++ for (unsigned i = 0; i < num_fields; i++) { ++ fields[i] = *glsl_get_struct_field_data(type, i); ++ fields[i].type = io_explicit_type(fields[i].type, false); ++ fields[i].offset = slot * 16; ++ slot += glsl_count_vec4_slots(glsl_get_struct_field(type, i), false, true); ++ } ++ const glsl_type *result = ++ glsl_struct_type_with_explicit_alignment(fields, num_fields, ++ glsl_get_type_name(type), ++ false, 16); ++ free(fields); ++ return result; ++ } ++ ++ return type; ++} ++ ++static void ++assign_slots(uint8_t *slot_map, unsigned *count, const nir_variable *var, ++ const glsl_type *elem) ++{ ++ unsigned slots = io_slots(var, elem); ++ for (unsigned i = 0; i < slots; i++) ++ slot_map[var->data.location + i] = 1; ++ *count = MAX2(*count, (unsigned) var->data.location + slots); ++} ++ ++static unsigned ++compact_slots(uint8_t *slot_map, unsigned max_location) ++{ ++ unsigned n = 0; ++ for (unsigned loc = 0; loc < max_location; loc++) ++ slot_map[loc] = slot_map[loc] ? n++ : UINT8_MAX; ++ for (unsigned loc = max_location; loc < VARYING_SLOT_MAX; loc++) ++ slot_map[loc] = UINT8_MAX; ++ return n; ++} ++ ++void ++tu_mesh_gather_io(const nir_shader *ms, struct tu_mesh_io *io) ++{ ++ memset(io, 0, sizeof(*io)); ++ ++ unsigned vertex_end = 0, prim_end = 0; ++ nir_foreach_shader_out_variable (var, ms) { ++ if (is_indices_or_cull(var)) ++ continue; ++ ++ assert(var->data.location < VARYING_SLOT_MAX); ++ const glsl_type *elem = glsl_get_array_element(var->type); ++ if (var->data.per_primitive) ++ assign_slots(io->prim_slot, &prim_end, var, elem); ++ else ++ assign_slots(io->vertex_slot, &vertex_end, var, elem); ++ ++ if (var->data.location == VARYING_SLOT_PRIMITIVE_ID) ++ io->writes_primitive_id = true; ++ } ++ ++ io->vertex_slots = compact_slots(io->vertex_slot, vertex_end); ++ io->prim_slots = compact_slots(io->prim_slot, prim_end); ++ ++ io->max_vertices = MAX2(ms->info.mesh.max_vertices_out, 1); ++ io->max_primitives = MAX2(ms->info.mesh.max_primitives_out, 1); ++ io->verts_per_prim = mesa_vertices_per_prim(ms->info.mesh.primitive_type); ++ ++ unsigned vertex_size = io->max_vertices * io->vertex_slots * 16; ++ io->prim_offset = vertex_size; ++ io->index_offset = ++ io->prim_offset + io->max_primitives * io->prim_slots * 16; ++ io->cull_offset = io->index_offset + ++ align(io->max_primitives * io->verts_per_prim * 4, 16); ++ io->stride = io->cull_offset + align(io->max_primitives * 4, 16); ++ io->chunk_workgroups = MIN2(TU_MESH_RECORD_SIZE / io->stride, 65535); ++} ++ ++static nir_def * ++addr_add(nir_builder *b, nir_def *addr, nir_def *offset) ++{ ++ return nir_iadd(b, addr, nir_u2u64(b, offset)); ++} ++ ++static nir_def * ++load_ring(nir_builder *b) ++{ ++ return nir_pack_64_2x32(b, nir_load_mesh_ring_ir3(b)); ++} ++ ++static unsigned ++var_record_offset(const struct tu_mesh_io *io, const nir_variable *var) ++{ ++ const uint8_t *slot_map = ++ var->data.per_primitive ? io->prim_slot : io->vertex_slot; ++ return slot_map[var->data.location] * 16 + var->data.location_frac * 4; ++} ++ ++/* Rebuilds every deref chain rooted at old onto new_parent. */ ++static void ++rewrite_deref_chain(nir_builder *b, nir_deref_instr *old, ++ nir_deref_instr *new_parent) ++{ ++ nir_foreach_use_safe (src, &old->def) { ++ nir_instr *user = nir_src_use_instr(src); ++ if (user->type == nir_instr_type_deref) { ++ nir_deref_instr *child = nir_instr_as_deref(user); ++ b->cursor = nir_before_instr(user); ++ rewrite_deref_chain(b, child, ++ nir_build_deref_follower(b, new_parent, child)); ++ } else { ++ nir_src_rewrite(src, &new_parent->def); ++ } ++ } ++} ++ ++/* Turns all accesses to var into global memory accesses at addr. */ ++static void ++retarget_var(nir_function_impl *impl, nir_variable *var, nir_def *addr, ++ const glsl_type *type, unsigned align_offset) ++{ ++ nir_builder b = nir_builder_create(impl); ++ ++ nir_foreach_block (block, impl) { ++ nir_foreach_instr_safe (instr, block) { ++ if (instr->type != nir_instr_type_deref) ++ continue; ++ ++ nir_deref_instr *deref = nir_instr_as_deref(instr); ++ if (deref->deref_type != nir_deref_type_var || deref->var != var) ++ continue; ++ ++ b.cursor = nir_before_instr(instr); ++ nir_deref_instr *cast = ++ nir_build_deref_cast_with_alignment(&b, addr, nir_var_mem_global, ++ type, 0, 16, align_offset); ++ rewrite_deref_chain(&b, deref, cast); ++ } ++ } ++} ++ ++static nir_def * ++record_array_addr(nir_builder *b, nir_def *rec, const struct tu_mesh_io *io, ++ const nir_variable *var) ++{ ++ unsigned base = var->data.per_primitive ? io->prim_offset : 0; ++ return addr_add(b, rec, nir_imm_int(b, base + var_record_offset(io, var))); ++} ++ ++static const glsl_type * ++record_array_type(const struct tu_mesh_io *io, const nir_variable *var) ++{ ++ unsigned stride = ++ (var->data.per_primitive ? io->prim_slots : io->vertex_slots) * 16; ++ return glsl_array_type(io_explicit_type(glsl_get_array_element(var->type), ++ var->data.compact), ++ glsl_get_length(var->type), stride); ++} ++ ++/* Copies an output element out of its record, one vector at a time. */ ++static void ++copy_output(nir_builder *b, nir_deref_instr *dst, nir_deref_instr *src) ++{ ++ if (glsl_type_is_vector_or_scalar(dst->type)) { ++ nir_store_deref(b, dst, nir_load_deref(b, src), ~0); ++ } else if (glsl_type_is_struct(dst->type)) { ++ for (unsigned i = 0; i < glsl_get_length(dst->type); i++) { ++ copy_output(b, nir_build_deref_struct(b, dst, i), ++ nir_build_deref_struct(b, src, i)); ++ } ++ } else { ++ for (unsigned i = 0; i < glsl_get_length(dst->type); i++) { ++ copy_output(b, nir_build_deref_array_imm(b, dst, i), ++ nir_build_deref_array_imm(b, src, i)); ++ } ++ } ++} ++ ++nir_shader * ++tu_mesh_build_vs(const nir_shader *ms, const struct tu_mesh_io *io, ++ const nir_shader_compiler_options *options) ++{ ++ nir_builder _b = ++ nir_builder_init_simple_shader(MESA_SHADER_VERTEX, options, "tu_mesh_vs"); ++ nir_builder *b = &_b; ++ nir_shader *vs = b->shader; ++ ++ vs->info.clip_distance_array_size = ms->info.clip_distance_array_size; ++ vs->info.cull_distance_array_size = ms->info.cull_distance_array_size; ++ ++ unsigned vpp = io->verts_per_prim; ++ nir_def *vertex = nir_load_vertex_id(b); ++ nir_def *prim = nir_udiv_imm(b, vertex, vpp); ++ nir_def *corner = nir_isub(b, vertex, nir_imul_imm(b, prim, vpp)); ++ nir_def *wg = nir_udiv_imm(b, prim, io->max_primitives); ++ nir_def *p = nir_isub(b, prim, nir_imul_imm(b, wg, io->max_primitives)); ++ ++ nir_def *rec = ++ addr_add(b, load_ring(b), ++ nir_iadd_imm(b, nir_imul_imm(b, wg, io->stride), ++ TU_MESH_RECORD_OFFSET)); ++ ++ nir_def *index_offset = ++ nir_iadd_imm(b, nir_imul_imm(b, nir_iadd(b, nir_imul_imm(b, p, vpp), ++ corner), 4), ++ io->index_offset); ++ nir_def *index = nir_load_global(b, 1, 32, addr_add(b, rec, index_offset), ++ .align_mul = 4); ++ nir_def *dead = nir_ieq_imm(b, index, TU_MESH_DEAD_INDEX); ++ index = nir_bcsel(b, dead, nir_imm_int(b, 0), index); ++ ++ nir_def *vertex_rec = ++ addr_add(b, rec, nir_imul_imm(b, index, io->vertex_slots * 16)); ++ nir_def *prim_rec = addr_add(b, rec, nir_imul_imm(b, p, io->prim_slots * 16)); ++ ++ bool writes_pos = false; ++ nir_foreach_shader_out_variable (ms_var, ms) { ++ if (is_indices_or_cull(ms_var)) ++ continue; ++ ++ nir_variable *var = nir_variable_clone(ms_var, vs); ++ var->type = glsl_get_array_element(ms_var->type); ++ if (var->data.per_primitive) { ++ var->data.per_primitive = false; ++ var->data.interpolation = INTERP_MODE_FLAT; ++ } ++ if (var->data.location == VARYING_SLOT_PRIMITIVE_ID) ++ var->data.location = TU_MESH_PRIMITIVE_ID_SLOT; ++ nir_shader_add_variable(vs, var); ++ ++ nir_def *base = ms_var->data.per_primitive ? prim_rec : vertex_rec; ++ unsigned offset = var_record_offset(io, ms_var) + ++ (ms_var->data.per_primitive ? io->prim_offset : 0); ++ nir_deref_instr *src = ++ nir_build_deref_cast_with_alignment( ++ b, addr_add(b, base, nir_imm_int(b, offset)), nir_var_mem_global, ++ io_explicit_type(var->type, var->data.compact), 0, 16, offset % 16); ++ nir_deref_instr *dst = nir_build_deref_var(b, var); ++ ++ if (var->data.location == VARYING_SLOT_POS) { ++ nir_def *pos = nir_load_deref(b, src); ++ nir_store_deref(b, dst, ++ nir_bcsel(b, dead, nir_imm_vec4(b, 2.0, 2.0, 2.0, 1.0), ++ pos), 0xf); ++ writes_pos = true; ++ } else { ++ copy_output(b, dst, src); ++ } ++ } ++ ++ if (!writes_pos) { ++ nir_variable *pos = nir_variable_create(vs, nir_var_shader_out, ++ glsl_vec4_type(), "gl_Position"); ++ pos->data.location = VARYING_SLOT_POS; ++ nir_store_var(b, pos, nir_imm_vec4(b, 2.0, 2.0, 2.0, 1.0), 0xf); ++ } ++ ++ nir_shader_gather_info(vs, nir_shader_get_entrypoint(vs)); ++ return vs; ++} ++ ++struct table_sysvals { ++ nir_intrinsic_instr *hw_workgroup_id; ++ nir_intrinsic_instr *hw_base_workgroup_id; ++ nir_def *l; ++ nir_def *workgroup_id; ++ nir_def *num_workgroups; ++ nir_def *draw_id; ++ nir_def *task_slot; ++}; ++ ++struct lower_sysvals_state { ++ struct table_sysvals table; ++ nir_variable *counts; ++ nir_def *task_header; ++}; ++ ++static bool ++lower_sysval_intrinsic(nir_builder *b, nir_intrinsic_instr *intr, void *data) ++{ ++ struct lower_sysvals_state *state = (struct lower_sysvals_state *) data; ++ struct table_sysvals *s = &state->table; ++ nir_def *value; ++ ++ if (intr == s->hw_workgroup_id || intr == s->hw_base_workgroup_id) ++ return false; ++ ++ b->cursor = nir_before_instr(&intr->instr); ++ ++ switch (intr->intrinsic) { ++ case nir_intrinsic_load_workgroup_id: ++ value = s->workgroup_id; ++ break; ++ case nir_intrinsic_load_base_workgroup_id: ++ value = nir_imm_zero(b, 3, intr->def.bit_size); ++ break; ++ case nir_intrinsic_load_num_workgroups: ++ value = s->num_workgroups; ++ break; ++ case nir_intrinsic_load_draw_id: ++ value = s->draw_id; ++ break; ++ case nir_intrinsic_load_global_invocation_id: { ++ const uint16_t *size = b->shader->info.workgroup_size; ++ value = nir_iadd(b, nir_imul(b, s->workgroup_id, ++ nir_imm_ivec3(b, size[0], size[1], size[2])), ++ nir_load_local_invocation_id(b)); ++ break; ++ } ++ case nir_intrinsic_set_vertex_and_primitive_count: { ++ nir_deref_instr *counts = nir_build_deref_var(b, state->counts); ++ nir_store_deref(b, nir_build_deref_array_imm(b, counts, 0), ++ intr->src[0].ssa, 0x1); ++ nir_store_deref(b, nir_build_deref_array_imm(b, counts, 1), ++ intr->src[1].ssa, 0x1); ++ nir_instr_remove(&intr->instr); ++ return true; ++ } ++ case nir_intrinsic_launch_mesh_workgroups: { ++ nir_push_if(b, nir_ieq_imm(b, nir_load_local_invocation_index(b), 0)); ++ nir_store_global(b, nir_vec4(b, nir_channel(b, intr->src[0].ssa, 0), ++ nir_channel(b, intr->src[0].ssa, 1), ++ nir_channel(b, intr->src[0].ssa, 2), ++ s->draw_id), ++ state->task_header, .align_mul = 16); ++ nir_pop_if(b, NULL); ++ nir_instr_remove(&intr->instr); ++ return true; ++ } ++ case nir_intrinsic_barrier: { ++ const nir_variable_mode ring_modes = ++ nir_var_shader_out | nir_var_mem_task_payload; ++ nir_variable_mode modes = nir_intrinsic_memory_modes(intr); ++ if (!(modes & ring_modes)) ++ return false; ++ nir_intrinsic_set_memory_modes( ++ intr, (nir_variable_mode)((modes & ~ring_modes) | nir_var_mem_global)); ++ return true; ++ } ++ default: ++ return false; ++ } ++ ++ nir_def_replace(&intr->def, nir_trim_vector(b, value, intr->def.num_components)); ++ return true; ++} ++ ++/* Finds the table entry holding linear workgroup l: the last entry whose ++ * first workgroup is not above l. ++ */ ++static nir_def * ++find_table_entry(nir_builder *b, nir_def *ring, unsigned table_offset, ++ nir_def *l) ++{ ++ nir_def *table = addr_add(b, ring, nir_imm_int(b, table_offset)); ++ nir_def *count = nir_load_global(b, 1, 32, table, .align_mul = 16); ++ ++ nir_variable *lo = nir_local_variable_create(b->impl, glsl_uint_type(), "lo"); ++ nir_variable *hi = nir_local_variable_create(b->impl, glsl_uint_type(), "hi"); ++ nir_store_var(b, lo, nir_imm_int(b, 0), 0x1); ++ nir_store_var(b, hi, nir_iadd_imm(b, nir_umax(b, count, nir_imm_int(b, 1)), -1), 0x1); ++ ++ nir_loop *loop = nir_push_loop(b); ++ { ++ nir_def *lo_v = nir_load_var(b, lo); ++ nir_def *hi_v = nir_load_var(b, hi); ++ nir_break_if(b, nir_uge(b, lo_v, hi_v)); ++ ++ nir_def *mid = nir_ushr_imm(b, nir_iadd_imm(b, nir_iadd(b, lo_v, hi_v), 1), 1); ++ nir_def *start = nir_load_global( ++ b, 1, 32, ++ addr_add(b, table, ++ nir_iadd_imm(b, nir_imul_imm(b, mid, TU_MESH_TABLE_ENTRY_SIZE), ++ TU_MESH_TABLE_ENTRIES)), ++ .align_mul = 16); ++ nir_def *take = nir_uge(b, l, start); ++ nir_store_var(b, lo, nir_bcsel(b, take, mid, lo_v), 0x1); ++ nir_store_var(b, hi, nir_bcsel(b, take, hi_v, nir_iadd_imm(b, mid, -1)), 0x1); ++ } ++ nir_pop_loop(b, loop); ++ ++ return addr_add(b, table, ++ nir_iadd_imm(b, nir_imul_imm(b, nir_load_var(b, lo), ++ TU_MESH_TABLE_ENTRY_SIZE), ++ TU_MESH_TABLE_ENTRIES)); ++} ++ ++/* Decodes the API workgroup of this dispatch from the table. The chunk's ++ * first linear workgroup is passed as the base workgroup. ++ */ ++static void ++load_table_sysvals(nir_builder *b, nir_def *ring, unsigned table_offset, ++ struct table_sysvals *s) ++{ ++ nir_def *hw_id = nir_load_workgroup_id(b); ++ nir_def *hw_base = nir_load_base_workgroup_id(b, 32); ++ s->hw_workgroup_id = nir_def_as_intrinsic(hw_id); ++ s->hw_base_workgroup_id = nir_def_as_intrinsic(hw_base); ++ s->l = nir_iadd(b, nir_channel(b, hw_id, 0), nir_channel(b, hw_base, 0)); ++ ++ nir_def *entry = find_table_entry(b, ring, table_offset, s->l); ++ nir_def *e0 = nir_load_global(b, 4, 32, entry, .align_mul = 16); ++ nir_def *e1 = nir_load_global(b, 4, 32, addr_add(b, entry, nir_imm_int(b, 16)), ++ .align_mul = 16); ++ ++ nir_def *gx = nir_channel(b, e0, 1), *gy = nir_channel(b, e0, 2); ++ nir_def *local = nir_isub(b, s->l, nir_channel(b, e0, 0)); ++ nir_def *row = nir_udiv(b, local, gx); ++ ++ s->workgroup_id = nir_vec3(b, nir_umod(b, local, gx), nir_umod(b, row, gy), ++ nir_udiv(b, row, gy)); ++ s->num_workgroups = nir_channels(b, e0, 0xe); ++ s->draw_id = nir_channel(b, e1, 0); ++ s->task_slot = nir_channel(b, e1, 1); ++} ++ ++static void ++emit_workgroup_loop(nir_builder *b, unsigned count, unsigned wg_size, ++ void (*body)(nir_builder *, nir_def *, void *), void *data) ++{ ++ nir_variable *i = nir_local_variable_create(b->impl, glsl_uint_type(), "i"); ++ nir_store_var(b, i, nir_load_local_invocation_index(b), 0x1); ++ ++ nir_loop *loop = nir_push_loop(b); ++ { ++ nir_def *i_v = nir_load_var(b, i); ++ nir_break_if(b, nir_uge_imm(b, i_v, count)); ++ body(b, i_v, data); ++ nir_store_var(b, i, nir_iadd_imm(b, i_v, wg_size), 0x1); ++ } ++ nir_pop_loop(b, loop); ++} ++ ++static void ++emit_workgroup_barrier(nir_builder *b) ++{ ++ nir_barrier(b, .execution_scope = SCOPE_WORKGROUP, ++ .memory_scope = SCOPE_WORKGROUP, ++ .memory_semantics = NIR_MEMORY_ACQ_REL, ++ .memory_modes = (nir_variable_mode)(nir_var_mem_shared | ++ nir_var_mem_global)); ++} ++ ++struct ms_epilogue { ++ const struct tu_mesh_io *io; ++ nir_def *rec; ++ nir_def *prim_count; ++ bool has_cull; ++}; ++ ++static void ++clear_cull(nir_builder *b, nir_def *p, void *data) ++{ ++ struct ms_epilogue *e = (struct ms_epilogue *) data; ++ nir_store_global(b, nir_imm_int(b, 0), ++ addr_add(b, e->rec, nir_iadd_imm(b, nir_imul_imm(b, p, 4), ++ e->io->cull_offset)), ++ .write_mask = 0x1, .align_mul = 4); ++} ++ ++static void ++write_indices(nir_builder *b, nir_def *p, void *data) ++{ ++ struct ms_epilogue *e = (struct ms_epilogue *) data; ++ const struct tu_mesh_io *io = e->io; ++ ++ nir_def *live = nir_ult(b, p, e->prim_count); ++ if (e->has_cull) { ++ nir_def *cull = nir_load_global( ++ b, 1, 32, ++ addr_add(b, e->rec, nir_iadd_imm(b, nir_imul_imm(b, p, 4), ++ io->cull_offset)), ++ .align_mul = 4); ++ live = nir_iand(b, live, nir_ieq_imm(b, cull, 0)); ++ } ++ ++ nir_def *addr = ++ addr_add(b, e->rec, ++ nir_iadd_imm(b, nir_imul_imm(b, p, io->verts_per_prim * 4), ++ io->index_offset)); ++ nir_def *index = nir_load_global(b, io->verts_per_prim, 32, addr, ++ .align_mul = 4); ++ index = nir_umin(b, index, nir_imm_int(b, io->max_vertices - 1)); ++ index = nir_bcsel(b, live, index, nir_imm_int(b, TU_MESH_DEAD_INDEX)); ++ nir_store_global(b, index, addr, ++ .write_mask = nir_component_mask(io->verts_per_prim), ++ .align_mul = 4); ++} ++ ++static nir_def * ++task_slot_addr(nir_builder *b, nir_def *ring, nir_def *slot, ++ unsigned payload_stride) ++{ ++ return addr_add(b, ring, ++ nir_iadd_imm(b, nir_imul_imm(b, slot, payload_stride), ++ TU_MESH_TASK_OFFSET)); ++} ++ ++static nir_def * ++task_payload_addr(nir_builder *b, nir_def *ring, nir_def *slot, ++ unsigned payload_stride) ++{ ++ return addr_add(b, task_slot_addr(b, ring, slot, payload_stride), ++ nir_imm_int(b, TU_MESH_TASK_HEADER_SIZE)); ++} ++ ++static bool ++lower_task_payload_access(nir_builder *b, nir_intrinsic_instr *intr, ++ void *data) ++{ ++ nir_def *base = (nir_def *) data; ++ ++ b->cursor = nir_before_instr(&intr->instr); ++ ++ switch (intr->intrinsic) { ++ case nir_intrinsic_load_task_payload: { ++ nir_def *addr = addr_add(b, base, nir_iadd_imm(b, intr->src[0].ssa, ++ nir_intrinsic_base(intr))); ++ nir_def_replace(&intr->def, ++ nir_load_global(b, intr->def.num_components, ++ intr->def.bit_size, addr, ++ .align_mul = nir_intrinsic_align_mul(intr), ++ .align_offset = nir_intrinsic_align_offset(intr))); ++ return true; ++ } ++ case nir_intrinsic_store_task_payload: { ++ nir_def *addr = addr_add(b, base, nir_iadd_imm(b, intr->src[1].ssa, ++ nir_intrinsic_base(intr))); ++ nir_store_global(b, intr->src[0].ssa, addr, ++ .write_mask = nir_intrinsic_write_mask(intr), ++ .align_mul = nir_intrinsic_align_mul(intr), ++ .align_offset = nir_intrinsic_align_offset(intr)); ++ nir_instr_remove(&intr->instr); ++ return true; ++ } ++ case nir_intrinsic_task_payload_atomic: ++ case nir_intrinsic_task_payload_atomic_swap: { ++ bool swap = intr->intrinsic == nir_intrinsic_task_payload_atomic_swap; ++ nir_def *addr = addr_add(b, base, nir_iadd_imm(b, intr->src[0].ssa, ++ nir_intrinsic_base(intr))); ++ nir_def *result = ++ swap ? nir_global_atomic_swap(b, intr->def.bit_size, addr, ++ intr->src[1].ssa, intr->src[2].ssa, ++ .atomic_op = nir_intrinsic_atomic_op(intr)) ++ : nir_global_atomic(b, intr->def.bit_size, addr, intr->src[1].ssa, ++ .atomic_op = nir_intrinsic_atomic_op(intr)); ++ nir_def_replace(&intr->def, result); ++ return true; ++ } ++ default: ++ return false; ++ } ++} ++ ++static void ++lower_task_payload_vars(nir_shader *nir) ++{ ++ NIR_PASS(_, nir, nir_lower_vars_to_explicit_types, nir_var_mem_task_payload, ++ glsl_get_natural_size_align_bytes); ++ NIR_PASS(_, nir, nir_lower_explicit_io, nir_var_mem_task_payload, ++ nir_address_format_32bit_offset); ++} ++ ++static void ++become_compute(nir_shader *nir) ++{ ++ nir->info.stage = MESA_SHADER_COMPUTE; ++ nir->info.next_stage = MESA_SHADER_NONE; ++ memset(&nir->info.cs, 0, sizeof(nir->info.cs)); ++ nir->info.inputs_read = 0; ++ nir->info.outputs_written = 0; ++ nir->info.outputs_read = 0; ++ nir->info.per_primitive_outputs = 0; ++ nir->info.clip_distance_array_size = 0; ++ nir->info.cull_distance_array_size = 0; ++} ++ ++void ++tu_mesh_lower_ms(nir_shader *ms, const struct tu_mesh_io *io, ++ unsigned task_payload_stride) ++{ ++ nir_function_impl *impl = nir_shader_get_entrypoint(ms); ++ nir_builder _b = nir_builder_at(nir_before_impl(impl)); ++ nir_builder *b = &_b; ++ ++ unsigned wg_size = ms->info.workgroup_size[0] * ms->info.workgroup_size[1] * ++ ms->info.workgroup_size[2]; ++ ++ nir_def *ring = load_ring(b); ++ ++ struct lower_sysvals_state sysvals = { ++ .counts = nir_variable_create(ms, nir_var_mem_shared, ++ glsl_array_type(glsl_uint_type(), 2, 4), ++ "tu_mesh_counts"), ++ }; ++ load_table_sysvals(b, ring, TU_MESH_MS_TABLE_OFFSET, &sysvals.table); ++ nir_def *l = sysvals.table.l; ++ ++ nir_def *rec = ++ addr_add(b, ring, ++ nir_iadd_imm(b, nir_imul_imm(b, nir_umod_imm(b, l, io->chunk_workgroups), ++ io->stride), ++ TU_MESH_RECORD_OFFSET)); ++ ++ struct ms_epilogue epilogue = { .io = io, .rec = rec }; ++ ++ nir_variable *indices = NULL; ++ nir_foreach_shader_out_variable (var, ms) { ++ if (var->data.location == VARYING_SLOT_PRIMITIVE_INDICES) ++ indices = var; ++ else if (var->data.location == VARYING_SLOT_CULL_PRIMITIVE) ++ epilogue.has_cull = true; ++ } ++ ++ nir_push_if(b, nir_ieq_imm(b, nir_load_local_invocation_index(b), 0)); ++ { ++ nir_store_array_var_imm(b, sysvals.counts, 0, nir_imm_int(b, 0), 0x1); ++ nir_store_array_var_imm(b, sysvals.counts, 1, nir_imm_int(b, 0), 0x1); ++ } ++ nir_pop_if(b, NULL); ++ if (epilogue.has_cull) ++ emit_workgroup_loop(b, io->max_primitives, wg_size, clear_cull, &epilogue); ++ emit_workgroup_barrier(b); ++ ++ nir_def *payload = ++ task_payload_addr(b, ring, sysvals.table.task_slot, task_payload_stride); ++ ++ nir_shader_intrinsics_pass(ms, lower_sysval_intrinsic, nir_metadata_none, ++ &sysvals); ++ ++ nir_foreach_variable_with_modes_safe (var, ms, nir_var_shader_out) { ++ nir_builder vb = nir_builder_at(nir_after_instr(nir_def_instr(rec))); ++ nir_def *addr; ++ const glsl_type *type; ++ unsigned align_offset = 0; ++ ++ if (var == indices) { ++ addr = addr_add(&vb, rec, nir_imm_int(&vb, io->index_offset)); ++ type = glsl_array_type(glsl_get_array_element(var->type), ++ glsl_get_length(var->type), ++ io->verts_per_prim * 4); ++ } else if (var->data.location == VARYING_SLOT_CULL_PRIMITIVE) { ++ addr = addr_add(&vb, rec, nir_imm_int(&vb, io->cull_offset)); ++ type = glsl_array_type(glsl_get_array_element(var->type), ++ glsl_get_length(var->type), 4); ++ } else { ++ addr = record_array_addr(&vb, rec, io, var); ++ type = record_array_type(io, var); ++ align_offset = var_record_offset(io, var) % 16; ++ } ++ ++ retarget_var(impl, var, addr, type, align_offset); ++ } ++ ++ b->cursor = nir_after_impl(impl); ++ emit_workgroup_barrier(b); ++ epilogue.prim_count = ++ nir_umin(b, nir_load_array_var_imm(b, sysvals.counts, 1), ++ nir_imm_int(b, io->max_primitives)); ++ emit_workgroup_loop(b, io->max_primitives, wg_size, write_indices, &epilogue); ++ ++ lower_task_payload_vars(ms); ++ nir_shader_intrinsics_pass(ms, lower_task_payload_access, ++ nir_metadata_control_flow, payload); ++ ++ NIR_PASS(_, ms, nir_remove_dead_derefs); ++ NIR_PASS(_, ms, nir_remove_dead_variables, ++ nir_var_shader_out | nir_var_mem_task_payload, NULL); ++ NIR_PASS(_, ms, nir_lower_vars_to_ssa); ++ NIR_PASS(_, ms, nir_opt_dce); ++ ++ become_compute(ms); ++} ++ ++unsigned ++tu_mesh_task_chunk(unsigned task_payload_stride) ++{ ++ return MIN2(TU_MESH_TASK_SIZE / task_payload_stride, ++ TU_MESH_TABLE_MAX_ENTRIES); ++} ++ ++unsigned ++tu_mesh_lower_ts(nir_shader *ts) ++{ ++ lower_task_payload_vars(ts); ++ nir_lower_task_shader_options options = {}; ++ NIR_PASS(_, ts, nir_lower_task_shader, options); ++ ++ unsigned stride = ++ align(TU_MESH_TASK_HEADER_SIZE + ts->info.task_payload_size, 16); ++ ++ nir_function_impl *impl = nir_shader_get_entrypoint(ts); ++ nir_builder _b = nir_builder_at(nir_before_impl(impl)); ++ nir_builder *b = &_b; ++ ++ nir_def *ring = load_ring(b); ++ ++ struct lower_sysvals_state sysvals = {}; ++ load_table_sysvals(b, ring, TU_MESH_TS_TABLE_OFFSET, &sysvals.table); ++ ++ nir_def *slot = nir_umod_imm(b, sysvals.table.l, tu_mesh_task_chunk(stride)); ++ sysvals.task_header = task_slot_addr(b, ring, slot, stride); ++ nir_def *payload = task_payload_addr(b, ring, slot, stride); ++ ++ /* A workgroup that launches nothing leaves an empty header. */ ++ nir_push_if(b, nir_ieq_imm(b, nir_load_local_invocation_index(b), 0)); ++ { ++ nir_store_global(b, nir_vec4(b, nir_imm_int(b, 0), nir_imm_int(b, 0), ++ nir_imm_int(b, 0), sysvals.table.draw_id), ++ sysvals.task_header, .align_mul = 16); ++ } ++ nir_pop_if(b, NULL); ++ ++ nir_shader_intrinsics_pass(ts, lower_sysval_intrinsic, nir_metadata_none, ++ &sysvals); ++ nir_shader_intrinsics_pass(ts, lower_task_payload_access, ++ nir_metadata_control_flow, payload); ++ ++ NIR_PASS(_, ts, nir_remove_dead_variables, nir_var_mem_task_payload, NULL); ++ NIR_PASS(_, ts, nir_lower_vars_to_ssa); ++ NIR_PASS(_, ts, nir_opt_dce); ++ ++ become_compute(ts); ++ return stride; ++} ++ ++void ++tu_mesh_lower_fs_inputs(nir_shader *fs, bool remap_primitive_id) ++{ ++ nir_foreach_shader_in_variable (var, fs) { ++ if (var->data.per_primitive) { ++ var->data.per_primitive = false; ++ var->data.interpolation = INTERP_MODE_FLAT; ++ } ++ ++ if (remap_primitive_id && ++ var->data.location == VARYING_SLOT_PRIMITIVE_ID) ++ var->data.location = TU_MESH_PRIMITIVE_ID_SLOT; ++ } ++ ++ fs->info.per_primitive_inputs = 0; ++} ++ ++struct setup_state { ++ nir_def *ring; ++ nir_def *table; ++ nir_def *params[TU_MESH_PARAM_NUM]; ++ nir_def *count; ++}; ++ ++/* Returns the launch dimensions and draw id of source item i. */ ++static nir_def * ++setup_load_item(nir_builder *b, struct setup_state *s, nir_def *i) ++{ ++ nir_def *task = nir_ieq_imm(b, s->params[TU_MESH_PARAM_SOURCE], ++ TU_MESH_SOURCE_TASK); ++ nir_def *draw = nir_iadd(b, s->params[TU_MESH_PARAM_FIRST], i); ++ ++ nir_def *indirect_value, *task_value; ++ nir_push_if(b, task); ++ { ++ nir_def *addr = ++ addr_add(b, s->ring, ++ nir_iadd_imm(b, nir_imul(b, i, s->params[TU_MESH_PARAM_STRIDE]), ++ TU_MESH_TASK_OFFSET)); ++ task_value = nir_load_global(b, 4, 32, addr, .align_mul = 16); ++ } ++ nir_push_else(b, NULL); ++ { ++ nir_def *src = nir_pack_64_2x32_split(b, s->params[TU_MESH_PARAM_SRC_LO], ++ s->params[TU_MESH_PARAM_SRC_HI]); ++ nir_def *addr = ++ nir_iadd(b, src, nir_imul(b, nir_u2u64(b, draw), ++ nir_u2u64(b, s->params[TU_MESH_PARAM_STRIDE]))); ++ nir_def *dims = nir_load_global(b, 3, 32, addr, .align_mul = 4); ++ indirect_value = nir_vec4(b, nir_channel(b, dims, 0), ++ nir_channel(b, dims, 1), ++ nir_channel(b, dims, 2), draw); ++ } ++ nir_pop_if(b, NULL); ++ ++ return nir_if_phi(b, task_value, indirect_value); ++} ++ ++static nir_def * ++item_groups(nir_builder *b, nir_def *item) ++{ ++ return nir_imul(b, nir_imul(b, nir_channel(b, item, 0), ++ nir_channel(b, item, 1)), ++ nir_channel(b, item, 2)); ++} ++ ++/* Runs body(i) for every item of this invocation's contiguous range. */ ++static void ++setup_range_loop(nir_builder *b, struct setup_state *s, nir_variable *start, ++ bool write) ++{ ++ nir_def *t = nir_load_local_invocation_index(b); ++ nir_def *per_thread = ++ nir_udiv_imm(b, nir_iadd_imm(b, s->count, TU_MESH_SETUP_WORKGROUP_SIZE - 1), ++ TU_MESH_SETUP_WORKGROUP_SIZE); ++ nir_def *begin = nir_umin(b, nir_imul(b, t, per_thread), s->count); ++ nir_def *end = nir_umin(b, nir_iadd(b, begin, per_thread), s->count); ++ ++ nir_variable *i = nir_local_variable_create(b->impl, glsl_uint_type(), "i"); ++ nir_store_var(b, i, begin, 0x1); ++ ++ nir_loop *loop = nir_push_loop(b); ++ { ++ nir_def *i_v = nir_load_var(b, i); ++ nir_break_if(b, nir_uge(b, i_v, end)); ++ ++ nir_def *item = setup_load_item(b, s, i_v); ++ nir_def *groups = item_groups(b, item); ++ nir_def *start_v = nir_load_var(b, start); ++ ++ if (write) { ++ nir_def *task = nir_ieq_imm(b, s->params[TU_MESH_PARAM_SOURCE], ++ TU_MESH_SOURCE_TASK); ++ nir_def *entry = ++ addr_add(b, s->table, ++ nir_iadd_imm(b, nir_imul_imm(b, i_v, TU_MESH_TABLE_ENTRY_SIZE), ++ TU_MESH_TABLE_ENTRIES)); ++ nir_store_global(b, nir_vec4(b, start_v, nir_channel(b, item, 0), ++ nir_channel(b, item, 1), ++ nir_channel(b, item, 2)), ++ entry, .align_mul = 16); ++ nir_store_global(b, nir_vec4(b, nir_channel(b, item, 3), ++ nir_bcsel(b, task, i_v, nir_imm_int(b, 0)), ++ nir_imm_int(b, 0), nir_imm_int(b, 0)), ++ addr_add(b, entry, nir_imm_int(b, 16)), ++ .align_mul = 16); ++ } ++ ++ nir_store_var(b, start, nir_iadd(b, start_v, groups), 0x1); ++ nir_store_var(b, i, nir_iadd_imm(b, i_v, 1), 0x1); ++ } ++ nir_pop_loop(b, loop); ++} ++ ++/* Builds a table (entries sorted by first workgroup) from indirect commands ++ * or task launches, and the per-chunk dispatch and draw arguments. ++ */ ++nir_shader * ++tu_mesh_build_setup_cs(const nir_shader_compiler_options *options) ++{ ++ nir_builder _b = ++ nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, options, ++ "tu_mesh_setup"); ++ nir_builder *b = &_b; ++ b->shader->info.workgroup_size[0] = TU_MESH_SETUP_WORKGROUP_SIZE; ++ b->shader->info.workgroup_size[1] = 1; ++ b->shader->info.workgroup_size[2] = 1; ++ ++ struct setup_state s = {}; ++ s.ring = load_ring(b); ++ for (unsigned i = 0; i < TU_MESH_PARAM_NUM; i += 4) { ++ nir_def *v = nir_load_global( ++ b, MIN2(4, TU_MESH_PARAM_NUM - i), 32, ++ addr_add(b, s.ring, nir_imm_int(b, TU_MESH_PARAMS_OFFSET + i * 4)), ++ .align_mul = 16); ++ for (unsigned c = 0; c < v->num_components; c++) ++ s.params[i + c] = nir_channel(b, v, c); ++ } ++ s.table = addr_add(b, s.ring, s.params[TU_MESH_PARAM_TABLE]); ++ ++ nir_def *first = s.params[TU_MESH_PARAM_FIRST]; ++ nir_def *count = s.params[TU_MESH_PARAM_COUNT]; ++ s.count = nir_umin(b, nir_bcsel(b, nir_ult(b, first, count), ++ nir_isub(b, count, first), nir_imm_int(b, 0)), ++ s.params[TU_MESH_PARAM_MAX_COUNT]); ++ ++ nir_variable *sums = nir_variable_create( ++ b->shader, nir_var_mem_shared, ++ glsl_array_type(glsl_uint_type(), TU_MESH_SETUP_WORKGROUP_SIZE + 1, 4), ++ "sums"); ++ nir_variable *start = nir_local_variable_create(b->impl, glsl_uint_type(), "start"); ++ nir_def *t = nir_load_local_invocation_index(b); ++ ++ nir_store_var(b, start, nir_imm_int(b, 0), 0x1); ++ setup_range_loop(b, &s, start, false); ++ nir_store_array_var(b, sums, t, nir_load_var(b, start), 0x1); ++ emit_workgroup_barrier(b); ++ ++ nir_push_if(b, nir_ieq_imm(b, t, 0)); ++ { ++ nir_variable *i = nir_local_variable_create(b->impl, glsl_uint_type(), "i"); ++ nir_variable *sum = nir_local_variable_create(b->impl, glsl_uint_type(), "sum"); ++ nir_store_var(b, i, nir_imm_int(b, 0), 0x1); ++ nir_store_var(b, sum, nir_imm_int(b, 0), 0x1); ++ ++ nir_loop *loop = nir_push_loop(b); ++ { ++ nir_def *i_v = nir_load_var(b, i); ++ nir_break_if(b, nir_uge_imm(b, i_v, TU_MESH_SETUP_WORKGROUP_SIZE)); ++ nir_def *sum_v = nir_load_var(b, sum); ++ nir_def *v = nir_load_array_var(b, sums, i_v); ++ nir_store_array_var(b, sums, i_v, sum_v, 0x1); ++ nir_store_var(b, sum, nir_iadd(b, sum_v, v), 0x1); ++ nir_store_var(b, i, nir_iadd_imm(b, i_v, 1), 0x1); ++ } ++ nir_pop_loop(b, loop); ++ ++ nir_store_array_var_imm(b, sums, TU_MESH_SETUP_WORKGROUP_SIZE, ++ nir_load_var(b, sum), 0x1); ++ } ++ nir_pop_if(b, NULL); ++ emit_workgroup_barrier(b); ++ ++ nir_store_var(b, start, nir_load_array_var(b, sums, t), 0x1); ++ setup_range_loop(b, &s, start, true); ++ ++ nir_def *total = nir_load_array_var_imm(b, sums, TU_MESH_SETUP_WORKGROUP_SIZE); ++ nir_push_if(b, nir_ieq_imm(b, t, 0)); ++ { ++ nir_store_global(b, s.count, s.table, .align_mul = 16); ++ } ++ nir_pop_if(b, NULL); ++ ++ nir_push_if(b, nir_ult(b, t, s.params[TU_MESH_PARAM_CHUNKS])); ++ { ++ nir_def *chunk = s.params[TU_MESH_PARAM_CHUNK]; ++ nir_def *base = nir_imul(b, t, chunk); ++ nir_def *left = nir_bcsel(b, nir_ult(b, base, total), ++ nir_isub(b, total, base), nir_imm_int(b, 0)); ++ nir_def *groups = nir_umin(b, left, chunk); ++ nir_def *args = ++ addr_add(b, s.table, ++ nir_iadd_imm(b, nir_imul_imm(b, t, TU_MESH_ARGS_SIZE), ++ TU_MESH_TABLE_ARGS)); ++ nir_store_global(b, nir_vec4(b, groups, nir_imm_int(b, 1), nir_imm_int(b, 1), ++ nir_b2i32(b, nir_ine_imm(b, groups, 0))), ++ args, .align_mul = 16); ++ nir_store_global(b, nir_vec4(b, nir_imul(b, groups, ++ s.params[TU_MESH_PARAM_VERTICES]), ++ nir_imm_int(b, 1), nir_imm_int(b, 0), ++ nir_imm_int(b, 0)), ++ addr_add(b, args, nir_imm_int(b, TU_MESH_ARG_DRAW)), ++ .align_mul = 16); ++ } ++ nir_pop_if(b, NULL); ++ ++ NIR_PASS(_, b->shader, nir_lower_vars_to_ssa); ++ return b->shader; ++} +diff --git a/src/freedreno/vulkan/tu_mesh.h b/src/freedreno/vulkan/tu_mesh.h +new file mode 100644 +index 0000000..1e9fda0 +--- /dev/null ++++ b/src/freedreno/vulkan/tu_mesh.h +@@ -0,0 +1,127 @@ ++/* ++ * Copyright © 2026 MaxsTechReview ++ * SPDX-License-Identifier: MIT ++ */ ++ ++#ifndef TU_MESH_H ++#define TU_MESH_H ++ ++#include "tu_common.h" ++ ++/* Mesh and task shaders run as compute dispatches that write into a ++ * device-wide ring. A generated vertex shader then pulls one vertex per ++ * primitive corner out of the ring in a non-indexed draw. ++ * ++ * Ring layout: ++ * - setup parameters (enum tu_mesh_param); ++ * - a mesh and a task table, each holding an entry count, per-chunk ++ * dispatch and draw arguments, and entries sorted by first linear ++ * workgroup (first workgroup, group counts, draw id, task slot); ++ * - task slots: launch dimensions and draw id, followed by the payload; ++ * - records: per mesh workgroup vertex, primitive and index data. ++ */ ++#define TU_MESH_RING_SIZE (64u << 20) ++ ++enum tu_mesh_param { ++ TU_MESH_PARAM_TABLE, ++ TU_MESH_PARAM_SOURCE, ++ TU_MESH_PARAM_SRC_LO, ++ TU_MESH_PARAM_SRC_HI, ++ TU_MESH_PARAM_STRIDE, ++ TU_MESH_PARAM_MAX_COUNT, ++ TU_MESH_PARAM_COUNT, ++ TU_MESH_PARAM_FIRST, ++ TU_MESH_PARAM_CHUNK, ++ TU_MESH_PARAM_VERTICES, ++ TU_MESH_PARAM_CHUNKS, ++ TU_MESH_PARAM_NUM, ++}; ++ ++enum tu_mesh_source { ++ TU_MESH_SOURCE_INDIRECT, ++ TU_MESH_SOURCE_TASK, ++}; ++ ++#define TU_MESH_PARAMS_OFFSET 0 ++#define TU_MESH_MAX_CHUNKS 32 ++#define TU_MESH_MAX_TASK_CHUNKS 4 ++#define TU_MESH_TABLE_MAX_ENTRIES 16384 ++#define TU_MESH_TABLE_ENTRY_SIZE 32 ++/* Chunk arguments: dispatch size, enable flag, then a VkDrawIndirectCommand. */ ++#define TU_MESH_ARGS_SIZE 32 ++#define TU_MESH_ARG_ENABLE 12 ++#define TU_MESH_ARG_DRAW 16 ++#define TU_MESH_TABLE_ARGS 64 ++#define TU_MESH_TABLE_ENTRIES \ ++ (TU_MESH_TABLE_ARGS + TU_MESH_MAX_CHUNKS * TU_MESH_ARGS_SIZE) ++#define TU_MESH_TABLE_SIZE \ ++ (TU_MESH_TABLE_ENTRIES + TU_MESH_TABLE_MAX_ENTRIES * TU_MESH_TABLE_ENTRY_SIZE) ++#define TU_MESH_MS_TABLE_OFFSET 64 ++#define TU_MESH_TS_TABLE_OFFSET (TU_MESH_MS_TABLE_OFFSET + TU_MESH_TABLE_SIZE) ++#define TU_MESH_TASK_OFFSET (TU_MESH_TS_TABLE_OFFSET + TU_MESH_TABLE_SIZE) ++#define TU_MESH_TASK_HEADER_SIZE 16 ++#define TU_MESH_TASK_SIZE (16u << 20) ++#define TU_MESH_RECORD_OFFSET (TU_MESH_TASK_OFFSET + TU_MESH_TASK_SIZE) ++#define TU_MESH_RECORD_SIZE (TU_MESH_RING_SIZE - TU_MESH_RECORD_OFFSET) ++ ++#define TU_MESH_SETUP_WORKGROUP_SIZE 128 ++#define TU_MESH_DEAD_INDEX 0xffffffffu ++ ++/* The fragment shader reads a mesh shader's PrimitiveId from this generic ++ * varying, since only a geometry shader may write the hardware one. ++ */ ++#define TU_MESH_PRIMITIVE_ID_SLOT VARYING_SLOT_TEX0 ++ ++struct tu_mesh_io { ++ uint8_t vertex_slot[VARYING_SLOT_MAX]; ++ uint8_t prim_slot[VARYING_SLOT_MAX]; ++ unsigned vertex_slots; ++ unsigned prim_slots; ++ ++ unsigned max_vertices; ++ unsigned max_primitives; ++ unsigned verts_per_prim; ++ ++ unsigned prim_offset; ++ unsigned index_offset; ++ unsigned cull_offset; ++ unsigned stride; ++ unsigned chunk_workgroups; ++ ++ bool writes_primitive_id; ++}; ++ ++/* State a mesh or task compute variant carries into the draw path. */ ++struct tu_mesh_state { ++ uint32_t stride; ++ uint32_t chunk_workgroups; ++ uint32_t task_payload_stride; ++ uint16_t max_primitives; ++ uint8_t verts_per_prim; ++ uint8_t topology; ++}; ++ ++void ++tu_mesh_gather_io(const nir_shader *ms, struct tu_mesh_io *io); ++ ++nir_shader * ++tu_mesh_build_vs(const nir_shader *ms, const struct tu_mesh_io *io, ++ const nir_shader_compiler_options *options); ++ ++void ++tu_mesh_lower_ms(nir_shader *ms, const struct tu_mesh_io *io, ++ unsigned task_payload_stride); ++ ++unsigned ++tu_mesh_lower_ts(nir_shader *ts); ++ ++unsigned ++tu_mesh_task_chunk(unsigned task_payload_stride); ++ ++void ++tu_mesh_lower_fs_inputs(nir_shader *fs, bool remap_primitive_id); ++ ++nir_shader * ++tu_mesh_build_setup_cs(const nir_shader_compiler_options *options); ++ ++#endif /* TU_MESH_H */ +diff --git a/src/freedreno/vulkan/tu_pipeline.cc b/src/freedreno/vulkan/tu_pipeline.cc +index 75c0782..949a4e0 100644 +--- a/src/freedreno/vulkan/tu_pipeline.cc ++++ b/src/freedreno/vulkan/tu_pipeline.cc +@@ -53,6 +53,17 @@ emit_load_state(struct tu_cs *cs, unsigned opcode, enum a6xx_state_type st, + tu_cs_emit_qw(cs, offset | (base << 28)); + } + ++/* Task and mesh shaders run as compute dispatches during the draw, so their ++ * descriptors are not prefetched with the graphics stages. ++ */ ++static VkShaderStageFlags ++load_state_stages(const struct tu_pipeline *pipeline, ++ const struct tu_descriptor_set_binding_layout *binding) ++{ ++ return pipeline->active_stages & binding->shader_stages & ++ ~(VK_SHADER_STAGE_TASK_BIT_EXT | VK_SHADER_STAGE_MESH_BIT_EXT); ++} ++ + static unsigned + tu6_load_state_size(struct tu_pipeline *pipeline, + struct tu_pipeline_layout *layout) +@@ -68,7 +79,7 @@ tu6_load_state_size(struct tu_pipeline *pipeline, + struct tu_descriptor_set_binding_layout *binding = &set_layout->binding[j]; + unsigned count = 0; + /* See comment in tu6_emit_load_state(). */ +- VkShaderStageFlags stages = pipeline->active_stages & binding->shader_stages; ++ VkShaderStageFlags stages = load_state_stages(pipeline, binding); + unsigned stage_count = util_bitcount(stages); + + if (!binding->array_size) +@@ -155,7 +166,7 @@ tu6_emit_load_state(struct tu_device *device, + * stages aren't present in a used pipeline. We don't want to emit + * loads for unused descriptors. + */ +- VkShaderStageFlags stages = pipeline->active_stages & binding->shader_stages; ++ VkShaderStageFlags stages = load_state_stages(pipeline, binding); + unsigned count = binding->array_size; + + /* If this is a variable-count descriptor, then the array_size is an +@@ -409,7 +420,7 @@ tu6_emit_xs_config(struct tu_crb &crb, struct tu_shader_stages stages) + } + TU_GENX(tu6_emit_xs_config); + +-static void ++void + tu6_emit_dynamic_offset(struct tu_cs *cs, + const struct ir3_shader_variant *xs, + const struct tu_shader *shader, +@@ -1561,7 +1572,7 @@ tu_hash_shaders(unsigned char *hash, + if (layout) + _mesa_blake3_update(&ctx, layout->blake3, sizeof(layout->blake3)); + +- for (int i = 0; i < MESA_SHADER_STAGES; ++i) { ++ for (int i = 0; i < MESA_SHADER_MESH_STAGES; ++i) { + if (stages[i] || nir[i]) { + tu_hash_stage(&ctx, pipeline_flags, stages[i], nir[i], &keys[i]); + } +@@ -1735,13 +1746,13 @@ tu_pipeline_builder_compile_shaders(struct tu_pipeline_builder *builder, + struct tu_pipeline *pipeline) + { + VkResult result = VK_SUCCESS; +- const VkPipelineShaderStageCreateInfo *stage_infos[MESA_SHADER_STAGES] = { ++ const VkPipelineShaderStageCreateInfo *stage_infos[MESA_SHADER_MESH_STAGES] = { + NULL + }; + VkPipelineCreationFeedback pipeline_feedback = { + .flags = VK_PIPELINE_CREATION_FEEDBACK_VALID_BIT, + }; +- VkPipelineCreationFeedback stage_feedbacks[MESA_SHADER_STAGES] = { 0 }; ++ VkPipelineCreationFeedback stage_feedbacks[MESA_SHADER_MESH_STAGES] = { 0 }; + + const bool executable_info = + builder->create_flags & +@@ -1753,6 +1764,12 @@ tu_pipeline_builder_compile_shaders(struct tu_pipeline_builder *builder, + + int64_t pipeline_start = os_time_get_nano(); + ++ if (builder->active_stages & VK_SHADER_STAGE_MESH_BIT_EXT) { ++ result = tu_init_mesh_shading(builder->device); ++ if (result != VK_SUCCESS) ++ return result; ++ } ++ + const VkPipelineCreationFeedbackCreateInfo *creation_feedback = + vk_find_struct_const(builder->create_info->pNext, PIPELINE_CREATION_FEEDBACK_CREATE_INFO); + +@@ -1963,9 +1980,11 @@ tu_pipeline_builder_compile_shaders(struct tu_pipeline_builder *builder, + unsigned char shader_blake3[BLAKE3_KEY_LEN + 1]; + memcpy(shader_blake3, pipeline_blake3, sizeof(pipeline_blake3)); + ++ bool has_mesh = stage_infos[MESA_SHADER_MESH] || nir[MESA_SHADER_MESH]; + for (mesa_shader_stage stage = MESA_SHADER_VERTEX; stage < ARRAY_SIZE(nir); + stage = (mesa_shader_stage) (stage + 1)) { +- if (stage_infos[stage] || nir[stage]) { ++ if (stage_infos[stage] || nir[stage] || ++ (stage == MESA_SHADER_VERTEX && has_mesh)) { + bool shader_application_cache_hit; + shader_blake3[BLAKE3_KEY_LEN] = (unsigned char) stage; + shaders[stage] = +diff --git a/src/freedreno/vulkan/tu_pipeline.h b/src/freedreno/vulkan/tu_pipeline.h +index 791a4fc..31bf5c4 100644 +--- a/src/freedreno/vulkan/tu_pipeline.h ++++ b/src/freedreno/vulkan/tu_pipeline.h +@@ -70,7 +70,7 @@ struct tu_nir_shaders + /* This is optional, and is only filled out when a library pipeline is + * compiled with RETAIN_LINK_TIME_OPTIMIZATION_INFO. + */ +- nir_shader *nir[MESA_SHADER_STAGES]; ++ nir_shader *nir[MESA_SHADER_MESH_STAGES]; + }; + + extern const struct vk_pipeline_cache_object_ops tu_nir_shaders_ops; +@@ -216,7 +216,7 @@ struct tu_pipeline + /* draw states for the pipeline */ + struct tu_draw_state load_state; + +- struct tu_shader *shaders[MESA_SHADER_STAGES]; ++ struct tu_shader *shaders[MESA_SHADER_MESH_STAGES]; + + struct tu_program_state program; + +@@ -243,7 +243,7 @@ struct tu_graphics_lib_pipeline { + struct { + nir_shader *nir; + struct tu_shader_key key; +- } shaders[MESA_SHADER_FRAGMENT + 1]; ++ } shaders[MESA_SHADER_MESH_STAGES]; + + /* Used to stitch together an overall layout for the final pipeline. */ + struct tu_descriptor_set_layout *layouts[MAX_SETS]; +@@ -314,6 +314,12 @@ template + void + tu6_emit_shared_consts_enable(struct tu_crb &crb, bool shared_consts_enable); + ++void ++tu6_emit_dynamic_offset(struct tu_cs *cs, ++ const struct ir3_shader_variant *xs, ++ const struct tu_shader *shader, ++ const struct tu_program_state *program); ++ + template + void + tu6_emit_vpc(struct tu_cs *cs, +diff --git a/src/freedreno/vulkan/tu_shader.cc b/src/freedreno/vulkan/tu_shader.cc +index 57c95e6..ad9d563 100644 +--- a/src/freedreno/vulkan/tu_shader.cc ++++ b/src/freedreno/vulkan/tu_shader.cc +@@ -1165,7 +1165,8 @@ lower_inline_ubo(nir_builder *b, nir_intrinsic_instr *intrin, void *cb_data) + params->dev->physical_device->info->props.load_inline_uniforms_via_preamble_ldgk; + + for (unsigned i = 0; i < const_state->num_inline_ubos; i++) { +- if (const_state->ubos[i].base == binding.desc_set && ++ if (!const_state->ubos[i].mesh_ring && ++ const_state->ubos[i].base == binding.desc_set && + const_state->ubos[i].offset == binding_layout->offset) { + range = const_state->ubos[i].size_vec4 * 16; + if (use_ldg_k) { +@@ -1219,6 +1220,46 @@ lower_inline_ubo(nir_builder *b, nir_intrinsic_instr *intrin, void *cb_data) + return true; + } + ++static bool ++shader_uses_mesh_ring(nir_shader *shader) ++{ ++ nir_foreach_function_impl (impl, shader) { ++ nir_foreach_block (block, impl) { ++ nir_foreach_instr (instr, block) { ++ if (instr->type == nir_instr_type_intrinsic && ++ nir_instr_as_intrinsic(instr)->intrinsic == ++ nir_intrinsic_load_mesh_ring_ir3) ++ return true; ++ } ++ } ++ } ++ return false; ++} ++ ++static bool ++lower_mesh_ring(nir_builder *b, nir_intrinsic_instr *intrin, void *cb_data) ++{ ++ if (intrin->intrinsic != nir_intrinsic_load_mesh_ring_ir3) ++ return false; ++ ++ struct lower_instr_params *params = (struct lower_instr_params *) cb_data; ++ struct tu_const_state *const_state = ¶ms->shader->const_state; ++ unsigned i = const_state->num_inline_ubos - 1; ++ assert(const_state->ubos[i].mesh_ring); ++ ++ b->cursor = nir_before_instr(&intrin->instr); ++ nir_def *addr; ++ if (params->dev->physical_device->info->props.load_inline_uniforms_via_preamble_ldgk) { ++ addr = ir3_load_driver_ubo(b, 2, &const_state->inline_uniforms_ubo, i * 2); ++ } else { ++ addr = nir_load_const_ir3(b, 2, 32, nir_imm_int(b, 0), ++ .base = const_state->ubos[i].const_offset_vec4 * 4); ++ } ++ ++ nir_def_replace(&intrin->def, addr); ++ return true; ++} ++ + /* Returns whether we can speculatively access through a Vulkan descriptor. + * Used for SSBO/UBO accesses, where vtn produces + * vulkan_resource_index/load_vulkan_descriptor. See build_bindless() for the +@@ -1553,6 +1594,20 @@ tu_lower_io(nir_shader *shader, struct tu_device *dev, + } + } + ++ if (shader_uses_mesh_ring(shader)) { ++ const_state->ubos[const_state->num_inline_ubos++] = ++ (struct tu_inline_ubo) { ++ .push_address = !use_ldg_k, ++ .mesh_ring = true, ++ .const_offset_vec4 = ++ const_allocs->max_const_offset_vec4 + ldgk_consts, ++ .size_vec4 = 1, ++ }; ++ ++ if (!use_ldg_k) ++ ldgk_consts += align(1, dev->compiler->const_upload_unit); ++ } ++ + ir3_const_alloc(const_allocs, IR3_CONST_ALLOC_INLINE_UNIFORM_ADDRS, ldgk_consts, 1); + + if (dev->physical_device->compiler_options.enable_ssbo_emulation) { +@@ -1584,6 +1639,9 @@ tu_lower_io(nir_shader *shader, struct tu_device *dev, + progress |= nir_shader_intrinsics_pass(shader, lower_inline_ubo, + nir_metadata_none, + ¶ms); ++ progress |= nir_shader_intrinsics_pass(shader, lower_mesh_ring, ++ nir_metadata_control_flow, ++ ¶ms); + } + + progress |= nir_shader_instructions_pass(shader, +@@ -3347,6 +3405,9 @@ tu_shader_serialize(struct vk_pipeline_cache_object *object, + case MESA_SHADER_FRAGMENT: + blob_write_bytes(blob, &shader->fs, sizeof(shader->fs)); + break; ++ case MESA_SHADER_COMPUTE: ++ blob_write_bytes(blob, &shader->mesh, sizeof(shader->mesh)); ++ break; + default: + break; + } +@@ -3388,6 +3449,9 @@ tu_shader_deserialize(struct vk_pipeline_cache *cache, + case MESA_SHADER_FRAGMENT: + blob_copy_bytes(blob, &shader->fs, sizeof(shader->fs)); + break; ++ case MESA_SHADER_COMPUTE: ++ blob_copy_bytes(blob, &shader->mesh, sizeof(shader->mesh)); ++ break; + default: + break; + } +@@ -3575,6 +3639,8 @@ tu_shader_create(struct tu_device *dev, + return VK_ERROR_OUT_OF_HOST_MEMORY; + + shader->per_layer_viewport = info->per_layer_viewport; ++ if (nir->info.stage == MESA_SHADER_COMPUTE) ++ shader->mesh = info->mesh; + + if (nir->info.stage == MESA_SHADER_FRAGMENT && + key->fdm_per_layer) { +@@ -3788,6 +3854,50 @@ tu6_get_tessmode(const struct nir_shader *shader) + } + } + ++/* Replaces the task and mesh shaders with compute shaders feeding a ++ * generated vertex shader. ++ */ ++static void ++tu_lower_mesh_pipeline(struct tu_device *dev, nir_shader **nir, ++ struct tu_shader_info *info, void *mem_ctx) ++{ ++ nir_shader *ms = nir[MESA_SHADER_MESH]; ++ struct tu_mesh_io io; ++ tu_mesh_gather_io(ms, &io); ++ ++ nir[MESA_SHADER_VERTEX] = ++ tu_mesh_build_vs(ms, &io, ir3_get_compiler_options(dev->compiler)); ++ ralloc_steal(mem_ctx, nir[MESA_SHADER_VERTEX]); ++ ++ unsigned payload_stride = 0; ++ if (nir[MESA_SHADER_TASK]) { ++ payload_stride = tu_mesh_lower_ts(nir[MESA_SHADER_TASK]); ++ info[MESA_SHADER_TASK].mesh = (struct tu_mesh_state) { ++ .chunk_workgroups = tu_mesh_task_chunk(payload_stride), ++ .task_payload_stride = payload_stride, ++ }; ++ } ++ ++ VkPrimitiveTopology topology = ++ io.verts_per_prim == 1 ? VK_PRIMITIVE_TOPOLOGY_POINT_LIST : ++ io.verts_per_prim == 2 ? VK_PRIMITIVE_TOPOLOGY_LINE_LIST : ++ VK_PRIMITIVE_TOPOLOGY_TRIANGLE_LIST; ++ ++ tu_mesh_lower_ms(ms, &io, payload_stride); ++ info[MESA_SHADER_MESH].mesh = (struct tu_mesh_state) { ++ .stride = io.stride, ++ .chunk_workgroups = io.chunk_workgroups, ++ .task_payload_stride = payload_stride, ++ .max_primitives = (uint16_t) io.max_primitives, ++ .verts_per_prim = (uint8_t) io.verts_per_prim, ++ .topology = (uint8_t) topology, ++ }; ++ ++ if (nir[MESA_SHADER_FRAGMENT]) ++ tu_mesh_lower_fs_inputs(nir[MESA_SHADER_FRAGMENT], ++ io.writes_primitive_id); ++} ++ + VkResult + tu_compile_shaders(struct tu_device *device, + VkPipelineCreateFlags2KHR pipeline_flags, +@@ -3803,11 +3913,13 @@ tu_compile_shaders(struct tu_device *device, + VkPipelineCreationFeedback *stage_feedbacks) + { + struct ir3_shader_key ir3_key = {}; +- struct tu_shader_info info[MESA_SHADER_STAGES] = {}; ++ struct tu_shader_info info[MESA_SHADER_MESH_STAGES] = {}; ++ nir_shader *mesh_nir[2] = {}; + VkResult result = VK_SUCCESS; + void *mem_ctx = ralloc_context(NULL); + +- for (mesa_shader_stage stage = MESA_SHADER_VERTEX; stage < MESA_SHADER_STAGES; ++ for (mesa_shader_stage stage = MESA_SHADER_VERTEX; ++ stage < MESA_SHADER_MESH_STAGES; + stage = (mesa_shader_stage) (stage + 1)) { + const VkPipelineShaderStageCreateInfo *stage_info = stage_infos[stage]; + if (!stage_info) +@@ -3831,7 +3943,7 @@ tu_compile_shaders(struct tu_device *device, + + if (nir_initial_disasm) { + for (mesa_shader_stage stage = MESA_SHADER_VERTEX; +- stage < MESA_SHADER_STAGES; ++ stage < MESA_SHADER_MESH_STAGES; + stage = (mesa_shader_stage) (stage + 1)) { + if (!nir[stage]) + continue; +@@ -3841,7 +3953,20 @@ tu_compile_shaders(struct tu_device *device, + } + } + +- for (mesa_shader_stage stage = MESA_SHADER_VERTEX; stage < MESA_SHADER_STAGES; ++ if (nir[MESA_SHADER_MESH] && ++ nir[MESA_SHADER_MESH]->info.stage == MESA_SHADER_MESH) { ++ if (nir_out) { ++ mesh_nir[0] = nir_shader_clone(NULL, nir[MESA_SHADER_MESH]); ++ if (nir[MESA_SHADER_TASK]) ++ mesh_nir[1] = nir_shader_clone(NULL, nir[MESA_SHADER_TASK]); ++ } ++ tu_lower_mesh_pipeline(device, nir, info, mem_ctx); ++ } else if (nir[MESA_SHADER_FRAGMENT]) { ++ tu_mesh_lower_fs_inputs(nir[MESA_SHADER_FRAGMENT], false); ++ } ++ ++ for (mesa_shader_stage stage = MESA_SHADER_VERTEX; ++ stage < MESA_SHADER_MESH_STAGES; + stage = (mesa_shader_stage) (stage + 1)) { + if (!nir[stage]) + continue; +@@ -3853,16 +3978,18 @@ tu_compile_shaders(struct tu_device *device, + stage_feedbacks[stage].duration += os_time_get_nano() - stage_start; + } + +- tu_link_shaders(device, nir, MESA_SHADER_STAGES); ++ tu_link_shaders(device, nir, MESA_SHADER_FRAGMENT + 1); + + if (nir_out) { + for (mesa_shader_stage stage = MESA_SHADER_VERTEX; +- stage < MESA_SHADER_STAGES; stage = (mesa_shader_stage) (stage + 1)) { +- if (!nir[stage]) ++ stage <= MESA_SHADER_FRAGMENT; stage = (mesa_shader_stage) (stage + 1)) { ++ if (!nir[stage] || (stage == MESA_SHADER_VERTEX && nir[MESA_SHADER_MESH])) + continue; + + nir_out[stage] = nir_shader_clone(NULL, nir[stage]); + } ++ nir_out[MESA_SHADER_MESH] = mesh_nir[0]; ++ nir_out[MESA_SHADER_TASK] = mesh_nir[1]; + } + + /* With pipelines, tessellation modes can be set on either shader, for +@@ -3915,7 +4042,8 @@ tu_compile_shaders(struct tu_device *device, + if (nir[MESA_SHADER_TESS_CTRL] && !nir[MESA_SHADER_FRAGMENT]) + ir3_key.tcs_store_primid = true; + +- for (mesa_shader_stage stage = MESA_SHADER_VERTEX; stage < MESA_SHADER_STAGES; ++ for (mesa_shader_stage stage = MESA_SHADER_VERTEX; ++ stage < MESA_SHADER_MESH_STAGES; + stage = (mesa_shader_stage) (stage + 1)) { + if (!nir[stage] || shaders[stage]) + continue; +@@ -3945,7 +4073,8 @@ tu_compile_shaders(struct tu_device *device, + fail: + ralloc_free(mem_ctx); + +- for (mesa_shader_stage stage = MESA_SHADER_VERTEX; stage < MESA_SHADER_STAGES; ++ for (mesa_shader_stage stage = MESA_SHADER_VERTEX; ++ stage < MESA_SHADER_MESH_STAGES; + stage = (mesa_shader_stage) (stage + 1)) { + if (shaders[stage]) { + tu_shader_destroy(device, shaders[stage]); +@@ -4125,6 +4254,45 @@ tu_destroy_empty_shaders(struct tu_device *dev) + vk_pipeline_cache_object_unref(&dev->vk, &dev->empty_fs_fdm->base); + } + ++static VkResult ++tu_mesh_setup_create(struct tu_device *dev, struct tu_shader **shader) ++{ ++ nir_shader *nir = ++ tu_mesh_build_setup_cs(ir3_get_compiler_options(dev->compiler)); ++ ++ struct tu_shader_key key = {}; ++ tu_shader_key_subgroup_size(&key, true, false, NULL, dev); ++ struct ir3_shader_key ir3_key = {}; ++ struct tu_shader_info info = {}; ++ struct tu_pipeline_layout layout = {}; ++ ++ tu_lower_nir(dev, nir, &key, &ir3_key, &info); ++ return tu_shader_create(dev, shader, nir, &key, &info, &ir3_key, NULL, 0, ++ &layout, false); ++} ++ ++VkResult ++tu_init_mesh_shading(struct tu_device *dev) ++{ ++ if (p_atomic_read(&dev->mesh_setup)) ++ return VK_SUCCESS; ++ ++ VkResult result = VK_SUCCESS; ++ mtx_lock(&dev->mutex); ++ if (!dev->mesh_ring) { ++ result = tu_bo_init_new(dev, NULL, &dev->mesh_ring, TU_MESH_RING_SIZE, ++ TU_BO_ALLOC_INTERNAL_RESOURCE, "mesh ring"); ++ } ++ if (result == VK_SUCCESS && !dev->mesh_setup) { ++ struct tu_shader *setup; ++ result = tu_mesh_setup_create(dev, &setup); ++ if (result == VK_SUCCESS) ++ p_atomic_set(&dev->mesh_setup, setup); ++ } ++ mtx_unlock(&dev->mutex); ++ return result; ++} ++ + void + tu_shader_destroy(struct tu_device *dev, + struct tu_shader *shader) +diff --git a/src/freedreno/vulkan/tu_shader.h b/src/freedreno/vulkan/tu_shader.h +index b19287b..5665d21 100644 +--- a/src/freedreno/vulkan/tu_shader.h ++++ b/src/freedreno/vulkan/tu_shader.h +@@ -14,6 +14,7 @@ + + #include "tu_cs.h" + #include "tu_descriptor_set.h" ++#include "tu_mesh.h" + #include "tu_suballoc.h" + + struct tu_inline_ubo +@@ -25,6 +26,9 @@ struct tu_inline_ubo + /* If true, push the base address instead */ + bool push_address; + ++ /* Push the mesh shading ring address instead of descriptor data */ ++ bool mesh_ring; ++ + /* Push it to this location in the const file, in vec4s */ + unsigned const_offset_vec4; + +@@ -49,7 +53,7 @@ struct tu_const_state + struct tu_push_constant_range push_consts; + uint32_t dynamic_offset_loc; + unsigned num_inline_ubos; +- struct tu_inline_ubo ubos[MAX_INLINE_UBOS]; ++ struct tu_inline_ubo ubos[MAX_INLINE_UBOS + 1]; + uint32_t num_bindless_base_addresses; + uint32_t bindless_base_const_offset_vec4; + +@@ -123,6 +127,9 @@ struct tu_shader + */ + bool read_only_input_attachments; + } fs; ++ ++ /* Mesh and task shaders compiled to compute. */ ++ struct tu_mesh_state mesh; + }; + }; + +@@ -154,6 +161,7 @@ struct tu_shader_key { + */ + struct tu_shader_info { + bool per_layer_viewport; ++ struct tu_mesh_state mesh; + }; + + extern const struct vk_pipeline_cache_object_ops tu_shader_ops; +@@ -265,6 +273,9 @@ tu_init_empty_shaders(struct tu_device *device); + void + tu_destroy_empty_shaders(struct tu_device *device); + ++VkResult ++tu_init_mesh_shading(struct tu_device *device); ++ + void + tu_shader_destroy(struct tu_device *dev, + struct tu_shader *shader); diff --git a/linux/patches/0002-tu-ir3-Support-a-required-subgroup-size-of-half-a-wa.patch b/linux/patches/0002-tu-ir3-Support-a-required-subgroup-size-of-half-a-wa.patch new file mode 100644 index 0000000..616c600 --- /dev/null +++ b/linux/patches/0002-tu-ir3-Support-a-required-subgroup-size-of-half-a-wa.patch @@ -0,0 +1,484 @@ +From 0000000000000000000000000000000000000000 Mon Sep 17 00:00:00 2001 +From: MaxsTechReview +Date: Sun, 27 Sep 2026 03:08:53 -0400 +Subject: [PATCH 2/5] tu,ir3: Support a required subgroup size of half a wave + +A8XX only runs 64-wide waves, so the smallest subgroup size was 64 and +vkd3d-proton rejected every compute shader declaring WaveSize(32). + +Report a minimum subgroup size of 32 and lower such shaders onto single +waves, each half of a wave acting as one subgroup. The lowering stays free +of divergent control flow: the scan macros lose invocations when placed in +both sides of a divergent branch. +--- + src/freedreno/ir3/ir3_lower_subgroups.c | 290 ++++++++++++++++++++++++ + src/freedreno/ir3/ir3_nir.c | 3 + + src/freedreno/ir3/ir3_nir.h | 1 + + src/freedreno/ir3/ir3_shader.c | 6 + + src/freedreno/ir3/ir3_shader.h | 2 + + src/freedreno/vulkan/tu_device.cc | 22 +- + src/freedreno/vulkan/tu_pipeline.cc | 8 +- + src/freedreno/vulkan/tu_pipeline.h | 1 + + src/freedreno/vulkan/tu_shader.cc | 6 +- + 9 files changed, 329 insertions(+), 10 deletions(-) + +diff --git a/src/freedreno/ir3/ir3_lower_subgroups.c b/src/freedreno/ir3/ir3_lower_subgroups.c +index c354289..cb9ad27 100644 +--- a/src/freedreno/ir3/ir3_lower_subgroups.c ++++ b/src/freedreno/ir3/ir3_lower_subgroups.c +@@ -876,3 +876,293 @@ ir3_nir_lower_shuffle(nir_shader *nir, struct ir3_shader *shader) + return nir_shader_lower_instructions(nir, filter_shuffle, lower_shuffle, + NULL); + } ++ ++/* Subgroups of half a wave: each half of a 64-wide wave is its own subgroup. ++ * Operations that depend on the subgroup size or span the subgroup are ++ * confined to the half of the invocation, leaving only whole-wave operations ++ * for the rest of the backend. ++ */ ++#define HALF_SUBGROUP_SIZE 32 ++ ++static nir_def * ++half_index(nir_builder *b) ++{ ++ return nir_ushr_imm(b, nir_load_subgroup_invocation(b), ++ util_logbase2(HALF_SUBGROUP_SIZE)); ++} ++ ++static nir_def * ++half_invocation(nir_builder *b) ++{ ++ return nir_iand_imm(b, nir_load_subgroup_invocation(b), ++ HALF_SUBGROUP_SIZE - 1); ++} ++ ++/* Ballot of the invocations in the half of the current invocation. */ ++static nir_def * ++half_ballot(nir_builder *b, nir_def *cond) ++{ ++ nir_def *ballot = nir_ballot(b, 2, 32, cond); ++ return nir_vector_extract(b, ballot, half_index(b)); ++} ++ ++static nir_def * ++to_ballot_type(nir_builder *b, nir_def *mask, const nir_def *def) ++{ ++ return nir_pad_vector_imm_int(b, nir_u2uN(b, mask, def->bit_size), 0, ++ def->num_components); ++} ++ ++static nir_def * ++from_ballot_type(nir_builder *b, nir_def *value) ++{ ++ return nir_u2u32(b, nir_channel(b, value, 0)); ++} ++ ++static nir_def * ++build_half_mask(nir_builder *b, nir_intrinsic_op op) ++{ ++ nir_def *id = half_invocation(b); ++ ++ switch (op) { ++ case nir_intrinsic_load_subgroup_eq_mask: ++ return nir_ishl(b, nir_imm_int(b, 1), id); ++ case nir_intrinsic_load_subgroup_ge_mask: ++ return nir_ishl(b, nir_imm_int(b, ~0), id); ++ case nir_intrinsic_load_subgroup_gt_mask: ++ return nir_ishl(b, nir_imm_int(b, ~1), id); ++ case nir_intrinsic_load_subgroup_le_mask: ++ return nir_inot(b, nir_ishl(b, nir_imm_int(b, ~1), id)); ++ case nir_intrinsic_load_subgroup_lt_mask: ++ return nir_inot(b, nir_ishl(b, nir_imm_int(b, ~0), id)); ++ default: ++ UNREACHABLE("not a subgroup mask"); ++ } ++} ++ ++/* Reads a value from one invocation of each half. The indices are uniform ++ * across the wave, so no half has to run on its own. ++ */ ++static nir_def * ++build_half_read(nir_builder *b, nir_def *val, nir_def *index[2]) ++{ ++ nir_def *lo = nir_read_invocation(b, val, index[0]); ++ nir_def *hi = nir_read_invocation(b, val, index[1]); ++ return nir_bcsel(b, nir_ieq_imm(b, half_index(b), 0), lo, hi); ++} ++ ++static void ++build_half_first(nir_builder *b, nir_def *first[2]) ++{ ++ nir_def *ballot = nir_ballot(b, 2, 32, nir_imm_true(b)); ++ for (unsigned i = 0; i < 2; i++) { ++ nir_def *lsb = nir_find_lsb(b, nir_channel(b, ballot, i)); ++ first[i] = nir_iadd_imm(b, nir_iand_imm(b, lsb, HALF_SUBGROUP_SIZE - 1), ++ i * HALF_SUBGROUP_SIZE); ++ } ++} ++ ++static nir_def * ++build_half_read_first(nir_builder *b, nir_def *val) ++{ ++ nir_def *first[2]; ++ build_half_first(b, first); ++ return build_half_read(b, val, first); ++} ++ ++static nir_def * ++build_half_read_invocation(nir_builder *b, nir_def *val, nir_def *index) ++{ ++ nir_def *first[2], *indices[2]; ++ build_half_first(b, first); ++ for (unsigned i = 0; i < 2; i++) { ++ nir_def *id = nir_read_invocation(b, index, first[i]); ++ indices[i] = nir_iadd_imm(b, nir_iand_imm(b, id, HALF_SUBGROUP_SIZE - 1), ++ i * HALF_SUBGROUP_SIZE); ++ } ++ return build_half_read(b, val, indices); ++} ++ ++/* The upper half scans with the lower half replaced by the identity, which ++ * leaves the lower half's own scan unchanged. ++ */ ++static nir_def * ++build_half_scan(nir_builder *b, nir_intrinsic_instr *intr) ++{ ++ nir_op op = nir_intrinsic_reduction_op(intr); ++ nir_def *val = intr->src[0].ssa; ++ nir_const_value identity = nir_alu_binop_identity(op, val->bit_size); ++ nir_def *identity_def = ++ nir_replicate(b, nir_build_imm(b, 1, val->bit_size, &identity), ++ val->num_components); ++ nir_def *upper = nir_ine_imm(b, half_index(b), 0); ++ nir_def *upper_val = nir_bcsel(b, upper, val, identity_def); ++ ++ nir_def *lo, *hi; ++ if (intr->intrinsic == nir_intrinsic_inclusive_scan) { ++ lo = nir_inclusive_scan(b, val, .reduction_op = op); ++ hi = nir_inclusive_scan(b, upper_val, .reduction_op = op); ++ } else { ++ lo = nir_exclusive_scan(b, val, .reduction_op = op); ++ hi = nir_exclusive_scan(b, upper_val, .reduction_op = op); ++ } ++ return nir_bcsel(b, upper, hi, lo); ++} ++ ++static nir_def * ++build_num_half_subgroups(nir_builder *b) ++{ ++ nir_def *count; ++ ++ if (b->shader->info.workgroup_size_variable) { ++ nir_def *size = nir_load_workgroup_size(b); ++ count = nir_imul(b, nir_imul(b, nir_channel(b, size, 0), ++ nir_channel(b, size, 1)), ++ nir_channel(b, size, 2)); ++ } else { ++ count = nir_imm_int(b, nir_static_workgroup_size(b->shader)); ++ } ++ ++ return nir_udiv_imm(b, nir_iadd_imm(b, count, HALF_SUBGROUP_SIZE - 1), ++ HALF_SUBGROUP_SIZE); ++} ++ ++static nir_def * ++build_half_vote_eq(nir_builder *b, nir_intrinsic_instr *intr) ++{ ++ nir_def *val = intr->src[0].ssa; ++ nir_def *first = build_half_read_first(b, val); ++ nir_def *differs = nir_imm_false(b); ++ ++ for (unsigned i = 0; i < val->num_components; i++) { ++ nir_def *x = nir_channel(b, val, i); ++ nir_def *y = nir_channel(b, first, i); ++ differs = nir_ior(b, differs, intr->intrinsic == nir_intrinsic_vote_feq ++ ? nir_fneu(b, x, y) ++ : nir_ine(b, x, y)); ++ } ++ ++ return nir_ieq_imm(b, half_ballot(b, differs), 0); ++} ++ ++static bool ++lower_half_subgroup_intrinsic(nir_builder *b, nir_intrinsic_instr *intr, ++ void *data) ++{ ++ b->cursor = nir_before_instr(&intr->instr); ++ ++ nir_def *res; ++ switch (intr->intrinsic) { ++ case nir_intrinsic_load_subgroup_size: ++ res = nir_imm_int(b, HALF_SUBGROUP_SIZE); ++ break; ++ case nir_intrinsic_load_subgroup_invocation: ++ res = half_invocation(b); ++ break; ++ case nir_intrinsic_load_subgroup_id: ++ res = nir_iadd(b, nir_imul_imm(b, nir_load_subgroup_id(b), 2), ++ half_index(b)); ++ break; ++ case nir_intrinsic_load_num_subgroups: ++ res = build_num_half_subgroups(b); ++ break; ++ case nir_intrinsic_load_subgroup_eq_mask: ++ case nir_intrinsic_load_subgroup_ge_mask: ++ case nir_intrinsic_load_subgroup_gt_mask: ++ case nir_intrinsic_load_subgroup_le_mask: ++ case nir_intrinsic_load_subgroup_lt_mask: ++ res = to_ballot_type(b, build_half_mask(b, intr->intrinsic), &intr->def); ++ break; ++ case nir_intrinsic_ballot: ++ res = to_ballot_type(b, half_ballot(b, intr->src[0].ssa), &intr->def); ++ break; ++ case nir_intrinsic_inverse_ballot: ++ res = nir_test_mask(b, nir_ushr(b, from_ballot_type(b, intr->src[0].ssa), ++ half_invocation(b)), 1); ++ break; ++ case nir_intrinsic_ballot_bitfield_extract: ++ res = nir_test_mask(b, nir_ushr(b, from_ballot_type(b, intr->src[0].ssa), ++ intr->src[1].ssa), 1); ++ break; ++ case nir_intrinsic_ballot_bit_count_reduce: ++ res = nir_bit_count(b, from_ballot_type(b, intr->src[0].ssa)); ++ break; ++ case nir_intrinsic_ballot_bit_count_inclusive: ++ case nir_intrinsic_ballot_bit_count_exclusive: { ++ nir_intrinsic_op mask = ++ intr->intrinsic == nir_intrinsic_ballot_bit_count_inclusive ++ ? nir_intrinsic_load_subgroup_le_mask ++ : nir_intrinsic_load_subgroup_lt_mask; ++ res = nir_bit_count(b, nir_iand(b, from_ballot_type(b, intr->src[0].ssa), ++ build_half_mask(b, mask))); ++ break; ++ } ++ case nir_intrinsic_ballot_find_lsb: ++ res = nir_find_lsb(b, from_ballot_type(b, intr->src[0].ssa)); ++ break; ++ case nir_intrinsic_ballot_find_msb: ++ res = nir_ufind_msb(b, from_ballot_type(b, intr->src[0].ssa)); ++ break; ++ case nir_intrinsic_vote_any: ++ res = nir_ine_imm(b, half_ballot(b, intr->src[0].ssa), 0); ++ break; ++ case nir_intrinsic_vote_all: ++ res = nir_ieq_imm(b, half_ballot(b, nir_inot(b, intr->src[0].ssa)), 0); ++ break; ++ case nir_intrinsic_vote_ieq: ++ case nir_intrinsic_vote_feq: ++ res = build_half_vote_eq(b, intr); ++ break; ++ case nir_intrinsic_elect: ++ res = nir_ieq(b, half_invocation(b), ++ nir_find_lsb(b, half_ballot(b, nir_imm_true(b)))); ++ break; ++ case nir_intrinsic_read_first_invocation: ++ res = build_half_read_first(b, intr->src[0].ssa); ++ break; ++ case nir_intrinsic_read_invocation: ++ res = build_half_read_invocation(b, intr->src[0].ssa, intr->src[1].ssa); ++ break; ++ case nir_intrinsic_inclusive_scan: ++ case nir_intrinsic_exclusive_scan: ++ res = build_half_scan(b, intr); ++ break; ++ case nir_intrinsic_shuffle: { ++ nir_def *base = nir_iand_imm(b, nir_load_subgroup_invocation(b), ++ ~(HALF_SUBGROUP_SIZE - 1)); ++ nir_def *id = nir_iand_imm(b, intr->src[1].ssa, HALF_SUBGROUP_SIZE - 1); ++ res = nir_shuffle(b, intr->src[0].ssa, nir_ior(b, base, id)); ++ break; ++ } ++ case nir_intrinsic_reduce: ++ case nir_intrinsic_rotate: { ++ unsigned cluster_size = nir_intrinsic_cluster_size(intr); ++ if (cluster_size && cluster_size <= HALF_SUBGROUP_SIZE) ++ return false; ++ nir_intrinsic_set_cluster_size(intr, HALF_SUBGROUP_SIZE); ++ return true; ++ } ++ default: ++ /* Relative shuffles and quad operations never leave the half for ++ * defined results. ++ */ ++ return false; ++ } ++ ++ nir_def_replace(&intr->def, res); ++ return true; ++} ++ ++bool ++ir3_nir_lower_half_subgroups(nir_shader *nir) ++{ ++ bool progress = nir_shader_intrinsics_pass(nir, lower_half_subgroup_intrinsic, ++ nir_metadata_control_flow, NULL); ++ ++ /* Subgroup intrinsics now refer to the whole wave. */ ++ nir->info.api_subgroup_size = 2 * HALF_SUBGROUP_SIZE; ++ nir->info.max_subgroup_size = 2 * HALF_SUBGROUP_SIZE; ++ nir->info.min_subgroup_size = 2 * HALF_SUBGROUP_SIZE; ++ ++ return progress; ++} +diff --git a/src/freedreno/ir3/ir3_nir.c b/src/freedreno/ir3/ir3_nir.c +index 554cf57..0edc5ea 100644 +--- a/src/freedreno/ir3/ir3_nir.c ++++ b/src/freedreno/ir3/ir3_nir.c +@@ -1054,6 +1054,9 @@ ir3_nir_post_finalize(struct ir3_shader *shader) + * it here. Beyond this point nir_intrinsic_load_subgroup_size will return + * the "real" subgroup size. + */ ++ if (shader->options.api_wavesize == IR3_HALF_ONLY) ++ OPT(s, ir3_nir_lower_half_subgroups); ++ + unsigned subgroup_size = 0, max_subgroup_size = 0; + ir3_shader_get_subgroup_size(compiler, &shader->options, s->info.stage, + &subgroup_size, &max_subgroup_size); +diff --git a/src/freedreno/ir3/ir3_nir.h b/src/freedreno/ir3/ir3_nir.h +index 19355dd..4cb1031 100644 +--- a/src/freedreno/ir3/ir3_nir.h ++++ b/src/freedreno/ir3/ir3_nir.h +@@ -118,6 +118,7 @@ nir_def *ir3_nir_try_propagate_bit_shift(nir_builder *b, + bool ir3_nir_lower_subgroups_filter(const nir_intrinsic_instr *intrin, const void *data); + bool ir3_nir_lower_shuffle(nir_shader *nir, struct ir3_shader *shader); + bool ir3_nir_opt_subgroups(nir_shader *nir, struct ir3_shader_variant *v); ++bool ir3_nir_lower_half_subgroups(nir_shader *nir); + + nir_def *ir3_get_shared_driver_ubo(nir_builder *b, + const struct ir3_driver_ubo *ubo); +diff --git a/src/freedreno/ir3/ir3_shader.c b/src/freedreno/ir3/ir3_shader.c +index 3745a4d..d729c19 100644 +--- a/src/freedreno/ir3/ir3_shader.c ++++ b/src/freedreno/ir3/ir3_shader.c +@@ -1329,6 +1329,12 @@ ir3_shader_get_subgroup_size(const struct ir3_compiler *compiler, + case IR3_DOUBLE_ONLY: + *subgroup_size = *max_subgroup_size = compiler->info->threadsize_base * 2; + break; ++ case IR3_HALF_ONLY: ++ /* Half-wave subgroups are lowered onto single waves before anything ++ * else looks at the subgroup size. ++ */ ++ *subgroup_size = *max_subgroup_size = compiler->info->threadsize_base; ++ break; + case IR3_SINGLE_OR_DOUBLE: + /* For vertex stages, we know the wavesize will never be doubled. + * Lower subgroup_size here, to avoid having to deal with it when +diff --git a/src/freedreno/ir3/ir3_shader.h b/src/freedreno/ir3/ir3_shader.h +index d52a106..e7c6617 100644 +--- a/src/freedreno/ir3/ir3_shader.h ++++ b/src/freedreno/ir3/ir3_shader.h +@@ -140,6 +140,8 @@ enum ir3_wavesize_option { + IR3_SINGLE_ONLY, + IR3_SINGLE_OR_DOUBLE, + IR3_DOUBLE_ONLY, ++ /* Subgroups of half a single wave, emulated on single waves. */ ++ IR3_HALF_ONLY, + }; + + /** +diff --git a/src/freedreno/vulkan/tu_device.cc b/src/freedreno/vulkan/tu_device.cc +index ff079f7..ee9d4ab 100644 +--- a/src/freedreno/vulkan/tu_device.cc ++++ b/src/freedreno/vulkan/tu_device.cc +@@ -222,6 +222,16 @@ tu_subgroup_size(const struct tu_physical_device *device) + (device->expose_double_threadsize ? 2 : 1); + } + ++/* Subgroups of half a wave are emulated on top of the hardware subgroup ++ * operations. ++ */ ++static uint32_t ++tu_min_subgroup_size(const struct tu_physical_device *device) ++{ ++ return device->info->threadsize_base / ++ (device->info->props.has_getfiberid ? 2 : 1); ++} ++ + static void + get_device_extensions(const struct tu_physical_device *device, + struct vk_device_extension_table *ext) +@@ -1141,16 +1151,16 @@ static void + tu_get_physical_device_properties_1_3(struct tu_physical_device *pdevice, + struct vk_properties *p) + { +- p->minSubgroupSize = pdevice->info->threadsize_base; ++ p->minSubgroupSize = tu_min_subgroup_size(pdevice); + p->maxSubgroupSize = tu_subgroup_size(pdevice); +- p->maxComputeWorkgroupSubgroups = pdevice->info->max_waves; ++ p->maxComputeWorkgroupSubgroups = pdevice->info->max_waves * ++ pdevice->info->threadsize_base / p->minSubgroupSize; + /* Only compute and fragment shaders can run with a doubled wave size, the +- * geometry stages always run at threadsize_base. So when we expose more +- * than one possible subgroup size we can't honor a required subgroup size +- * in those stages. ++ * geometry stages always run at threadsize_base. So when we expose the ++ * doubled size we can't honor a required subgroup size in those stages. + */ + p->requiredSubgroupSizeStages = +- p->minSubgroupSize == p->maxSubgroupSize ++ !pdevice->expose_double_threadsize + ? VK_SHADER_STAGE_ALL + : (VK_SHADER_STAGE_COMPUTE_BIT | VK_SHADER_STAGE_FRAGMENT_BIT | + VK_SHADER_STAGE_TASK_BIT_EXT | VK_SHADER_STAGE_MESH_BIT_EXT); +diff --git a/src/freedreno/vulkan/tu_pipeline.cc b/src/freedreno/vulkan/tu_pipeline.cc +index 949a4e0..7343480 100644 +--- a/src/freedreno/vulkan/tu_pipeline.cc ++++ b/src/freedreno/vulkan/tu_pipeline.cc +@@ -1525,6 +1525,10 @@ tu_append_executable(struct tu_pipeline *pipeline, + struct tu_pipeline_executable exe = { + .stage = variant->type, + .stats = variant->info, ++ .subgroup_size = ++ variant->shader_options.api_wavesize == IR3_HALF_ONLY ++ ? variant->compiler->info->threadsize_base / 2 ++ : variant->info.subgroup_size, + .is_binning = variant->binning_pass, + .nir_from_spirv = nir_from_spirv, + .nir_final = ralloc_strdup(pipeline->executables_mem_ctx, variant->disasm_info.nir), +@@ -5196,7 +5200,6 @@ tu_GetPipelineExecutablePropertiesKHR( + uint32_t* pExecutableCount, + VkPipelineExecutablePropertiesKHR* pProperties) + { +- VK_FROM_HANDLE(tu_device, dev, _device); + VK_FROM_HANDLE(tu_pipeline, pipeline, pPipelineInfo->pipeline); + VK_OUTARRAY_MAKE_TYPED(VkPipelineExecutablePropertiesKHR, out, + pProperties, pExecutableCount); +@@ -5213,8 +5216,7 @@ tu_GetPipelineExecutablePropertiesKHR( + + VK_COPY_STR(props->description, _mesa_shader_stage_to_string(stage)); + +- props->subgroupSize = +- dev->compiler->info->threadsize_base * (exe->stats.double_threadsize ? 2 : 1); ++ props->subgroupSize = exe->subgroup_size; + } + } + +diff --git a/src/freedreno/vulkan/tu_pipeline.h b/src/freedreno/vulkan/tu_pipeline.h +index 31bf5c4..b3a9305 100644 +--- a/src/freedreno/vulkan/tu_pipeline.h ++++ b/src/freedreno/vulkan/tu_pipeline.h +@@ -163,6 +163,7 @@ struct tu_pipeline_executable { + mesa_shader_stage stage; + + struct ir3_info stats; ++ uint32_t subgroup_size; + bool is_binning; + + char *nir_from_spirv; +diff --git a/src/freedreno/vulkan/tu_shader.cc b/src/freedreno/vulkan/tu_shader.cc +index ad9d563..8e7194f 100644 +--- a/src/freedreno/vulkan/tu_shader.cc ++++ b/src/freedreno/vulkan/tu_shader.cc +@@ -4095,7 +4095,11 @@ tu_shader_key_subgroup_size(struct tu_shader_key *key, + struct tu_device *dev) + { + enum ir3_wavesize_option api_wavesize, real_wavesize; +- if (!dev->physical_device->expose_double_threadsize) { ++ if (subgroup_info && ++ subgroup_info->requiredSubgroupSize < dev->compiler->info->threadsize_base) { ++ api_wavesize = IR3_HALF_ONLY; ++ real_wavesize = IR3_SINGLE_ONLY; ++ } else if (!dev->physical_device->expose_double_threadsize) { + api_wavesize = IR3_SINGLE_ONLY; + real_wavesize = IR3_SINGLE_ONLY; + } else { diff --git a/linux/patches/0003-ir3-Sanitize-cube-map-directions-on-A8XX.patch b/linux/patches/0003-ir3-Sanitize-cube-map-directions-on-A8XX.patch new file mode 100644 index 0000000..8e9c957 --- /dev/null +++ b/linux/patches/0003-ir3-Sanitize-cube-map-directions-on-A8XX.patch @@ -0,0 +1,121 @@ +From 0000000000000000000000000000000000000000 Mon Sep 17 00:00:00 2001 +From: MaxsTechReview +Date: Mon, 28 Sep 2026 03:03:45 -0400 +Subject: [PATCH 3/5] ir3: Sanitize cube map directions on A8XX + +A8XX can hang when a cube map is sampled with a direction that has no +defined face projection. FINAL FANTASY VII REBIRTH does this while loading: +its lighting pass normalizes a zero vector, which gives NaN. + +Such directions give undefined results, so on A8XX replace NaN components +with zero, clamp to finite values and turn the zero vector into +Z. +--- + src/freedreno/common/freedreno_dev_info.h | 5 ++ + src/freedreno/common/freedreno_devices.py | 1 + + src/freedreno/ir3/ir3_nir.c | 58 +++++++++++++++++++++++ + 3 files changed, 64 insertions(+) + +diff --git a/src/freedreno/common/freedreno_dev_info.h b/src/freedreno/common/freedreno_dev_info.h +index 34fb9e6..c720898 100644 +--- a/src/freedreno/common/freedreno_dev_info.h ++++ b/src/freedreno/common/freedreno_dev_info.h +@@ -516,6 +516,11 @@ struct fd_dev_info { + + /* On a750+ SUBPASS_FENCE also implicitly does CCU_RESOLVE_CLEAN */ + bool subpass_fence_cleans_resolve; ++ ++ /* On a8xx sampling a cube map with a direction that has no defined ++ * face projection (NaN, the zero vector) may hang the GPU. ++ */ ++ bool cube_coord_hang_quirk; + } props; + }; + +diff --git a/src/freedreno/common/freedreno_devices.py b/src/freedreno/common/freedreno_devices.py +index 3033957..9fcc077 100644 +--- a/src/freedreno/common/freedreno_devices.py ++++ b/src/freedreno/common/freedreno_devices.py +@@ -1313,6 +1313,7 @@ a8xx_base = GPUProps( + max_texel_buffer_range_elements = (1 << 29) - 1, + max_storage_buffer_range_bytes = (1 << 31) - 1, + alias_mova_quirk = False, ++ cube_coord_hang_quirk = True, + ) + + # For a8xx, the chicken bit and most other non-ctx reg +diff --git a/src/freedreno/ir3/ir3_nir.c b/src/freedreno/ir3/ir3_nir.c +index 0edc5ea..1c50220 100644 +--- a/src/freedreno/ir3/ir3_nir.c ++++ b/src/freedreno/ir3/ir3_nir.c +@@ -717,6 +717,61 @@ ir3_nir_min_lod_workaround(nir_shader *shader) + nir_metadata_control_flow, NULL); + } + ++/* Sampling a cube map with a direction that has no defined face projection ++ * (a NaN component, the zero vector, or more than one infinite component) ++ * can hang the GPU. Such directions give undefined results, so replace NaN ++ * with zero, clamp to finite values and turn the zero vector into +Z. ++ */ ++static bool ++ir3_nir_lower_cube_coord_cb(struct nir_builder *b, nir_tex_instr *tex, ++ void *_data) ++{ ++ if (tex->sampler_dim != GLSL_SAMPLER_DIM_CUBE) ++ return false; ++ ++ int src_idx = nir_tex_instr_src_index(tex, nir_tex_src_coord); ++ if (src_idx < 0) ++ return false; ++ ++ b->cursor = nir_before_instr(&tex->instr); ++ ++ nir_def *coord = tex->src[src_idx].src.ssa; ++ unsigned bit_size = coord->bit_size; ++ nir_def *dir = nir_trim_vector(b, coord, 3); ++ nir_def *max = nir_imm_floatN_t(b, bit_size == 16 ? 65504.0 : FLT_MAX, ++ bit_size); ++ ++ dir = nir_bcsel(b, nir_fisnan(b, dir), nir_imm_zero(b, 3, bit_size), dir); ++ dir = nir_fclamp(b, dir, nir_fneg(b, max), max); ++ ++ nir_def *abs = nir_fabs(b, dir); ++ nir_def *major = nir_fmax(b, nir_fmax(b, nir_channel(b, abs, 0), ++ nir_channel(b, abs, 1)), ++ nir_channel(b, abs, 2)); ++ nir_def *z = nir_bcsel(b, nir_feq_imm(b, major, 0.0), ++ nir_imm_floatN_t(b, 1.0, bit_size), ++ nir_channel(b, dir, 2)); ++ ++ nir_def *comps[4] = { ++ nir_channel(b, dir, 0), ++ nir_channel(b, dir, 1), ++ z, ++ }; ++ if (coord->num_components > 3) ++ comps[3] = nir_channel(b, coord, 3); ++ ++ nir_src_rewrite(&tex->src[src_idx].src, ++ nir_vec(b, comps, coord->num_components)); ++ return true; ++} ++ ++static bool ++ir3_nir_lower_cube_coord(nir_shader *shader) ++{ ++ return nir_shader_tex_pass(shader, ir3_nir_lower_cube_coord_cb, ++ nir_metadata_control_flow, NULL); ++} ++ + void + ir3_finalize_nir(struct ir3_compiler *compiler, + const struct ir3_shader_nir_options *options, +@@ -787,6 +842,9 @@ ir3_finalize_nir(struct ir3_compiler *compiler, + NIR_PASS(_, s, ir3_nir_lower_sparse_residency); + NIR_PASS(_, s, ir3_nir_min_lod_workaround); + ++ if (compiler->info->props.cube_coord_hang_quirk) ++ NIR_PASS(_, s, ir3_nir_lower_cube_coord); ++ + /* for opencl kernels, TPL1_MODE_CNTL should be configured for + * isammode=CL and .arraycoordroundmode = ROUND_NEAREST_EVEN, + * as opposed to gl/vk compute shaders which follow GL rules: diff --git a/linux/patches/0004-tu-Invalidate-bindless-descriptors-through-the-A8XX-.patch b/linux/patches/0004-tu-Invalidate-bindless-descriptors-through-the-A8XX-.patch new file mode 100644 index 0000000..70cd329 --- /dev/null +++ b/linux/patches/0004-tu-Invalidate-bindless-descriptors-through-the-A8XX-.patch @@ -0,0 +1,100 @@ +From 0000000000000000000000000000000000000000 Mon Sep 17 00:00:00 2001 +From: MaxsTechReview +Date: Mon, 28 Sep 2026 03:03:46 -0400 +Subject: [PATCH 4/5] tu: Invalidate bindless descriptors through the A8XX + registers + +A8XX dropped the bindless bits from SP_UPDATE_CNTL, so writing them left +stale bindless descriptors cached after descriptor sets changed. Use +SP_GFX_BINDLESS_INVALIDATE and SP_CS_BINDLESS_INVALIDATE instead, and also +invalidate when the hardware is initialized and before blits. +--- + src/freedreno/vulkan/tu_clear_blit.cc | 2 ++ + src/freedreno/vulkan/tu_cmd_buffer.cc | 15 +++++---------- + src/freedreno/vulkan/tu_cmd_buffer.h | 20 ++++++++++++++++++++ + 3 files changed, 27 insertions(+), 10 deletions(-) + +diff --git a/src/freedreno/vulkan/tu_clear_blit.cc b/src/freedreno/vulkan/tu_clear_blit.cc +index 7f58b35..2befb7c 100644 +--- a/src/freedreno/vulkan/tu_clear_blit.cc ++++ b/src/freedreno/vulkan/tu_clear_blit.cc +@@ -1036,6 +1036,8 @@ r3d_common(struct tu_cmd_buffer *cmd, struct tu_cs *cs, enum r3d_type type, + .gfx_shared_const = true, + .cs_bindless = CHIP == A6XX ? 0x1f : 0xff, + .gfx_bindless = CHIP == A6XX ? 0x1f : 0xff,)); ++ if (CHIP >= A8XX) ++ tu_emit_bindless_invalidate(cs, true, true); + + with_crb (cs, 2 * 5 + 2 * 12) { + tu6_emit_xs_config(crb, { .vs = vs, .fs = fs }); +diff --git a/src/freedreno/vulkan/tu_cmd_buffer.cc b/src/freedreno/vulkan/tu_cmd_buffer.cc +index f87b236..00da4fe 100644 +--- a/src/freedreno/vulkan/tu_cmd_buffer.cc ++++ b/src/freedreno/vulkan/tu_cmd_buffer.cc +@@ -390,12 +390,8 @@ tu6_emit_flushes(struct tu_cmd_buffer *cmd_buffer, + tu_emit_event_write(cmd_buffer, cs, FD_CACHE_CLEAN); + if (flushes & TU_CMD_FLAG_CACHE_INVALIDATE) + tu_emit_event_write(cmd_buffer, cs, FD_CACHE_INVALIDATE); +- if (flushes & TU_CMD_FLAG_BINDLESS_DESCRIPTOR_INVALIDATE) { +- tu_cs_emit_regs(cs, SP_UPDATE_CNTL(CHIP, +- .cs_bindless = CHIP == A6XX ? 0x1f : 0xff, +- .gfx_bindless = CHIP == A6XX ? 0x1f : 0xff, +- )); +- } ++ if (flushes & TU_CMD_FLAG_BINDLESS_DESCRIPTOR_INVALIDATE) ++ tu_emit_bindless_invalidate(cs, true, true); + + /* SUBPASS_SLICE_FENCE is a weaker version of: + * - CACHE_INVALIDATE (only invalidate UCHE GMEM aperture) +@@ -2506,6 +2502,8 @@ tu_init_hw(struct tu_cmd_buffer *cmd, struct tu_cs *cs) + .gfx_shared_const = true, + .cs_bindless = CHIP == A6XX ? 0x1f : 0xff, + .gfx_bindless = CHIP == A6XX ? 0x1f : 0xff,)); ++ if (CHIP >= A8XX) ++ tu_emit_bindless_invalidate(cs, true, true); + + tu_cs_emit_wfi(cs); + +@@ -4818,10 +4816,7 @@ tu6_emit_bindless_bases(struct tu_cmd_buffer *cmd, struct tu_cs *cs, + } + } + +- tu_cs_emit_regs(cs, SP_UPDATE_CNTL(CHIP, +- .cs_bindless = compute ? CHIP == A6XX ? 0x1f : 0xff : 0, +- .gfx_bindless = !compute ? CHIP == A6XX ? 0x1f : 0xff : 0, +- )); ++ tu_emit_bindless_invalidate(cs, !compute, compute); + } + + template +diff --git a/src/freedreno/vulkan/tu_cmd_buffer.h b/src/freedreno/vulkan/tu_cmd_buffer.h +index 3372fb3..b8c72be 100644 +--- a/src/freedreno/vulkan/tu_cmd_buffer.h ++++ b/src/freedreno/vulkan/tu_cmd_buffer.h +@@ -901,6 +901,26 @@ void tu6_emit_blit_scissor(struct tu_cmd_buffer *cmd, struct tu_cs *cs, + + void tu_disable_draw_states(struct tu_cmd_buffer *cmd, struct tu_cs *cs); + ++/* A8XX dropped the bindless bits from SP_UPDATE_CNTL, the bindless descriptor ++ * caches are invalidated through dedicated registers instead. ++ */ ++template ++static inline void ++tu_emit_bindless_invalidate(struct tu_cs *cs, bool gfx, bool compute) ++{ ++ if (CHIP >= A8XX) { ++ if (gfx) ++ tu_cs_emit_regs(cs, A6XX_SP_GFX_BINDLESS_INVALIDATE(.dword = 1)); ++ if (compute) ++ tu_cs_emit_regs(cs, A6XX_SP_CS_BINDLESS_INVALIDATE(.dword = 1)); ++ } else { ++ tu_cs_emit_regs(cs, SP_UPDATE_CNTL(CHIP, ++ .cs_bindless = compute ? CHIP == A6XX ? 0x1f : 0xff : 0, ++ .gfx_bindless = gfx ? CHIP == A6XX ? 0x1f : 0xff : 0, ++ )); ++ } ++} ++ + void tu6_apply_depth_bounds_workaround(struct tu_device *device, + uint32_t *rb_depth_cntl); + diff --git a/linux/patches/0005-tu-kgsl-Fetch-A8XX-command-streams-through-a-virtual.patch b/linux/patches/0005-tu-kgsl-Fetch-A8XX-command-streams-through-a-virtual.patch new file mode 100644 index 0000000..323eb49 --- /dev/null +++ b/linux/patches/0005-tu-kgsl-Fetch-A8XX-command-streams-through-a-virtual.patch @@ -0,0 +1,153 @@ +From 0000000000000000000000000000000000000000 Mon Sep 17 00:00:00 2001 +From: MaxsTechReview +Date: Mon, 28 Sep 2026 03:03:46 -0400 +Subject: [PATCH 5/5] tu/kgsl: Fetch A8XX command streams through a virtual BO + alias + +On A8XX, freeing command stream memory with GPUMEM_FREE_ID leads to GPU +hangs, even when the GPU is idle. FINAL FANTASY VII REBIRTH hits this within +seconds to minutes of play. + +Map every BO that holds IBs through a VBO alias and hand the GPU the alias +address. Before the BO is freed, unbind the alias range, then free the +alias. +--- + src/freedreno/vulkan/tu_knl.h | 3 + + src/freedreno/vulkan/tu_knl_kgsl.cc | 96 +++++++++++++++++++++++++++++ + 2 files changed, 99 insertions(+) + +diff --git a/src/freedreno/vulkan/tu_knl.h b/src/freedreno/vulkan/tu_knl.h +index 27f035f..6e72ffa 100644 +--- a/src/freedreno/vulkan/tu_knl.h ++++ b/src/freedreno/vulkan/tu_knl.h +@@ -57,6 +57,9 @@ struct tu_bo { + * export file descriptor. + */ + int shared_fd; ++ ++ /* Virtual BO the GPU accesses this BO through, or 0 if none. */ ++ uint32_t alias_handle; + #endif + + bool implicit_sync : 1; +diff --git a/src/freedreno/vulkan/tu_knl_kgsl.cc b/src/freedreno/vulkan/tu_knl_kgsl.cc +index a999807..d76a9f9 100644 +--- a/src/freedreno/vulkan/tu_knl_kgsl.cc ++++ b/src/freedreno/vulkan/tu_knl_kgsl.cc +@@ -295,6 +295,95 @@ kgsl_align_flag(uint64_t align) + return 0; + } + ++/* On A8XX, freeing IB memory with a plain GPUMEM_FREE_ID leads to GPU hangs, ++ * even when the GPU is idle. Fetching IBs through a virtual BO alias whose ++ * range is unbound before the memory is freed avoids them. ++ */ ++static bool ++kgsl_bo_needs_alias(struct tu_device *dev, enum tu_bo_alloc_flags flags) ++{ ++ return (flags & TU_BO_ALLOC_NO_32B_ROLLOVER) && ++ dev->physical_device->has_sparse && ++ dev->physical_device->info->chip >= A8XX; ++} ++ ++static VkResult ++kgsl_bo_create_alias(struct tu_device *dev, struct tu_bo *bo, uint64_t align) ++{ ++ int fd = dev->physical_device->local_fd; ++ ++ struct kgsl_gpuobj_alloc alloc = { ++ .size = bo->size, ++ .flags = KGSL_MEMFLAGS_VBO | KGSL_MEMFLAGS_VBO_NO_MAP_ZERO | ++ kgsl_align_flag(align), ++ }; ++ ++ if (safe_ioctl(fd, IOCTL_KGSL_GPUOBJ_ALLOC, &alloc)) { ++ return vk_errorf(dev, VK_ERROR_OUT_OF_DEVICE_MEMORY, ++ "GPUOBJ_ALLOC failed (%s)", strerror(errno)); ++ } ++ ++ struct kgsl_gpumem_get_info info = { ++ .id = alloc.id, ++ }; ++ ++ struct kgsl_gpumem_bind_range range = { ++ .length = bo->size, ++ .child_id = bo->gem_handle, ++ .op = KGSL_GPUMEM_RANGE_OP_BIND, ++ }; ++ ++ struct kgsl_gpumem_bind_ranges bind = { ++ .ranges = (uint64_t)(uintptr_t)&range, ++ .ranges_nents = 1, ++ .ranges_size = sizeof(range), ++ .id = alloc.id, ++ }; ++ ++ if (safe_ioctl(fd, IOCTL_KGSL_GPUMEM_GET_INFO, &info) || ++ safe_ioctl(fd, IOCTL_KGSL_GPUMEM_BIND_RANGES, &bind)) { ++ VkResult result = ++ vk_errorf(dev, VK_ERROR_OUT_OF_DEVICE_MEMORY, ++ "binding IB alias failed (%s)", strerror(errno)); ++ ++ struct kgsl_gpumem_free_id req = { ++ .id = alloc.id, ++ }; ++ safe_ioctl(fd, IOCTL_KGSL_GPUMEM_FREE_ID, &req); ++ return result; ++ } ++ ++ bo->alias_handle = alloc.id; ++ bo->iova = info.gpuaddr; ++ return VK_SUCCESS; ++} ++ ++static void ++kgsl_bo_destroy_alias(struct tu_device *dev, struct tu_bo *bo) ++{ ++ int fd = dev->physical_device->local_fd; ++ ++ struct kgsl_gpumem_bind_range range = { ++ .length = bo->size, ++ .child_id = bo->gem_handle, ++ .op = KGSL_GPUMEM_RANGE_OP_UNBIND, ++ }; ++ ++ struct kgsl_gpumem_bind_ranges unbind = { ++ .ranges = (uint64_t)(uintptr_t)&range, ++ .ranges_nents = 1, ++ .ranges_size = sizeof(range), ++ .id = bo->alias_handle, ++ }; ++ ++ safe_ioctl(fd, IOCTL_KGSL_GPUMEM_BIND_RANGES, &unbind); ++ ++ struct kgsl_gpumem_free_id req = { ++ .id = bo->alias_handle, ++ }; ++ safe_ioctl(fd, IOCTL_KGSL_GPUMEM_FREE_ID, &req); ++} ++ + static VkResult + kgsl_bo_init(struct tu_device *dev, + struct vk_object_base *base, +@@ -386,6 +475,10 @@ kgsl_bo_init(struct tu_device *dev, + result = kgsl_sparse_vma_map(dev, lazy_vma, bo, 0); + } else if (flags & TU_BO_ALLOC_REPLAYABLE) { + result = kgsl_bo_user_map(dev, bo, client_iova); ++ } else if (kgsl_bo_needs_alias(dev, flags)) { ++ result = kgsl_bo_create_alias(dev, bo, align); ++ if (result != VK_SUCCESS) ++ kgsl_bo_finish(dev, bo); + } + + if (result != VK_SUCCESS) +@@ -522,6 +615,9 @@ kgsl_bo_finish(struct tu_device *dev, struct tu_bo *bo) + + TU_RMV(bo_destroy, dev, bo); + ++ if (bo->alias_handle) ++ kgsl_bo_destroy_alias(dev, bo); ++ + struct kgsl_gpumem_free_id req = { + .id = bo->gem_handle + }; diff --git a/linux/release_version.py b/linux/release_version.py new file mode 100644 index 0000000..3ca4d18 --- /dev/null +++ b/linux/release_version.py @@ -0,0 +1,20 @@ +#!/usr/bin/env python3 +import json +import re +import sys + + +def next_version(releases): + versions = [] + for release in releases: + match = re.fullmatch(r'linux-v(\d+)\.(\d+)\.(\d+)', release['tag_name']) + if match and not release['draft'] and not release['prerelease']: + versions.append(tuple(map(int, match.groups()))) + if not versions: + return '0.1.0' + major, minor, patch = max(versions) + return '.'.join(map(str, max((0, 1, 0), (major, minor, patch + 1)))) + + +if __name__ == '__main__': + print(next_version([release for page in json.load(sys.stdin) for release in page])) diff --git a/linux/scheduler/schedule-linux-turnip.yml b/linux/scheduler/schedule-linux-turnip.yml new file mode 100644 index 0000000..23da3ec --- /dev/null +++ b/linux/scheduler/schedule-linux-turnip.yml @@ -0,0 +1,21 @@ +name: Weekly Linux Turnip + +on: + schedule: + - cron: '0 12 * * 3' + workflow_dispatch: + inputs: + publish: + description: Publish the Linux driver release + type: boolean + default: false + +permissions: + contents: write + +jobs: + linux: + uses: WinNative-Emu/Drivers/.github/workflows/build-linux-turnip.yml@feature/linux-drivers + with: + source_ref: feature/linux-drivers + publish: ${{ github.event_name == 'schedule' || inputs.publish }} diff --git a/linux/test_release_version.py b/linux/test_release_version.py new file mode 100644 index 0000000..45f5b77 --- /dev/null +++ b/linux/test_release_version.py @@ -0,0 +1,28 @@ +import unittest +from release_version import next_version + + +class ReleaseVersionTest(unittest.TestCase): + def release(self, tag, draft=False, prerelease=False): + return dict(tag_name=tag, draft=draft, prerelease=prerelease) + + def test_initial_version_ignores_android_and_drafts(self): + self.assertEqual(next_version([]), '0.1.0') + self.assertEqual(next_version([ + self.release('v1.16'), + self.release('linux-v0.1.0', draft=True), + self.release('linux-v0.2.0', prerelease=True), + self.release('linux-v20260921'), + ]), '0.1.0') + + def test_published_versions_sort_numerically(self): + self.assertEqual(next_version([ + self.release('linux-v0.1.9'), + self.release('linux-v0.1.10'), + self.release('linux-v0.1.11', draft=True), + ]), '0.1.11') + self.assertEqual(next_version([self.release('linux-v1.2.3')]), '1.2.4') + + +if __name__ == '__main__': + unittest.main() diff --git a/linux/verify_packages.py b/linux/verify_packages.py new file mode 100644 index 0000000..3dc6326 --- /dev/null +++ b/linux/verify_packages.py @@ -0,0 +1,41 @@ +#!/usr/bin/env python3 +import hashlib +import json +from pathlib import Path +import subprocess +import sys +import tempfile +import zipfile + +hashes = set() +commits = set() +for value in sys.argv[1:]: + archive = Path(value if value.endswith('.zip') else value + '_Axxx.zip') + with zipfile.ZipFile(archive) as package: + meta = json.loads(package.read('meta.json')) + library = package.read(meta['libraryName']) + if meta['name'] != f'WN Linux Turnip {meta["driverVersion"]}' or not meta['driverVersion'].endswith('-' + meta['variant']): + raise SystemExit('Incorrect driver label') + if meta['platform'] != 'linux' or meta['architecture'] != 'aarch64' or meta['libc'] != 'glibc': + raise SystemExit('Incorrect package platform') + digest = hashlib.sha256(library).hexdigest() + if digest != meta['librarySha256'] or digest in hashes: + raise SystemExit('Invalid or identical variant libraries') + hashes.add(digest) + commits.add(meta['mesaCommit']) + with tempfile.NamedTemporaryFile() as file: + file.write(library) + file.flush() + header = subprocess.check_output(['readelf', '-h', file.name], text=True) + dynamic = subprocess.check_output(['readelf', '-d', file.name], text=True) + if 'AArch64' not in header or '[libc.so.6]' not in dynamic or '[libc.so]' in dynamic: + raise SystemExit('Incorrect ELF ABI') + if b'vkCreateWaylandSurfaceKHR' not in library or b'vkCreateXcbSurfaceKHR' not in library: + raise SystemExit('Missing Linux WSI') + if b'drirc.d' in library: + raise SystemExit('Driver defaults read from disk instead of built in') + if (b'Failed to set initial PWR_MAX constraint' in library) != (meta['variant'] == 'p'): + raise SystemExit('Incorrect power variant') + print(f'{archive.name}: ARM64 glibc, Wayland/X11, {meta["variant"]}, Mesa {meta["mesaCommit"]}') +if len(commits) != 1: + raise SystemExit('Variants use different Mesa commits') diff --git a/linux/verify_patches.py b/linux/verify_patches.py new file mode 100644 index 0000000..6000376 --- /dev/null +++ b/linux/verify_patches.py @@ -0,0 +1,37 @@ +#!/usr/bin/env python3 +from pathlib import Path +import sys + +source, variant, logfile = sys.argv[1:] +log = Path(logfile).read_text() +for line in log.splitlines(): + if 'drawcall anchor absent' in line: + continue + if any(word in line.lower() for word in ('warning:', 'anchor absent', 'anchor not', 'skipping', 'fatal:')): + raise SystemExit(f'Patch did not apply: {line}') +root = Path(source) +kgsl = (root / 'src/freedreno/vulkan/tu_knl_kgsl.cc').read_text() +autotune = (root / 'src/freedreno/vulkan/tu_autotune.cc').read_text() +device = (root / 'src/freedreno/vulkan/tu_device.cc').read_text() +if not (root / 'src/freedreno/vulkan/tu_mesh.cc').exists() or '.EXT_mesh_shader = tu_has_mesh_shader(device)' not in device: + raise SystemExit('Missing mesh shader emulation') +ir3_nir = (root / 'src/freedreno/ir3/ir3_nir.c').read_text() +if 'ir3_nir_lower_half_subgroups' not in ir3_nir: + raise SystemExit('Missing half-wave subgroups') +if 'ir3_nir_lower_cube_coord' not in ir3_nir: + raise SystemExit('Missing cube coordinate sanitizing') +if 'SP_GFX_BINDLESS_INVALIDATE' not in (root / 'src/freedreno/vulkan/tu_cmd_buffer.h').read_text(): + raise SystemExit('Missing A8XX bindless invalidation') +if 'kgsl_bo_create_alias' not in kgsl: + raise SystemExit('Missing command stream BO alias') +checks = ['has_local = device->has_master = true'] +if variant == 'p': + checks += ['wnturnip_set_pwr_max_constraint(', 'count % 1000 == 0', 'KGSL_CONTEXT_PWR_CONSTRAINT', 'KGSL_CMDBATCH_PWR_CONSTRAINT'] +elif 'wnturnip_set_pwr_max_constraint(' in kgsl: + raise SystemExit('Balanced driver contains performance power constraints') +for marker in checks: + if marker not in kgsl: + raise SystemExit(f'Missing patch marker: {marker}') +if variant == 'b' and 'gmem_bandwidth * 10 + total_draw_call_bandwidth' not in autotune: + raise SystemExit('Missing balanced autotuner') +print(f'{variant}: Linux and variant patches verified')