From 629bc35e9dcc693a3435c7840f2b8387fa25ba7e Mon Sep 17 00:00:00 2001 From: ayo <93222141+countervolts@users.noreply.github.com> Date: Sun, 27 Sep 2026 16:53:22 -0400 Subject: [PATCH] ZLUDA: read null texture objects as zero (fixes Ghost of Tsushima GPU fault) --- README.md | 2 +- docs/building.md | 5 ++- packaging/RELEASE_NOTES.md | 2 +- packaging/SOURCES.txt | 2 +- .../0005-null-texture-objects-read-zero.patch | 40 +++++++++++++++++++ 5 files changed, 46 insertions(+), 5 deletions(-) create mode 100644 patches/zluda/0005-null-texture-objects-read-zero.patch diff --git a/README.md b/README.md index decea6f..ce08a92 100644 --- a/README.md +++ b/README.md @@ -78,7 +78,7 @@ The only other requirement is ROCm's HIP runtime 7.x, installed from your distri The full sequence is in [docs/building.md](docs/building.md). In short: -1. Build ZLUDA `ee2f25a` with `patches/zluda/0002`–`0004` applied. +1. Build ZLUDA `ee2f25a` with `patches/zluda/0002`–`0005` applied. 2. Build the patched vkd3d-proton: `scripts/build_vkd3d_proton_d4r.sh OUT_DIR`. 3. Build the shim and bridge: `scripts/build_d4r_nvngx_shim.sh`, `scripts/build_wine_nvcuda_bridge.sh`. 4. Stage the runtime with your NVIDIA files: `scripts/install_d4r_runtime.sh _nvngx.dll nvngx_dlss.dll`. diff --git a/docs/building.md b/docs/building.md index 8b45ef1..37d9176 100644 --- a/docs/building.md +++ b/docs/building.md @@ -4,7 +4,7 @@ These steps produce the pieces the launcher needs: a patched ZLUDA, a patched vk | Component | Tested version | |---|---| -| ZLUDA | `ee2f25a` (upstream), plus `patches/zluda/0002`, `0003`, `0004` in that order | +| ZLUDA | `ee2f25a` (upstream), plus `patches/zluda/0002`, `0003`, `0004`, `0005` in that order | | vkd3d-proton | `3dfc6f07` (the base GE-Proton11-3 ships), plus `patches/vkd3d-proton/0001` | | ROCm | 7.2 (HIP runtime, clang, device libraries) | | Proton | GE-Proton11-3 (with its OptiScaler integration) | @@ -17,7 +17,7 @@ git clone https://github.com/vosen/ZLUDA zluda && cd zluda git checkout ee2f25a git submodule update --init --recursive git lfs pull -for p in 0002 0003 0004; do git apply /path/to/d4r/patches/zluda/$p-*.patch; done +for p in 0002 0003 0004 0005; do git apply /path/to/d4r/patches/zluda/$p-*.patch; done # rebuild the device helpers the patches changed (ptx/lib/zluda_ptx_impl*.bc) ZLUDA_SOURCE_ROOT=$PWD ROCM_ROOT=/opt/rocm /path/to/d4r/scripts/build_zluda_ptx_helpers.sh LIBRARY_PATH=/opt/rocm/lib cargo build --release -p zluda @@ -33,6 +33,7 @@ What the patches add: - **0002**: the CUDA driver API surface NGX and DLSS use (arrays, textures, surfaces, 1010102 formats, launch parameter buffers), the PTX features DLSS kernels need, and module dumps for debugging (`D4R_ZLUDA_DUMP_DIR`). - **0003**: an implicit 256-thread launch bound for kernels without PTX bounds (removes massive register spilling, `D4R_ZLUDA_IMPLICIT_MAX_BLOCK`) and f16 tensor-core MMA on RDNA3 WMMA (`D4R_ZLUDA_WMMA`). - **0004**: the native kernel override hook (`D4R_ZLUDA_NATIVE_DIR`, see [native-kernels.md](native-kernels.md)), FP8 MMA on WMMA (`D4R_ZLUDA_WMMA_FP8`), optional elision of per-instruction denormal mode switches (`D4R_ZLUDA_IGNORE_DENORMAL`), inlined image helpers, and linking extra bitcode into a PTX module (`D4R_ZLUDA_EXTRA_BC`, used by the texture-kernel build). +- **0005**: a null texture object (CUDA handle 0) reads as zeros, as on NVIDIA GPUs, instead of faulting the GPU. DLSS samples absent optional inputs that way in some configurations (low-resolution motion vectors without HDR, as in Ghost of Tsushima). ## 2. vkd3d-proton diff --git a/packaging/RELEASE_NOTES.md b/packaging/RELEASE_NOTES.md index a08e18d..2e63e91 100644 --- a/packaging/RELEASE_NOTES.md +++ b/packaging/RELEASE_NOTES.md @@ -60,4 +60,4 @@ d4r is Apache 2.0. The zip also contains ZLUDA (Apache 2.0 or MIT), vkd3d-proton NVIDIA's files in the zip, and the kernels built from NVIDIA's code, belong to NVIDIA and are not covered by any of these licenses. d4r is not affiliated with NVIDIA, AMD or the OptiScaler project. -**SHA-256** `d4r-0.1.0.zip`: `c26670538cdd122f9766bb30c6099c0636e4f8365f6ef82109c3671eea2f8564` +**SHA-256** `d4r-0.1.0.zip`: `ddaa0866273c6e31f2f1282b2cc2ac370c804b04e8895cdda27c096747edf5a1` diff --git a/packaging/SOURCES.txt b/packaging/SOURCES.txt index cfaf878..8489644 100644 --- a/packaging/SOURCES.txt +++ b/packaging/SOURCES.txt @@ -10,7 +10,7 @@ Built from source by d4r (Apache License 2.0, d4r/licenses/d4r-LICENSE.txt): Third-party software, unmodified or patched by d4r: d4r/zluda/libcuda.so ZLUDA (Apache License 2.0 or MIT), https://github.com/vosen/ZLUDA - commit @ZLUDA_COMMIT@ with d4r/source/patches/zluda/0002, 0003 and 0004 applied. + commit @ZLUDA_COMMIT@ with d4r/source/patches/zluda/0002 to 0005 applied. It contains LLVM (Apache License 2.0 with LLVM exceptions). Licenses: d4r/licenses/ZLUDA-LICENSE-*.txt, d4r/licenses/LLVM-LICENSE.txt d3d12.dll, d3d12core.dll diff --git a/patches/zluda/0005-null-texture-objects-read-zero.patch b/patches/zluda/0005-null-texture-objects-read-zero.patch new file mode 100644 index 0000000..9081d2a --- /dev/null +++ b/patches/zluda/0005-null-texture-objects-read-zero.patch @@ -0,0 +1,40 @@ +From: d4r +Subject: [PATCH] null texture objects read as zero + +CUDA texture object 0 reads as zeros on NVIDIA GPUs, and DLSS samples absent optional inputs +(its exposure scales, with the flags Ghost of Tsushima uses) that way. A ZLUDA texture object is a +descriptor pointer, so handle 0 made the GPU read address 0 and fault. Handle 0 now points at an +all-zero descriptor, which AMD hardware treats as a null resource. Rebuild ptx/lib/*.bc after +applying (scripts/build_zluda_ptx_helpers.sh). +--- +--- a/ptx/lib/zluda_ptx_impl.cpp 2026-09-27 16:47:02.815623944 -0400 ++++ b/ptx/lib/zluda_ptx_impl.cpp 2026-09-27 16:47:02.815623944 -0400 +@@ -1725,9 +1725,18 @@ + c; + } + ++ // d4r: CUDA texture object 0 reads as zeros on NVIDIA GPUs, and DLSS samples absent optional inputs ++ // (exposure scales) that way. A ZLUDA texture object is a descriptor pointer, so 0 is redirected to ++ // an all-zero descriptor, which AMD hardware treats as a null resource (reads return 0). ++ __device__ static unsigned int d4r_null_texture[32]; ++ static uint64_t texture_or_null(uint64_t texobj) __device__ ++ { ++ return texobj != 0 ? texobj : uint64_t(uintptr_t(d4r_null_texture)); ++ } ++ + static std::pair get_image_and_sampler(uint64_t texobj) __device__ + { +- unsigned int GLOBAL_SPACE *image = (unsigned int GLOBAL_SPACE *)texobj; ++ unsigned int GLOBAL_SPACE *image = (unsigned int GLOBAL_SPACE *)texture_or_null(texobj); + unsigned int GLOBAL_SPACE *sampler = image + HIP_SAMPLER_OBJECT_OFFSET_DWORD; + return {image, sampler}; + } +@@ -1868,7 +1877,7 @@ + #define tex_2d_gather4(COMPONENT) \ + v4f32 FUNC(texobj_gather4##COMPONENT##_2d_v4_f32_f32)(uint64_t texobj, v2f32 coord) \ + { \ +- auto image = reinterpret_cast(static_cast(texobj)); \ ++ auto image = reinterpret_cast(static_cast(texture_or_null(texobj))); \ + auto sampler = image + HIP_SAMPLER_OBJECT_OFFSET_DWORD; \ + float2 coordinates{coord.x, coord.y}; \ + auto result = __ockl_image_gather4##COMPONENT##_2D(image, sampler, get_native_vector(coordinates)); \