diff --git a/CMakeLists.txt b/CMakeLists.txt index b992245..8ef40c8 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -21,6 +21,14 @@ add_custom_command(OUTPUT "${probe_object}" -o "${probe_object}" DEPENDS tools/windows/probe_kernel.hip VERBATIM) add_custom_target(d4r_probe_code_object ALL DEPENDS "${probe_object}") +set(pixel_object "${CMAKE_CURRENT_BINARY_DIR}/pixel_convert_gfx1201.hsaco") +add_custom_command(OUTPUT "${pixel_object}" + COMMAND "${D4R_HIP_CLANG}" -x hip -std=c++17 --offload-arch=gfx1201 + --offload-device-only --no-gpu-bundle-output -mno-wavefrontsize64 + -nogpuinc -nogpulib -O2 "${CMAKE_CURRENT_SOURCE_DIR}/tools/windows/pixel_convert_gfx1201.hip" + -o "${pixel_object}" + DEPENDS tools/windows/pixel_convert_gfx1201.hip tools/windows/pixel_format.h kernels/common/hip_device_minimal.h VERBATIM) +add_custom_target(d4r_pixel_code_object ALL DEPENDS "${pixel_object}") set(wmma_object "${CMAKE_CURRENT_BINARY_DIR}/wmma_gfx1201.hsaco") add_custom_command(OUTPUT "${wmma_object}" COMMAND "${D4R_HIP_CLANG}" -x hip --offload-arch=gfx1201 @@ -68,7 +76,7 @@ foreach(family k m) endforeach() endforeach() add_custom_target(d4r_native_gfx1201_compile ALL DEPENDS ${native_objects}) -foreach(probe hip_gfx1201_probe gfx12_wmma_probe k_module_probe native_replay_probe cuda_driver_probe cuda_mma_probe cuda_image_probe d3d12_command_probe d3d12_hip_interop_probe hip_stream_lifecycle_probe ngx_cuda_init_probe d3d12_dlss_probe) +foreach(probe hip_gfx1201_probe gfx12_wmma_probe k_module_probe native_replay_probe cuda_driver_probe cuda_mma_probe cuda_image_probe d3d12_command_probe d3d12_pixel_probe d3d12_hip_interop_probe hip_stream_lifecycle_probe ngx_cuda_init_probe d3d12_dlss_probe) add_executable(d4r_${probe} tools/windows/${probe}.cpp) target_compile_definitions(d4r_${probe} PRIVATE WIN32_LEAN_AND_MEAN NOMINMAX __HIP_PLATFORM_AMD__) target_include_directories(d4r_${probe} PRIVATE "${D4R_HIP_ROOT}/include") @@ -90,6 +98,7 @@ if(NOT EXISTS "${D4R_MINHOOK_ROOT}/include/MinHook.h") endif() add_subdirectory("${D4R_MINHOOK_ROOT}" "${CMAKE_CURRENT_BINARY_DIR}/minhook" EXCLUDE_FROM_ALL) add_library(d4r_nvngx_windows SHARED tools/windows/d4r_nvngx_windows.cpp tools/windows/d3d12_command_hooks.cpp) +add_dependencies(d4r_nvngx_windows d4r_pixel_code_object) set_target_properties(d4r_nvngx_windows PROPERTIES PREFIX "" OUTPUT_NAME "d4r_nvngx") target_compile_definitions(d4r_nvngx_windows PRIVATE WIN32_LEAN_AND_MEAN NOMINMAX __HIP_PLATFORM_AMD__) target_include_directories(d4r_nvngx_windows PRIVATE "${D4R_HIP_ROOT}/include") @@ -139,6 +148,8 @@ install(TARGETS d4r_nvapi_compat RUNTIME DESTINATION nvapi-compat) target_link_libraries(d4r_d3d12_hip_interop_probe PRIVATE d3d12 dxgi dxguid) target_link_libraries(d4r_d3d12_dlss_probe PRIVATE d3d12 dxgi dxguid) target_link_libraries(d4r_d3d12_command_probe PRIVATE d3d12 dxgi dxguid) +target_link_libraries(d4r_d3d12_pixel_probe PRIVATE d3d12 dxgi dxguid) +add_dependencies(d4r_d3d12_pixel_probe d4r_pixel_code_object) add_dependencies(d4r_d3d12_command_probe d4r_nvngx_windows) set(D4R_AGILITY_ROOT "" CACHE PATH "Optional unpacked Microsoft.Direct3D.D3D12 1.619.5 debug SDK") if(EXISTS "${D4R_AGILITY_ROOT}/build/native/bin/x64/d3d12SDKLayers.dll") @@ -158,6 +169,7 @@ add_dependencies(d4r_hip_gfx1201_probe d4r_probe_code_object) add_dependencies(d4r_gfx12_wmma_probe d4r_wmma_code_object) add_dependencies(d4r_k_module_probe d4r_native_gfx1201_compile) install(FILES "${probe_object}" DESTINATION bin) +install(FILES "${pixel_object}" DESTINATION bin) install(FILES "${wmma_object}" DESTINATION bin) install(FILES scripts/windows/test-windows-rdna4.ps1 DESTINATION .) install(FILES scripts/windows/test-k-replay.ps1 DESTINATION .) @@ -177,6 +189,9 @@ install(FILES kernels/tools/pwin_model.py kernels/tools/pwin_check.py kernels/to install(FILES kernels/tools/model_enc3.py kernels/tools/swin_model.py kernels/tools/swin_check.py DESTINATION reference) install(FILES docs/windows-rdna4-port.md DESTINATION docs) include(CTest) +add_test(NAME d3d12_gpu_pixel_formats COMMAND d4r_d3d12_pixel_probe + --hip-root "${D4R_HIP_ROOT}" --module "${pixel_object}" --iterations 2) +set_tests_properties(d3d12_gpu_pixel_formats PROPERTIES TIMEOUT 120 LABELS "hardware;gfx1201;formats") add_test(NAME d3d12_command_backend COMMAND d4r_d3d12_command_probe --hip-root "${D4R_HIP_ROOT}" --module "$" --interop-mode indirect --iterations 32) set_tests_properties(d3d12_command_backend PROPERTIES TIMEOUT 120 LABELS "hardware;d3d12;command_lists;lifetime") diff --git a/README.md b/README.md index 6e72f24..544fdb2 100644 --- a/README.md +++ b/README.md @@ -6,16 +6,17 @@ d4r runs NVIDIA's official DLSS Super Resolution library (`nvngx_dlss.dll`) in W **Proof of concept:** d4r shows that DLSS can run on an AMD GPU, but it is not really that practical for everyday use yet. It has been tested on one GPU in a handful of games and depends on unreleased patches to ZLUDA and vkd3d-proton. -**Windows / RDNA4 development:** the `windows-rdna4` branch adds native Windows -gfx1201 HIP/ZLUDA diagnostics, validated D3D12/HIP interop, an NGX init probe, -and a CPU-checked gfx12 WMMA backend. K `enc1` and `enc2` execute on -gfx1201 and match the existing numpy model on synthetic nonzero fixtures. -This branch incorporates upstream RDNA4 commit `dbef4b2`; native Windows CMake -builds all 11 K and 5 M modules with the upstream gfx12 layout. Other layers, -full K/M image validation and the Windows game runtime are still in progress. -Build commands, hardware results and remaining -porting gates are in [docs/windows-rdna4-port.md](docs/windows-rdna4-port.md). -This branch does not yet provide a Windows DLSS runtime. +**Windows / RDNA4 development:** `windows-rdna4` runs native K and M networks +in a standalone Windows D3D12 harness on an RX 9070 XT (gfx1201). All K layers +and all spatial blocks in forty temporal M launches pass NumPy/PTX validation. +The Windows shim imports D3D12 VRAM through HIP, executes DLSS at command-list +submission, and returns output in the same frame without CPU image copies. +RGBA16/32F, R11G11B10, RGBA/BGRA8, depth, exposure and motion formats have +independent D3D12 typed SRV/UAV conversion tests. Four-frame K/M runs with packed +output match the FP16 baseline exactly after resource-format quantisation. +The build incorporates upstream RDNA4 commit `dbef4b2`. OptiScaler/game +integration and performance profiling remain in progress; no game support is +claimed yet. See [build instructions and hardware results](docs/windows-rdna4-port.md). See [supported games](SUPPORTED_GAMES.md) for the tested games and DLSS models. diff --git a/docs/windows-rdna4-port.md b/docs/windows-rdna4-port.md index 699fcea..e20ec38 100644 --- a/docs/windows-rdna4-port.md +++ b/docs/windows-rdna4-port.md @@ -2,7 +2,8 @@ Branch: `windows-rdna4`. Target: Windows 11 x64, RX 9070 XT, **gfx1201**. Priority: correct K, correct M, native Windows, same-frame output, then speed. -This is a work log for an implementation in progress, not a claim of DLSS support. +This is a work log for an implementation in progress. Standalone DLSS K/M +works on the real Windows GPU; OptiScaler/game support remains unverified. ## Milestone and gates @@ -37,6 +38,67 @@ Do not integrate NGX until the integer PTX workload is stable on the real GPU. | M8: standalone Windows D3D12 NGX harness | PASS for ordinary recorded EvaluateFeature and explicit boundaries: K/M, four frames each, no shim CPU image copies; both bit-exact against explicit boundary baselines | | M9: OptiScaler integration, profiling, installation | Pending | +## Native GPU resource-format conversion (2026-10-01) + +`pixel_convert_gfx1201.hsaco` converts imported D3D12 copy footprints to NGX's +RGBA16F/RG16F/R32F arrays and converts output back to the original resource +format. Canonical formats retain the existing direct path. Other formats use +pitched HIP allocations and native decode/encode kernels; images stay in VRAM. +Completion currently synchronizes the host before NGX uses its other streams. +The module is installed beside `d4r_nvngx.dll`; `D4R_FORMAT_MODULE` can override +its absolute path. The diagnostic runner copies it into its owned temporary +NGX runtime alongside the shim. No NVIDIA binaries are installed in packages. + +Supported storage: RGBA16/32F, RGBA/BGRA8 UNORM, R11G11B10 FLOAT, RGB10A2 UNORM, +RG16/32F, RG16 SNORM/UNORM, R16F exposure, R32F/D32 depth, D24/R24 depth and +D16/R16 UNORM depth. Typeless R16 means floating-point exposure or UNORM depth, +according to the NGX plane. sRGB and D32S8 conversion are explicitly rejected +until independently validated. Optional mask textures are not supported yet. + +The public `d4r_d3d12_pixel_probe` tests nineteen plane/format combinations at +259x17, with padded pitches, partial workgroups and row guards. D3D12 typed SRV +loads provide independent float32 decode values, with a mathematical binary16 +nearest-value search for CUDA RNE conversion. HLSL `f32tof16` is not used as +an RNE reference: the current driver truncates that operation. Color encode +is compared bit-for-bit to typed UAV stores read through a D3D12 consumer SRV. +The reference includes HDR limits, subnormals, signed motion endpoints, UNORM +rounding/clamping and packed channel permutations. An initial RNE R11 encoder +failed this test; truncation now matches the driver, including finite HDR limits. +See [Microsoft's format-conversion specification](https://microsoft.github.io/DirectX-Specs/d3d/archive/D3D11_3_FunctionalSpec.htm). + +Results on RX 9070 XT / TheRock: + +* `pixel-final-ctest.log`: 14/14 PASS (13 hardware tests and one WARP test). +* `pixel-milestone-diagnostics.zip`: all fourteen public diagnostic gates PASS. +* `d3d12-m-packed-formats.zip`: four M frames, forty native network launches, + R11 color/output, RGBA32 motion, R16 exposure, zero shim CPU image copies. +* `d3d12-m-unorm-formats.zip`: four M frames with BGRA8 output, RG16 SNORM + motion and RGBA32 exposure; forty native launches, frame age zero. +* `d3d12-k-packed-formats.zip`: four K frames and forty-four native launches. +* Each of those three complete runs matches its FP16 baseline after the + independent D3D12 resource-format quantisation: RGB max absolute/relative + error 0, PSNR infinite. No tolerance was relaxed for packed output. + +Reproduce a complete packed M run with locally supplied NVIDIA DLLs: + +```powershell +$env:D4R_ZLUDA_NATIVE_DIR="$PWD/build/native-m-gfx1201" +$env:ZLUDA_CACHE_DIR="$PWD/build/zluda-cache-windows" +powershell -NoProfile -File scripts/windows/test-windows-rdna4.ps1 -RuntimeProfile therock -ZludaRoot "$PWD/dist/zluda-windows-f16-reference" -PackageRoot "$PWD/dist/windows-rdna4-command-list" -NgxCore "$PWD/_nvngx.dll" -DlssDll "$PWD/nvngx_dlss.dll" -NgxMode d3d12 -Preset 13 -NgxOnly -RequireNativeNetwork -CommandListBackend -PixelProfile packed -Trace -Iterations 4 -OutputDirectory "$PWD/test-results/d3d12-m-packed-formats" +``` + +One public diagnostic command, without NVIDIA DLLs: + +```powershell +powershell -NoProfile -File scripts/windows/test-windows-rdna4.ps1 -RuntimeProfile therock -ZludaRoot "$PWD/dist/zluda-windows-f16-reference" -PackageRoot "$PWD/dist/windows-rdna4-command-list" +``` + +The exact complete-frame comparison uses `frame_compare.py --exact` with +`--output-storage r11g11b10f` or `bgra8`. Ordinary comparisons still use the +existing FP16 output gates. Game integration next needs command-list initial +PSO handling, enhanced-layout tracking, optional mask import and OptiScaler +API/state-hook compatibility. The game directory is still untouched. + ## M temporal precision regression (2026-09-30) The earlier finite-frame and sampled-layer results did not cover every spatial diff --git a/scripts/windows/test-windows-rdna4.ps1 b/scripts/windows/test-windows-rdna4.ps1 index f54aa96..d462405 100644 --- a/scripts/windows/test-windows-rdna4.ps1 +++ b/scripts/windows/test-windows-rdna4.ps1 @@ -13,6 +13,7 @@ param( [switch]$Trace, [switch]$RequireNativeNetwork, [switch]$CommandListBackend, + [ValidateSet('baseline','packed','unorm')][string]$PixelProfile = 'baseline', [int]$Iterations = 32, [int]$TimeoutSeconds = 180 ) @@ -212,8 +213,11 @@ try { $mapOk = $interopOk = $true if (!$NgxOnly) { $commandOk = Invoke-Probe 'd3d12-command-backend' (Join-Path $bin 'd4r_d3d12_command_probe.exe') @( - '--hip-root', $HipRoot, '--module', (Join-Path $bin 'd4r_nvngx.dll'), '--iterations', "$Iterations") + '--hip-root', $HipRoot, '--module', (Join-Path $bin 'd4r_nvngx.dll'), '--interop-mode', 'indirect', '--iterations', "$Iterations") if (!$commandOk) { throw 'D3D12 command-list order/state/lifetime validation failed.' } + $pixelsOk = Invoke-Probe 'd3d12-gpu-pixel-formats' (Join-Path $bin 'd4r_d3d12_pixel_probe.exe') @( + '--hip-root', $HipRoot, '--module', (Join-Path $bin 'pixel_convert_gfx1201.hsaco'), '--iterations', '2') + if (!$pixelsOk) { throw 'D3D12 typed SRV/UAV format conversion validation failed.' } $mapOk = Invoke-Probe 'interop-map-lifetime' (Join-Path $bin 'd4r_d3d12_hip_interop_probe.exe') @( '--hip-root', $HipRoot, '--interop-mode', 'map', '--iterations', '64') $interopOk = Invoke-Probe 'interop-roundtrip' (Join-Path $bin 'd4r_d3d12_hip_interop_probe.exe') @( @@ -253,7 +257,8 @@ try { # alongside the private user-supplied DLLs for this test. $localShim = Join-Path $ngxRuntimeDirectory 'd4r_nvngx.dll' Copy-Item -LiteralPath (Join-Path $bin 'd4r_nvngx.dll') -Destination $localShim - $ngxArguments += @('--module', $localShim) + Copy-Item -LiteralPath (Join-Path $bin 'pixel_convert_gfx1201.hsaco') -Destination $ngxRuntimeDirectory + $ngxArguments += @('--module', $localShim, '--pixel-profile', $PixelProfile) if ($CommandListBackend) { $ngxArguments += @('--interop-mode', 'command-list') } } $ngxName = if ($NgxMode -eq 'evaluate') { "ngx-evaluate-preset-$Preset" } elseif ($NgxMode -eq 'd3d12') { "d3d12-evaluate-preset-$Preset" } else { 'ngx-init' } diff --git a/tools/windows/d3d12_dlss_probe.cpp b/tools/windows/d3d12_dlss_probe.cpp index 84de1f4..c7c60e0 100644 --- a/tools/windows/d3d12_dlss_probe.cpp +++ b/tools/windows/d3d12_dlss_probe.cpp @@ -3,6 +3,7 @@ #include "ngx_cuda_evaluate.h" #include #include +#include int main(int argc, char** argv) { using namespace d4r::diag; @@ -17,6 +18,8 @@ int main(int argc, char** argv) { const std::filesystem::path directory(diagDir ? wide(diagDir) : std::filesystem::current_path().wstring()); std::filesystem::create_directories(directory); const bool commandBackend = args.interop_mode == "command-list"; + if (args.interop_mode == "images" && args.pixel_profile != "baseline") throw std::runtime_error("Raw CUDA image diagnostic requires baseline formats"); + std::printf("D3D12_PIXEL_PROFILE %s\n", args.pixel_profile.c_str()); if (commandBackend && _putenv_s("D4R_D3D12_COMMAND_BACKEND", "1")) throw std::runtime_error("Cannot enable command backend"); for (const auto& pair : {std::make_pair("D4R_HIP_ROOT", args.hip_root), std::make_pair("D4R_NVCUDA_DLL", args.cuda_dll), std::make_pair("D4R_NVAPI_DLL", args.nvapi_dll), @@ -63,6 +66,10 @@ int main(int argc, char** argv) { desc.Height = i == 4 ? 1 : i == 3 ? outHeight : height; desc.DepthOrArraySize = desc.MipLevels = desc.SampleDesc.Count = 1; desc.Format = i == 0 || i == 3 ? DXGI_FORMAT_R16G16B16A16_FLOAT : i == 2 ? DXGI_FORMAT_R16G16_FLOAT : DXGI_FORMAT_R32_FLOAT; + if (args.pixel_profile == "packed") desc.Format = i == 0 || i == 3 ? DXGI_FORMAT_R11G11B10_FLOAT : + i == 2 ? DXGI_FORMAT_R32G32B32A32_FLOAT : i == 4 ? DXGI_FORMAT_R16_FLOAT : DXGI_FORMAT_R32_FLOAT; + if (args.pixel_profile == "unorm") desc.Format = i == 3 ? DXGI_FORMAT_B8G8R8A8_UNORM : + i == 2 ? DXGI_FORMAT_R16G16_SNORM : i == 4 ? DXGI_FORMAT_R32G32B32A32_FLOAT : desc.Format; if (i == 3) desc.Flags = D3D12_RESOURCE_FLAG_ALLOW_UNORDERED_ACCESS; D3D12_HEAP_PROPERTIES heap{}; heap.Type = D3D12_HEAP_TYPE_DEFAULT; heap.CreationNodeMask = heap.VisibleNodeMask = 1; @@ -79,11 +86,16 @@ int main(int argc, char** argv) { if (i == 0) { uint16_t pixel[] = {uint16_t(((x / 16 + y / 16) & 1) ? 0x3a00 : 0x3400), uint16_t((x & 32) ? 0x3800 : 0x3400), uint16_t((y & 32) ? 0x3a00 : 0x3800), 0x3c00}; - std::memcpy(row + x * 8, pixel, sizeof(pixel)); + if (desc.Format == DXGI_FORMAT_R11G11B10_FLOAT) { + const uint32_t packed = uint32_t(pixel[0] >> 4) | (uint32_t(pixel[1] >> 4) << 11) | (uint32_t(pixel[2] >> 5) << 22); + std::memcpy(row + x * 4, &packed, 4); + } else std::memcpy(row + x * 8, pixel, sizeof(pixel)); } else if (i == 1 || i == 4) { - const float value = i == 1 ? .5f : 1.f; std::memcpy(row + x * 4, &value, 4); + if (desc.Format == DXGI_FORMAT_R16_FLOAT) { const uint16_t one = 0x3c00; std::memcpy(row + x * 2, &one, 2); } + else { const float value = i == 1 ? .5f : 1.f; std::memcpy(row + x * 4, &value, 4); } } else if (i == 3) { - const uint16_t pixel[] = {0x7e00, 0x7e00, 0x7e00, 0x7e00}; std::memcpy(row + x * 8, pixel, 8); + if (desc.Format == DXGI_FORMAT_R16G16B16A16_FLOAT) { const uint16_t pixel[] = {0x7e00, 0x7e00, 0x7e00, 0x7e00}; std::memcpy(row + x * 8, pixel, 8); } + else { const uint32_t sentinel = 0xffffffff; std::memcpy(row + x * 4, &sentinel, 4); } } } texture.upload->Unmap(0, nullptr); @@ -206,8 +218,30 @@ int main(int argc, char** argv) { void* raw = nullptr; D3D12_RANGE range{0, size_t(textures[3].bytes)}; dx(readback->Map(0, &range, &raw), "Map(verification only)"); std::vector pixels(size_t(outWidth) * outHeight * 4); - for (unsigned y = 0; y < outHeight; ++y) std::memcpy(pixels.data() + size_t(y) * outWidth * 4, - static_cast(raw) + y * textures[3].footprint.Footprint.RowPitch, outWidth * 8); + const auto format = resource_desc(textures[3].image.Get()).Format; + if (format == DXGI_FORMAT_R16G16B16A16_FLOAT) { + for (unsigned y = 0; y < outHeight; ++y) std::memcpy(pixels.data() + size_t(y) * outWidth * 4, + static_cast(raw) + y * textures[3].footprint.Footprint.RowPitch, outWidth * 8); + } else { + std::ofstream native(directory / ("output-" + std::to_string(frame) + (args.pixel_profile == "packed" ? ".r11g11b10f" : ".bgra8")), std::ios::binary); + // Independently decode the diagnostic readback for previews and + // frame comparisons. No CPU image copies are in the backend. + static const auto unormHalf = [] { + std::array lut{}; + for (unsigned b = 0; b < 256; ++b) { double best = INFINITY; for (unsigned h = 0; h <= 0x3c00; ++h) { double error = std::abs(double(half_value(uint16_t(h))) - double(b)/255.); if (error < best || (error == best && !(h & 1))) { best = error; lut[b] = uint16_t(h); } } } + return lut; + }(); + for (unsigned y = 0; y < outHeight; ++y) { + const auto* row = static_cast(raw) + y * textures[3].footprint.Footprint.RowPitch; + native.write(reinterpret_cast(row), outWidth * 4); + for (unsigned x = 0; x < outWidth; ++x) { + auto* p = pixels.data() + (size_t(y) * outWidth + x) * 4; + if (format == DXGI_FORMAT_R11G11B10_FLOAT) { uint32_t value = 0; std::memcpy(&value,row + x * 4,4); p[0] = uint16_t((value & 2047) << 4); p[1] = uint16_t(((value >> 11) & 2047) << 4); p[2] = uint16_t((value >> 22) << 5); p[3] = 0x3c00; } + else for (unsigned c = 0; c < 4; ++c) p[c] = unormHalf[row[x * 4 + (c == 0 ? 2 : c == 2 ? 0 : c)]]; + } + } + if (!native) throw std::runtime_error("Saving native output format failed"); + } D3D12_RANGE noWrite{}; readback->Unmap(0, &noWrite); double sum = 0, square = 0; for (size_t i = 0; i < pixels.size(); ++i) { diff --git a/tools/windows/d3d12_pixel_probe.cpp b/tools/windows/d3d12_pixel_probe.cpp new file mode 100644 index 0000000..30f9c26 --- /dev/null +++ b/tools/windows/d3d12_pixel_probe.cpp @@ -0,0 +1,233 @@ +// Independent reference: D3D12 typed SRV loads and typed UAV stores. +#include "pixel_conversion.h" +#include +#include +#include +#include +using namespace d4r::diag; +using namespace d4r::win; +namespace { +constexpr unsigned width = 259, height = 17, count = width * height; +struct Case { DXGI_FORMAT resource, view; unsigned plane; const char* name; }; +const Case cases[] = { + {DXGI_FORMAT_R16G16B16A16_FLOAT, DXGI_FORMAT_R16G16B16A16_FLOAT, 0, "rgba16f"}, + {DXGI_FORMAT_R32G32B32A32_FLOAT, DXGI_FORMAT_R32G32B32A32_FLOAT, 0, "rgba32f"}, + {DXGI_FORMAT_R8G8B8A8_UNORM, DXGI_FORMAT_R8G8B8A8_UNORM, 0, "rgba8"}, + {DXGI_FORMAT_B8G8R8A8_UNORM, DXGI_FORMAT_B8G8R8A8_UNORM, 0, "bgra8"}, + {DXGI_FORMAT_R11G11B10_FLOAT, DXGI_FORMAT_R11G11B10_FLOAT, 0, "r11g11b10f"}, + {DXGI_FORMAT_R10G10B10A2_UNORM, DXGI_FORMAT_R10G10B10A2_UNORM, 0, "rgb10a2"}, + {DXGI_FORMAT_R16G16_FLOAT, DXGI_FORMAT_R16G16_FLOAT, 2, "rg16f"}, + {DXGI_FORMAT_R32G32_FLOAT, DXGI_FORMAT_R32G32_FLOAT, 2, "rg32f"}, + {DXGI_FORMAT_R16G16_SNORM, DXGI_FORMAT_R16G16_SNORM, 2, "rg16snorm"}, + {DXGI_FORMAT_R16G16_UNORM, DXGI_FORMAT_R16G16_UNORM, 2, "rg16unorm"}, + {DXGI_FORMAT_R16G16B16A16_FLOAT, DXGI_FORMAT_R16G16B16A16_FLOAT, 2, "motion-rgba16f"}, + {DXGI_FORMAT_R32G32B32A32_FLOAT, DXGI_FORMAT_R32G32B32A32_FLOAT, 2, "motion-rgba32f"}, + {DXGI_FORMAT_R32_FLOAT, DXGI_FORMAT_R32_FLOAT, 1, "depth-r32f"}, + {DXGI_FORMAT_R16_TYPELESS, DXGI_FORMAT_R16_UNORM, 1, "depth-r16unorm"}, + {DXGI_FORMAT_R24G8_TYPELESS, DXGI_FORMAT_R24_UNORM_X8_TYPELESS, 1, "depth24"}, + {DXGI_FORMAT_R16_TYPELESS, DXGI_FORMAT_R16_FLOAT, 4, "exposure-r16f"}, + {DXGI_FORMAT_R32_FLOAT, DXGI_FORMAT_R32_FLOAT, 4, "exposure-r32f"}, + {DXGI_FORMAT_R16G16B16A16_FLOAT, DXGI_FORMAT_R16G16B16A16_FLOAT, 4, "exposure-rgba16f"}, + {DXGI_FORMAT_R32G32B32A32_FLOAT, DXGI_FORMAT_R32G32B32A32_FLOAT, 4, "exposure-rgba32f"}, +}; +struct List { + ComPtr alloc; + ComPtr list; + explicit List(ID3D12Device* device) { + dx(device->CreateCommandAllocator(D3D12_COMMAND_LIST_TYPE_DIRECT, IID_PPV_ARGS(alloc.GetAddressOf())), "Pixel allocator"); + dx(device->CreateCommandList(0, D3D12_COMMAND_LIST_TYPE_DIRECT, alloc.Get(), nullptr, IID_PPV_ARGS(list.GetAddressOf())), "Pixel command list"); + } + void submit(ID3D12CommandQueue* queue, SharedTimeline& timeline) { + dx(list->Close(), "Close pixel list"); ID3D12CommandList* lists[] = {list.Get()}; + queue->ExecuteCommandLists(1, lists); timeline.drain(queue); + } +}; +ComPtr make_texture(ID3D12Device* device, DXGI_FORMAT fmt, bool uav = false) { + D3D12_RESOURCE_DESC desc{}; desc.Dimension = D3D12_RESOURCE_DIMENSION_TEXTURE2D; + desc.Width = width; desc.Height = height; desc.DepthOrArraySize = desc.MipLevels = desc.SampleDesc.Count = 1; + desc.Format = fmt; desc.Flags = uav ? D3D12_RESOURCE_FLAG_ALLOW_UNORDERED_ACCESS : D3D12_RESOURCE_FLAG_NONE; + D3D12_HEAP_PROPERTIES heap{}; heap.Type = D3D12_HEAP_TYPE_DEFAULT; heap.CreationNodeMask = heap.VisibleNodeMask = 1; + ComPtr result; + dx(device->CreateCommittedResource(&heap, D3D12_HEAP_FLAG_NONE, &desc, D3D12_RESOURCE_STATE_COMMON, + nullptr, IID_PPV_ARGS(result.GetAddressOf())), "Pixel texture"); return result; +} +ComPtr uav_buffer(ID3D12Device* device) { + D3D12_RESOURCE_DESC desc{}; desc.Dimension = D3D12_RESOURCE_DIMENSION_BUFFER; desc.Width = count * 16; + desc.Height = desc.DepthOrArraySize = desc.MipLevels = desc.SampleDesc.Count = 1; + desc.Layout = D3D12_TEXTURE_LAYOUT_ROW_MAJOR; desc.Flags = D3D12_RESOURCE_FLAG_ALLOW_UNORDERED_ACCESS; + D3D12_HEAP_PROPERTIES heap{}; heap.Type = D3D12_HEAP_TYPE_DEFAULT; heap.CreationNodeMask = heap.VisibleNodeMask = 1; + ComPtr result; + dx(device->CreateCommittedResource(&heap, D3D12_HEAP_FLAG_NONE, &desc, D3D12_RESOURCE_STATE_UNORDERED_ACCESS, + nullptr, IID_PPV_ARGS(result.GetAddressOf())), "Pixel reference buffer"); return result; +} +void fixture(uint8_t* row, unsigned y, PixelSpec spec) { + for (unsigned x = 0; x < width; ++x) { + unsigned i = x + y * width; auto* p = row + x * spec.bytes; + const uint16_t half[] = {0, 1, 0x3ff, 0x400, 0x1400, 0x3555, 0x3800, 0x3bff, 0x3c00, 0x3c01, 0x5c07, 0x7bff, 0x8001, 0xbc00}; + switch (spec.storage) { + case d4r::pixel::rgba16f: case d4r::pixel::rg16f: case d4r::pixel::r16f: + for (unsigned c = 0; c < spec.bytes / 2; ++c) reinterpret_cast(p)[c] = half[(i + c * 3) % std::size(half)]; break; + case d4r::pixel::rgba32f: case d4r::pixel::rg32f: case d4r::pixel::r32f: + for (unsigned c = 0; c < spec.bytes / 4; ++c) reinterpret_cast(p)[c] = float(int((i * 23 + c * 13) % 4097) - 2048) / 4096.f; break; + case d4r::pixel::r11g11b10f: + *reinterpret_cast(p) = (i % 1984) | (((i * 13) % 1984) << 11) | (((i * 7) % 992) << 22); break; + case d4r::pixel::rgb10a2: + *reinterpret_cast(p) = (i % 1024) | (((i * 13) % 1024) << 10) | (((i * 7) % 1024) << 20) | ((i % 4) << 30); break; + case d4r::pixel::rgba8: case d4r::pixel::bgra8: + for (unsigned c = 0; c < 4; ++c) p[c] = uint8_t(i * (2 * c + 1) + c * 43); break; + case d4r::pixel::rg16snorm: case d4r::pixel::rg16unorm: case d4r::pixel::r16unorm: + for (unsigned c = 0; c < spec.bytes / 2; ++c) reinterpret_cast(p)[c] = uint16_t(i < 4 ? i * 32767 : i * 193 + c * 37); break; + case d4r::pixel::depth24: *reinterpret_cast(p) = (i * 3791) | ((i % 256) << 24); break; + default: throw std::runtime_error("Missing diagnostic fixture"); + } + } +} +uint16_t reference_half(uint32_t bits) { + // Independent nearest-value search over the binary16 number line. HLSL + // f32tof16 may truncate; it is not the CUDA RNE conversion being tested. + static const std::vector values = [] { + std::vector v(31744); + for (unsigned h = 0; h < v.size(); ++h) v[h] = h < 1024 ? std::ldexp(double(h), -24) : + std::ldexp(1. + double(h & 1023) / 1024., int(h >> 10) - 15); + return v; + }(); + float value = 0; std::memcpy(&value, &bits, 4); + double magnitude = std::abs(double(value)); + if (!std::isfinite(magnitude) || magnitude > values.back()) throw std::runtime_error("Unexpected non-finite decode reference"); + auto upper = std::lower_bound(values.begin(), values.end(), magnitude); + unsigned h = unsigned(upper - values.begin()); + if (h && (magnitude - values[h - 1] < values[h] - magnitude || + (magnitude - values[h - 1] == values[h] - magnitude && (h & 1)))) --h; + return uint16_t(h | ((bits >> 16) & 0x8000)); +} +struct Shaders { + ComPtr root; + ComPtr decode, encode; + explicit Shaders(ID3D12Device* device) { + wchar_t system[32768]{}; if (!GetSystemDirectoryW(system, 32768)) throw std::runtime_error("System path"); + Library compiler(std::filesystem::path(system) / L"d3dcompiler_47.dll"); + auto compile = compiler.symbol("D3DCompile"); + const char* source = R"( +Texture2D Input : register(t0); +StructuredBuffer Canonical : register(t1); +RWStructuredBuffer Output : register(u0); +RWTexture2D Expected : register(u1); +cbuffer Constants : register(b0) { uint Width, Height, Plane; }; +[numthreads(64,1,1)] void read(uint3 t:SV_DispatchThreadID) { + if (t.x >= Width*Height) return; + float4 v = Input.Load(int3(t.x % Width,t.x / Width,0)); + Output[t.x] = asuint(v); +} +[numthreads(64,1,1)] void write(uint3 t:SV_DispatchThreadID) { + if (t.x < Width*Height) Expected[uint2(t.x % Width,t.x / Width)] = f16tof32(Canonical[t.x]); +})"; + D3D12_DESCRIPTOR_RANGE range[3]{}; + for (unsigned i = 0; i < 3; ++i) { range[i].RangeType = i == 2 ? D3D12_DESCRIPTOR_RANGE_TYPE_UAV : D3D12_DESCRIPTOR_RANGE_TYPE_SRV; range[i].NumDescriptors = 1; range[i].BaseShaderRegister = i == 2 ? 1 : i; } + D3D12_ROOT_PARAMETER parameters[5]{}; + for (unsigned i = 0; i < 3; ++i) { parameters[i].ParameterType = D3D12_ROOT_PARAMETER_TYPE_DESCRIPTOR_TABLE; parameters[i].DescriptorTable.NumDescriptorRanges = 1; parameters[i].DescriptorTable.pDescriptorRanges = &range[i]; } + parameters[3].ParameterType = D3D12_ROOT_PARAMETER_TYPE_UAV; + parameters[4].ParameterType = D3D12_ROOT_PARAMETER_TYPE_32BIT_CONSTANTS; parameters[4].Constants.Num32BitValues = 3; + D3D12_ROOT_SIGNATURE_DESC sd{}; sd.NumParameters = 5; sd.pParameters = parameters; + ComPtr blob, errors; dx(D3D12SerializeRootSignature(&sd, D3D_ROOT_SIGNATURE_VERSION_1, blob.GetAddressOf(), errors.GetAddressOf()), "Pixel root serialize"); + dx(device->CreateRootSignature(0, blob->GetBufferPointer(), blob->GetBufferSize(), IID_PPV_ARGS(root.GetAddressOf())), "Pixel root"); + for (unsigned i = 0; i < 2; ++i) { + if (FAILED(compile(source, std::strlen(source), nullptr, nullptr, nullptr, i ? "write" : "read", "cs_5_1", D3DCOMPILE_OPTIMIZATION_LEVEL3, + 0, blob.ReleaseAndGetAddressOf(), errors.ReleaseAndGetAddressOf()))) throw std::runtime_error(errors ? std::string(static_cast(errors->GetBufferPointer())) : "Pixel HLSL compilation"); + D3D12_COMPUTE_PIPELINE_STATE_DESC pd{}; pd.pRootSignature = root.Get(); pd.CS = {blob->GetBufferPointer(), blob->GetBufferSize()}; + dx(device->CreateComputePipelineState(&pd, IID_PPV_ARGS((i ? encode : decode).GetAddressOf())), "Pixel PSO"); + } + } +}; +std::vector read_reference(ID3D12Device* device, ID3D12CommandQueue* queue, SharedTimeline& timeline, + Shaders& shader, ID3D12DescriptorHeap* heap, D3D12_GPU_DESCRIPTOR_HANDLE srv, ID3D12Resource* reference, unsigned plane) { + auto readback = make_buffer(device, count * 16, D3D12_HEAP_TYPE_READBACK, D3D12_HEAP_FLAG_NONE, D3D12_RESOURCE_STATE_COPY_DEST); + List list(device); list.list->SetPipelineState(shader.decode.Get()); list.list->SetComputeRootSignature(shader.root.Get()); + list.list->SetDescriptorHeaps(1, &heap); list.list->SetComputeRootDescriptorTable(0, srv); + list.list->SetComputeRootUnorderedAccessView(3, reference->GetGPUVirtualAddress()); + const unsigned constants[] = {width, height, plane}; list.list->SetComputeRoot32BitConstants(4, 3, constants, 0); + list.list->Dispatch((count + 63) / 64, 1, 1); + transition(list.list.Get(), reference, D3D12_RESOURCE_STATE_UNORDERED_ACCESS, D3D12_RESOURCE_STATE_COPY_SOURCE); + list.list->CopyResource(readback.Get(), reference); + transition(list.list.Get(), reference, D3D12_RESOURCE_STATE_COPY_SOURCE, D3D12_RESOURCE_STATE_UNORDERED_ACCESS); + list.submit(queue, timeline); + void* data = nullptr; D3D12_RANGE range{0, count * 16}; dx(readback->Map(0, &range, &data), "Pixel readback"); + std::vector result(count * 4); std::memcpy(result.data(), data, count * 16); D3D12_RANGE noWrite{}; readback->Unmap(0, &noWrite); return result; +} +} +int main(int argc, char** argv) { + start(); + try { + Args args(argc, argv); HipApi hip(args.hip_root); hipDeviceProp_t props{}; hip.select_gfx1201(args.device, props); + ExternalApi external{hip}; PixelProgram program(hip, wide(args.module)); + ComPtr factory; dx(CreateDXGIFactory1(IID_PPV_ARGS(factory.GetAddressOf())), "Pixel DXGI"); + ComPtr adapter; + for (UINT i = 0; ; ++i) { dx(factory->EnumAdapters1(i, adapter.ReleaseAndGetAddressOf()), "Pixel adapter"); DXGI_ADAPTER_DESC1 d{}; dx(adapter->GetDesc1(&d), "Pixel adapter desc"); if (!std::memcmp(&d.AdapterLuid, props.luid, sizeof(LUID))) break; } + ComPtr device; dx(D3D12CreateDevice(adapter.Get(), D3D_FEATURE_LEVEL_12_0, IID_PPV_ARGS(device.GetAddressOf())), "Pixel D3D12"); + ComPtr queue; D3D12_COMMAND_QUEUE_DESC qd{}; qd.Type = D3D12_COMMAND_LIST_TYPE_DIRECT; dx(device->CreateCommandQueue(&qd, IID_PPV_ARGS(queue.GetAddressOf())), "Pixel queue"); + SharedTimeline timeline(external, device.Get()); Shaders shader(device.Get()); auto reference = uav_buffer(device.Get()); + unsigned verified = 0; + for (unsigned iteration = 0; iteration < args.iterations; ++iteration) for (const auto& item : cases) { + PixelSpec spec = pixel_spec(item.plane, item.resource); if (!pixel_supported(item.plane, spec.storage)) throw std::runtime_error("Unsupported fixture plane"); + auto input = make_texture(device.Get(), item.resource); SharedPlane shared(external, device.Get(), resource_desc(input.Get())); + auto upload = make_buffer(device.Get(), shared.bytes, D3D12_HEAP_TYPE_UPLOAD, D3D12_HEAP_FLAG_NONE, D3D12_RESOURCE_STATE_GENERIC_READ); + void* data = nullptr; D3D12_RANGE noRead{}; dx(upload->Map(0, &noRead, &data), "Pixel fixture upload"); + std::memset(data, 0x5a, size_t(shared.bytes)); for (unsigned y = 0; y < height; ++y) fixture(static_cast(data) + y * shared.footprint.Footprint.RowPitch, y, spec); upload->Unmap(0, nullptr); + { List list(device.Get()); D3D12_TEXTURE_COPY_LOCATION src{}, dst{}; src.pResource = upload.Get(); src.Type = D3D12_TEXTURE_COPY_TYPE_PLACED_FOOTPRINT; src.PlacedFootprint = shared.footprint; dst.pResource = input.Get(); dst.Type = D3D12_TEXTURE_COPY_TYPE_SUBRESOURCE_INDEX; + transition(list.list.Get(), input.Get(), D3D12_RESOURCE_STATE_COMMON, D3D12_RESOURCE_STATE_COPY_DEST); list.list->CopyTextureRegion(&dst, 0, 0, 0, &src, nullptr); transition(list.list.Get(), input.Get(), D3D12_RESOURCE_STATE_COPY_DEST, D3D12_RESOURCE_STATE_NON_PIXEL_SHADER_RESOURCE); shared.copy_input(list.list.Get(), input.Get(), D3D12_RESOURCE_STATE_NON_PIXEL_SHADER_RESOURCE); list.submit(queue.Get(), timeline); } + D3D12_DESCRIPTOR_HEAP_DESC hd{}; hd.Type = D3D12_DESCRIPTOR_HEAP_TYPE_CBV_SRV_UAV; hd.NumDescriptors = 4; hd.Flags = D3D12_DESCRIPTOR_HEAP_FLAG_SHADER_VISIBLE; + ComPtr heap; dx(device->CreateDescriptorHeap(&hd, IID_PPV_ARGS(heap.GetAddressOf())), "Pixel descriptors"); + D3D12_CPU_DESCRIPTOR_HANDLE cpu{}, initial{}; D3D12_GPU_DESCRIPTOR_HANDLE gpu{}; +#if defined(__MINGW32__) + heap->GetCPUDescriptorHandleForHeapStart(&cpu); heap->GetGPUDescriptorHandleForHeapStart(&gpu); +#else + cpu = heap->GetCPUDescriptorHandleForHeapStart(); gpu = heap->GetGPUDescriptorHandleForHeapStart(); +#endif + initial = cpu; unsigned stride = device->GetDescriptorHandleIncrementSize(hd.Type); + D3D12_SHADER_RESOURCE_VIEW_DESC srv{}; srv.Format = item.view; srv.ViewDimension = D3D12_SRV_DIMENSION_TEXTURE2D; srv.Shader4ComponentMapping = D3D12_DEFAULT_SHADER_4_COMPONENT_MAPPING; srv.Texture2D.MipLevels = 1; device->CreateShaderResourceView(input.Get(), &srv, cpu); + auto expected = read_reference(device.Get(), queue.Get(), timeline, shader, heap.Get(), gpu, reference.Get(), item.plane); + unsigned channels = item.plane == 2 ? 2 : item.plane == 0 ? 4 : 1; unsigned element = item.plane == 0 || item.plane == 2 ? 2 : 4; + PixelAllocation canonical(hip, width, height, channels * element); + std::vector guard(size_t(canonical.pitch * height), 0xcd); + hip.check(hip.hipMemcpy(canonical.data, guard.data(), guard.size(), hipMemcpyHostToDevice), "Diagnostic pitch guard initialization"); + program.convert(false, shared.mapped, shared.footprint.Footprint.RowPitch, canonical.data, canonical.pitch, width, height, spec.storage, item.plane); + std::vector actual(size_t(canonical.pitch * height)); hip.check(hip.hipMemcpy(actual.data(), canonical.data, actual.size(), hipMemcpyDeviceToHost), "Diagnostic canonical readback"); + unsigned differs = 0; double maxAbs = 0; + for (unsigned y = 0; y < height; ++y) for (uint64_t offset = width * channels * element; offset < canonical.pitch; ++offset) + if (actual[size_t(y * canonical.pitch + offset)] != 0xcd) throw std::runtime_error("GPU decode wrote past the row/pixel range"); + for (unsigned y = 0; y < height; ++y) for (unsigned x = 0; x < width; ++x) for (unsigned c = 0; c < channels; ++c) { + unsigned i = (y * width + x) * 4 + c; const auto* p = actual.data() + y * canonical.pitch + (x * channels + c) * element; uint32_t got = 0; std::memcpy(&got, p, element); + if (element == 2) expected[i] = reference_half(expected[i]); + bool mismatch = got != expected[i]; + if (element == 4) { float a = 0, b = 0; std::memcpy(&a, &got, 4); std::memcpy(&b, &expected[i], 4); double error = std::abs(double(a) - b); maxAbs = std::max(maxAbs, error); mismatch = !std::isfinite(a) || !std::isfinite(b) || error > 1.2e-7 * std::max(1., std::abs(double(b))); } + if (mismatch && differs++ < 5) std::fprintf(stderr, "PIXEL_DIFFERENT format=%s pixel=%u channel=%u expected=%08x actual=%08x\n", item.name, y * width + x, c, expected[i], got); + } + if (differs) throw std::runtime_error(std::string("D3D12 typed decode mismatch: ") + item.name + " count=" + std::to_string(differs)); + if (item.plane == 0) { + // BGRA uses the identical UNORM converter through an RGBA UAV; + // the consumer SRV independently tests the channel permutation. + DXGI_FORMAT encodedFormat = spec.storage == d4r::pixel::bgra8 ? DXGI_FORMAT_R8G8B8A8_UNORM : item.view; + auto encoded = make_texture(device.Get(), encodedFormat, true); auto actualTexture = make_texture(device.Get(), item.resource); + auto canonicalUpload = make_buffer(device.Get(), count * 16, D3D12_HEAP_TYPE_UPLOAD, D3D12_HEAP_FLAG_NONE, D3D12_RESOURCE_STATE_GENERIC_READ); + dx(canonicalUpload->Map(0, &noRead, &data), "Encode fixture upload"); + auto* values = static_cast(data); std::vector halves(count * 4); + // Cover finite half values, rounding ties, negatives and overflow. + for (unsigned i = 0; i < count; ++i) for (unsigned c = 0; c < 4; ++c) { uint16_t v = uint16_t(((i * 37 + c * 7919) % 31744) | (i % 11 == 0 ? 0x8000 : 0)); values[i * 4 + c] = v; halves[i * 4 + c] = v; } + canonicalUpload->Unmap(0, nullptr); + std::vector padded(size_t(canonical.pitch * height), 0x5a); for (unsigned y = 0; y < height; ++y) std::memcpy(padded.data() + y * canonical.pitch, halves.data() + y * width * 4, width * 8); + hip.check(hip.hipMemcpy(canonical.data, padded.data(), padded.size(), hipMemcpyHostToDevice), "Diagnostic encode input"); + program.convert(true, canonical.data, canonical.pitch, shared.mapped, shared.footprint.Footprint.RowPitch, width, height, spec.storage, 3); + cpu.ptr = initial.ptr + stride; srv = {}; srv.Shader4ComponentMapping = D3D12_DEFAULT_SHADER_4_COMPONENT_MAPPING; srv.ViewDimension = D3D12_SRV_DIMENSION_BUFFER; srv.Buffer.NumElements = count; srv.Buffer.StructureByteStride = 16; device->CreateShaderResourceView(canonicalUpload.Get(), &srv, cpu); + cpu.ptr += stride; D3D12_UNORDERED_ACCESS_VIEW_DESC uv{}; uv.Format = encodedFormat; uv.ViewDimension = D3D12_UAV_DIMENSION_TEXTURE2D; device->CreateUnorderedAccessView(encoded.Get(), nullptr, &uv, cpu); + { List list(device.Get()); transition(list.list.Get(), encoded.Get(), D3D12_RESOURCE_STATE_COMMON, D3D12_RESOURCE_STATE_UNORDERED_ACCESS); list.list->SetPipelineState(shader.encode.Get()); list.list->SetComputeRootSignature(shader.root.Get()); ID3D12DescriptorHeap* heaps[] = {heap.Get()}; list.list->SetDescriptorHeaps(1, heaps); auto handle = gpu; handle.ptr += stride; list.list->SetComputeRootDescriptorTable(1, handle); handle.ptr += stride; list.list->SetComputeRootDescriptorTable(2, handle); unsigned constants[] = {width,height,0}; list.list->SetComputeRoot32BitConstants(4,3,constants,0); list.list->Dispatch((count+63)/64,1,1); transition(list.list.Get(), encoded.Get(), D3D12_RESOURCE_STATE_UNORDERED_ACCESS, D3D12_RESOURCE_STATE_NON_PIXEL_SHADER_RESOURCE); shared.copy_output(list.list.Get(), actualTexture.Get(), D3D12_RESOURCE_STATE_COMMON); transition(list.list.Get(), actualTexture.Get(), D3D12_RESOURCE_STATE_COMMON, D3D12_RESOURCE_STATE_NON_PIXEL_SHADER_RESOURCE); list.submit(queue.Get(),timeline); } + srv = {}; srv.Format = item.view; srv.ViewDimension = D3D12_SRV_DIMENSION_TEXTURE2D; srv.Shader4ComponentMapping = D3D12_DEFAULT_SHADER_4_COMPONENT_MAPPING; srv.Texture2D.MipLevels = 1; device->CreateShaderResourceView(actualTexture.Get(), &srv, initial); + auto encodedActual = read_reference(device.Get(), queue.Get(), timeline, shader, heap.Get(), gpu, reference.Get(), 4); + srv.Format = encodedFormat; device->CreateShaderResourceView(encoded.Get(), &srv, initial); + auto encodedExpected = read_reference(device.Get(), queue.Get(), timeline, shader, heap.Get(), gpu, reference.Get(), 4); + for (unsigned i = 0; i < count * 4; ++i) if (encodedActual[i] != encodedExpected[i]) { if (differs++ < 5) std::fprintf(stderr, "ENCODE_DIFFERENT format=%s pixel=%u channel=%u expected=%08x actual=%08x\n", item.name, i/4,i%4,encodedExpected[i],encodedActual[i]); } + if (differs) throw std::runtime_error(std::string("D3D12 typed encode mismatch: ") + item.name + " count=" + std::to_string(differs)); + } + ++verified; std::printf("PASS PIXEL_FORMAT name=%s plane=%u pixels=%u decode_max_abs=%.9g encode=%u\n", item.name,item.plane,count,maxAbs,item.plane==0); + } + std::printf("PASS D3D12_GPU_FORMATS architecture=gfx1201 cases=%u iterations=%u width=%u height=%u typed_srv_reference=1 typed_uav_reference=1\n", verified,args.iterations,width,height); return 0; + } catch (const std::exception& error) { std::fprintf(stderr,"FAIL D3D12_GPU_FORMATS %s\n",error.what()); loaded_modules(); return 1; } +} diff --git a/tools/windows/diagnostic.h b/tools/windows/diagnostic.h index 312bfb3..ce92398 100644 --- a/tools/windows/diagnostic.h +++ b/tools/windows/diagnostic.h @@ -110,6 +110,7 @@ struct Args { std::string hip_root, cuda_dll, module, context = "primary"; std::string stream_lifetime = "persistent"; std::string interop_mode = "roundtrip"; + std::string pixel_profile = "baseline"; std::string ngx_core, dlss_dll; std::string nvapi_dll; std::string ngx_abi = "driver"; @@ -131,6 +132,10 @@ struct Args { else if (key == "--context") context = value; else if (key == "--stream-lifetime") stream_lifetime = value; else if (key == "--interop-mode") interop_mode = value; + else if (key == "--pixel-profile") { + if (value != "baseline" && value != "packed" && value != "unorm") throw std::runtime_error("pixel-profile must be baseline, packed or unorm"); + pixel_profile = value; + } else if (key == "--ngx-core") ngx_core = value; else if (key == "--dlss-dll") dlss_dll = value; else if (key == "--nvapi-dll") nvapi_dll = value; diff --git a/tools/windows/frame_compare.py b/tools/windows/frame_compare.py index 01f8bbe..95408f4 100644 --- a/tools/windows/frame_compare.py +++ b/tools/windows/frame_compare.py @@ -20,6 +20,16 @@ def main(args): raise RuntimeError(f'Output size mismatch for {path.name}') if not np.isfinite(reference).all() or not np.isfinite(output).all(): raise RuntimeError(f'NaN/Inf in {path.name}') + if args.output_storage == 'r11g11b10f': + # D3D float narrowing truncates. Compare after the independent + # resource-format quantisation, rather than relaxing the PSNR gate. + bits = reference.reshape(-1, 4).view(np.uint16).copy() + bits[bits & 0x8000 != 0] = 0 + bits[:, :2] &= 0x7ff0 + bits[:, 2] &= 0x7fe0 + reference = bits.view(np.float16).reshape(-1) + elif args.output_storage == 'bgra8': + reference = (np.rint(np.clip(reference.astype(np.float32), 0, 1) * 255) / 255).astype(np.float16) x = reference.reshape(-1, 4)[:, :3].astype(np.float64) y = output.reshape(-1, 4)[:, :3].astype(np.float64) error = np.abs(x - y) @@ -30,7 +40,7 @@ def main(args): print(f'FRAME_REFERENCE name={path.name} max_abs={error.max():.9g} ' f'max_rel={(error / np.maximum(np.abs(x), .001)).max():.9g} rms={rms:.9g} psnr_db={psnr:.6g} ' f'fraction_over_1pct={outliers:.9g}') - if psnr < args.min_psnr or error.max() > .025 * peak or outliers > .0001: + if (args.exact and np.any(error != 0)) or psnr < args.min_psnr or error.max() > .025 * peak or outliers > .0001: raise RuntimeError(f'Complete pipeline mismatch in {path.name}') print(f'PASS FRAME_REFERENCE frames={len(frames)} finite=1 min_psnr_db={args.min_psnr}') @@ -42,6 +52,8 @@ if __name__ == '__main__': parser.add_argument('--width', type=int, default=512) parser.add_argument('--height', type=int, default=288) parser.add_argument('--min-psnr', type=float, default=60.) + parser.add_argument('--output-storage', choices=('rgba16f', 'r11g11b10f', 'bgra8'), default='rgba16f') + parser.add_argument('--exact', action='store_true') args = parser.parse_args() try: main(args) diff --git a/tools/windows/ngx_windows_runtime.h b/tools/windows/ngx_windows_runtime.h index bdf5a29..55a4a31 100644 --- a/tools/windows/ngx_windows_runtime.h +++ b/tools/windows/ngx_windows_runtime.h @@ -1,5 +1,6 @@ #pragma once #include "d3d12_external.h" +#include "pixel_conversion.h" #include "cuda_image_api.h" #include "ngx_parameters.h" #include "d3d12_command_hooks.h" @@ -18,6 +19,15 @@ inline void ngx_check(unsigned result, const char* operation) { std::printf("D4R_NGX %s result=0x%08x\n", operation, result); if (result != 1) throw std::runtime_error(std::string(operation) + " NGX=" + std::to_string(result)); } +inline std::filesystem::path pixel_module_path() { + if (const char* path = std::getenv("D4R_FORMAT_MODULE")) return diag::wide(env_path("D4R_FORMAT_MODULE")); + HMODULE module = nullptr; + if (!GetModuleHandleExW(GET_MODULE_HANDLE_EX_FLAG_FROM_ADDRESS | GET_MODULE_HANDLE_EX_FLAG_UNCHANGED_REFCOUNT, + reinterpret_cast(&pixel_module_path), &module)) throw std::runtime_error("Locate GPU conversion module directory"); + wchar_t path[32768]{}; const DWORD length = GetModuleFileNameW(module, path, 32768); + if (!length || length == 32768) throw std::runtime_error("GPU conversion module path is invalid"); + return std::filesystem::path(path).parent_path() / L"pixel_convert_gfx1201.hsaco"; +} struct Runtime { diag::HipApi hip; @@ -32,6 +42,7 @@ struct Runtime { CUcontext context = nullptr; bool initialized = false; std::unique_ptr timeline; + std::unique_ptr pixels; std::mutex mutex; using Init = unsigned(*)(unsigned long long, const wchar_t*, unsigned); using Shutdown = unsigned(*)(); @@ -105,6 +116,7 @@ struct Runtime { if (context) (void)cuda.cuCtxSetCurrent(context); if (initialized) { (void)shutdown(); initialized = false; } timeline.reset(); + pixels.reset(); if (context) { (void)cuda.cuCtxSynchronize(); (void)cuda.cuCtxSetCurrent(nullptr); (void)cuda.cuDevicePrimaryCtxRelease_v2(ordinal); context = nullptr; @@ -123,13 +135,15 @@ class Feature { D3D12_RESOURCE_DESC desc{}; std::unique_ptr shared; std::unique_ptr image; + std::unique_ptr canonical; + PixelSpec spec{}; } planes_[5]; void cleanup() noexcept { (void)rt_->cuda.cuCtxSetCurrent(rt_->context); (void)rt_->cuda.cuCtxSynchronize(); if (handle_) { (void)rt_->release(handle_); handle_ = nullptr; } for (auto& plane : planes_) { - plane.image.reset(); plane.shared.reset(); plane.texture.Reset(); + plane.image.reset(); plane.canonical.reset(); plane.shared.reset(); plane.texture.Reset(); } if (scratch_) { (void)rt_->cuda.cuMemFree_v2(scratch_); scratch_ = 0; } if (parameters_) { (void)rt_->destroy(parameters_); parameters_ = nullptr; } @@ -209,17 +223,22 @@ private: throw std::runtime_error("Interop baseline requires single-sample 2D textures"); const bool half = i == 0 || i == 2 || i == 3; const unsigned channels = i == 0 || i == 3 ? 4 : i == 2 ? 2 : 1; - const DXGI_FORMAT format = channels == 4 ? DXGI_FORMAT_R16G16B16A16_FLOAT : - channels == 2 ? DXGI_FORMAT_R16G16_FLOAT : DXGI_FORMAT_R32_FLOAT; - if (desc.Format != format && !(i == 1 && (desc.Format == DXGI_FORMAT_R32_TYPELESS || desc.Format == DXGI_FORMAT_D32_FLOAT))) - throw std::runtime_error(std::string("GPU conversion not yet implemented for ") + names[i] + " format=" + std::to_string(desc.Format)); + if (desc.Width > UINT32_MAX) throw std::runtime_error("Texture width exceeds GPU conversion ABI"); + const PixelSpec spec = pixel_spec(i, desc.Format); + if (!pixel_supported(i, spec.storage)) throw std::runtime_error(std::string("Unsupported ") + names[i] + " DXGI format=" + std::to_string(desc.Format)); auto& plane = planes_[i]; if (!plane.shared || plane.desc.Width != desc.Width || plane.desc.Height != desc.Height || plane.desc.Format != desc.Format) { rt_->cuda.check(rt_->cuda.cuCtxSynchronize(), "Synchronize(resize interop)"); - plane.image.reset(); plane.shared.reset(); + plane.image.reset(); plane.canonical.reset(); plane.shared.reset(); plane.shared = std::make_unique(rt_->external, rt_->device.Get(), desc); + if (!spec.direct) { + if (!rt_->pixels) rt_->pixels = std::make_unique(rt_->hip, pixel_module_path()); + plane.canonical = std::make_unique(rt_->hip, unsigned(desc.Width), desc.Height, channels * (half ? 2 : 4)); + } plane.image = std::make_unique(rt_->images, unsigned(desc.Width), desc.Height, half ? 16 : 32, channels, i == 3, i == 0 ? 1 : 0); - plane.desc = desc; + plane.desc = desc; plane.spec = spec; + std::printf("D4R_FORMAT plane=%s dxgi=%u canonical_channels=%u canonical_bits=%u gpu_conversion=%u cpu_copy=0\n", + names[i], unsigned(desc.Format), channels, half ? 16u : 32u, !spec.direct); } plane.texture = texture; d4r_ngx_set_void(parameters_, names[i], &plane.image->object); @@ -240,7 +259,14 @@ private: rt_->timeline->wait_input(queue); rt_->cuda.check(rt_->cuda.cuCtxSynchronize(), "Interop input producer completion"); for (unsigned i : {0u, 1u, 2u, 4u}) if (planes_[i].texture) { - planes_[i].image->upload_device(reinterpret_cast(planes_[i].shared->mapped), planes_[i].shared->footprint.Footprint.RowPitch); + auto& plane = planes_[i]; + void* data = plane.shared->mapped; uint64_t pitch = plane.shared->footprint.Footprint.RowPitch; + if (plane.canonical) { + rt_->pixels->convert(false, data, pitch, plane.canonical->data, plane.canonical->pitch, + unsigned(plane.desc.Width), plane.desc.Height, plane.spec.storage, i); + data = plane.canonical->data; pitch = plane.canonical->pitch; + } + plane.image->upload_device(reinterpret_cast(data), pitch); } if (std::getenv("D4R_INTEROP_VERIFY")) { for (unsigned i : {0u, 1u, 2u, 4u}) if (planes_[i].texture) { @@ -259,8 +285,12 @@ private: // NGX can use internal nonblocking streams. The default stream alone // does not establish completion of those streams. rt_->cuda.check(rt_->cuda.cuCtxSynchronize(), "NGX all-stream completion"); - planes_[3].image->download_device(reinterpret_cast(planes_[3].shared->mapped), planes_[3].shared->footprint.Footprint.RowPitch); + auto& output = planes_[3]; + output.image->download_device(reinterpret_cast(output.canonical ? output.canonical->data : output.shared->mapped), + output.canonical ? output.canonical->pitch : output.shared->footprint.Footprint.RowPitch); rt_->cuda.check(rt_->cuda.cuCtxSynchronize(), "VRAM output completion"); + if (output.canonical) rt_->pixels->convert(true, output.canonical->data, output.canonical->pitch, output.shared->mapped, + output.shared->footprint.Footprint.RowPitch, unsigned(output.desc.Width), output.desc.Height, output.spec.storage, 3); rt_->timeline->signal_output(queue); std::printf("D4R_FRAME cpu_frame_copies=%u frame_age=0 interop_ngx_ms=%.3f\n", std::getenv("D4R_INTEROP_VERIFY") ? 1u : 0u, std::chrono::duration(std::chrono::steady_clock::now() - begin).count()); diff --git a/tools/windows/pixel_conversion.h b/tools/windows/pixel_conversion.h new file mode 100644 index 0000000..f18b75b --- /dev/null +++ b/tools/windows/pixel_conversion.h @@ -0,0 +1,86 @@ +#pragma once +#include "d3d12_external.h" +#include "pixel_format.h" + +namespace d4r::win { +struct PixelSpec { + pixel::Storage storage; + unsigned bytes; + bool direct; +}; +inline PixelSpec pixel_spec(unsigned plane, DXGI_FORMAT format) { + using namespace pixel; + // Typeless resources use the corresponding floating-point/UNORM view, + // as NGX's resource-only ABI provides no SRV descriptor to disambiguate. + switch (format) { + case DXGI_FORMAT_R16G16B16A16_FLOAT: case DXGI_FORMAT_R16G16B16A16_TYPELESS: + return {rgba16f, 8, plane == color || plane == output}; + case DXGI_FORMAT_R32G32B32A32_FLOAT: return {rgba32f, 16, false}; + case DXGI_FORMAT_R11G11B10_FLOAT: return {r11g11b10f, 4, false}; + case DXGI_FORMAT_R10G10B10A2_UNORM: case DXGI_FORMAT_R10G10B10A2_TYPELESS: return {rgb10a2, 4, false}; + case DXGI_FORMAT_R8G8B8A8_UNORM: case DXGI_FORMAT_R8G8B8A8_TYPELESS: return {rgba8, 4, false}; + case DXGI_FORMAT_B8G8R8A8_UNORM: case DXGI_FORMAT_B8G8R8A8_TYPELESS: return {bgra8, 4, false}; + case DXGI_FORMAT_R16G16_FLOAT: case DXGI_FORMAT_R16G16_TYPELESS: return {rg16f, 4, plane == motion}; + case DXGI_FORMAT_R32G32_FLOAT: case DXGI_FORMAT_R32G32_TYPELESS: return {rg32f, 8, false}; + case DXGI_FORMAT_R16G16_SNORM: return {rg16snorm, 4, false}; + case DXGI_FORMAT_R16G16_UNORM: return {rg16unorm, 4, false}; + case DXGI_FORMAT_R32_FLOAT: case DXGI_FORMAT_R32_TYPELESS: case DXGI_FORMAT_D32_FLOAT: + return {r32f, 4, plane == depth || plane == exposure}; + case DXGI_FORMAT_R16_FLOAT: return {r16f, 2, false}; + case DXGI_FORMAT_D24_UNORM_S8_UINT: case DXGI_FORMAT_R24G8_TYPELESS: case DXGI_FORMAT_R24_UNORM_X8_TYPELESS: + return {depth24, 4, false}; + case DXGI_FORMAT_R16_TYPELESS: return {plane == exposure ? r16f : r16unorm, 2, false}; + case DXGI_FORMAT_D16_UNORM: case DXGI_FORMAT_R16_UNORM: + return {r16unorm, 2, false}; + default: throw std::runtime_error("No validated GPU conversion for DXGI format=" + std::to_string(format)); + } +} +inline bool pixel_supported(unsigned plane, pixel::Storage storage) { + using namespace pixel; + switch (plane) { + case color: case output: return storage <= rgb10a2; + case motion: return storage == rgba16f || storage == rgba32f || (storage >= rg16f && storage <= rg16unorm); + case depth: return storage == r32f || storage == depth24 || storage == r16unorm; + case exposure: return storage == r32f || storage == r16f || storage == rgba16f || storage == rgba32f; + default: return false; + } +} +class PixelProgram { + diag::HipApi& hip_; + hipModule_t module_ = nullptr; + hipFunction_t decode_ = nullptr, encode_ = nullptr; +public: + PixelProgram(diag::HipApi& hip, const std::filesystem::path& module) : hip_(hip) { + try { + hip.check(hip.hipModuleLoad(&module_, diag::utf8(module.c_str()).c_str()), "Load native GPU format conversion"); + hip.check(hip.hipModuleGetFunction(&decode_, module_, "d4r_pixel_decode"), "Resolve GPU pixel decode"); + hip.check(hip.hipModuleGetFunction(&encode_, module_, "d4r_pixel_encode"), "Resolve GPU pixel encode"); + } catch (...) { if (module_) (void)hip.hipModuleUnload(module_); throw; } + } + ~PixelProgram() { if (module_) (void)hip_.hipModuleUnload(module_); } + PixelProgram(const PixelProgram&) = delete; + void convert(bool encode, const void* source, uint64_t sourcePitch, void* dest, uint64_t destPitch, + unsigned width, unsigned height, pixel::Storage format, unsigned plane) { + if (!width || !height || uint64_t(width) * height > UINT32_MAX) throw std::runtime_error("Invalid GPU conversion dimensions"); + unsigned storage = unsigned(format); + void* args[] = {&source, &sourcePitch, &dest, &destPitch, &width, &height, &storage, &plane}; + hip_.check(hip_.hipModuleLaunchKernel(encode ? encode_ : decode_, unsigned((uint64_t(width) * height + 255) / 256), 1, 1, + 256, 1, 1, 0, nullptr, args, nullptr), encode ? "GPU pixel encode" : "GPU pixel decode"); + // CUDA may use other nonblocking streams. Establish completion before + // giving the converted allocation to that API; images remain in VRAM. + hip_.check(hip_.hipDeviceSynchronize(), "GPU conversion completion"); + } +}; +class PixelAllocation { + diag::HipApi& hip_; +public: + void* data = nullptr; + uint64_t pitch; + PixelAllocation(diag::HipApi& hip, unsigned width, unsigned height, unsigned canonicalBytes) : hip_(hip) { + pitch = (uint64_t(width) * canonicalBytes + 255) & ~uint64_t(255); + hip.check(hip.hipMalloc(&data, size_t(pitch * height)), "Allocate canonical GPU pixels"); + } + ~PixelAllocation() { if (data) (void)hip_.hipFree(data); } + PixelAllocation(const PixelAllocation&) = delete; +}; +} diff --git a/tools/windows/pixel_convert_gfx1201.hip b/tools/windows/pixel_convert_gfx1201.hip new file mode 100644 index 0000000..b4d8907 --- /dev/null +++ b/tools/windows/pixel_convert_gfx1201.hip @@ -0,0 +1,133 @@ +#include "../../kernels/common/hip_device_minimal.h" +#include "pixel_format.h" + +namespace { +using namespace d4r::pixel; +using u8 = unsigned char; +using u16 = unsigned short; +using u32 = unsigned; +using u64 = unsigned long long; +__device__ inline float as_float(u32 bits) { return __builtin_bit_cast(float, bits); } +__device__ inline u32 as_uint(float value) { return __builtin_bit_cast(u32, value); } +__device__ inline float from_half(u16 bits) { + u32 sign = u32(bits & 0x8000u) << 16, exp = (bits >> 10) & 31, m = bits & 1023; + if (!exp) return as_float(as_uint(float(m) * 0x1p-24f) | sign); + return as_float(sign | (exp == 31 ? 0x7f800000u : (exp + 112u) << 23) | (m << 13)); +} +__device__ inline u16 to_half(float value) { + u32 bits = as_uint(value), sign = (bits >> 16) & 0x8000u, exponent = (bits >> 23) & 255u; + u32 mantissa = bits & 0x7fffffu; + if (exponent == 255) return u16(sign | 0x7c00u | (mantissa ? 0x200u : 0)); + int e = int(exponent) - 112; + if (e >= 31) return u16(sign | 0x7c00u); + unsigned shift = 13; + if (e <= 0) { + if (e < -10) return u16(sign); + mantissa |= 0x800000u; shift = unsigned(14 - e); e = 0; + } + u32 kept = mantissa >> shift, remainder = mantissa & ((1u << shift) - 1), midpoint = 1u << (shift - 1); + u32 result = sign | (u32(e) << 10); + result += kept & (e == 0 ? 0x7ffu : 0x3ffu); + return u16(result + (remainder > midpoint || (remainder == midpoint && (kept & 1)))); +} +__device__ inline u32 quantize(float value, unsigned maximum) { + if (!(value > 0)) return 0; + if (value >= 1) return maximum; + float scaled = value * float(maximum); + u32 integer = u32(scaled); float remainder = scaled - float(integer); + return integer + (remainder > .5f || (remainder == .5f && (integer & 1))); +} +__device__ inline u32 half_to_unsigned(u16 h, unsigned mantissa_bits) { + if ((h & 0x7c00u) == 0x7c00u && (h & 1023u)) return (31u << mantissa_bits) | (1u << (mantissa_bits - 1)); + if (h & 0x8000u) return 0; + // D3D floating-point format narrowing truncates. RNE here differs from + // the driver's typed UAV store and can turn finite HDR values into Inf. + return h >> (10 - mantissa_bits); +} +__device__ inline unsigned pixel_bytes(Storage fmt) { + switch (fmt) { + case rgba32f: return 16; + case rgba16f: case rg32f: return 8; + case r16f: case r16unorm: return 2; + default: return 4; + } +} +__device__ inline void decode(const u8* p, Storage fmt, float* value) { + const auto* halves = reinterpret_cast(p); + const auto* singles = reinterpret_cast(p); + value[0] = value[1] = value[2] = 0; value[3] = 1; + switch (fmt) { + case rgba16f: for (unsigned i = 0; i < 4; ++i) value[i] = from_half(halves[i]); break; + case rgba32f: for (unsigned i = 0; i < 4; ++i) value[i] = singles[i]; break; + case rgba8: case bgra8: + for (unsigned i = 0; i < 4; ++i) value[i] = float(p[i]) / 255.f; + if (fmt == bgra8) { float swap = value[0]; value[0] = value[2]; value[2] = swap; } break; + case r11g11b10f: { + u32 packed = *reinterpret_cast(p); + value[0] = from_half(u16((packed & 2047) << 4)); + value[1] = from_half(u16(((packed >> 11) & 2047) << 4)); + value[2] = from_half(u16((packed >> 22) << 5)); break; + } + case rgb10a2: { + u32 packed = *reinterpret_cast(p); + for (unsigned i = 0; i < 3; ++i) value[i] = float((packed >> (10 * i)) & 1023) / 1023.f; + value[3] = float(packed >> 30) / 3.f; break; + } + case rg16f: value[0] = from_half(halves[0]); value[1] = from_half(halves[1]); break; + case rg32f: value[0] = singles[0]; value[1] = singles[1]; break; + case rg16snorm: + for (unsigned i = 0; i < 2; ++i) { + int v = reinterpret_cast(p)[i]; + value[i] = v == -32768 ? -1.f : float(v) / 32767.f; + } break; + case rg16unorm: for (unsigned i = 0; i < 2; ++i) value[i] = float(halves[i]) / 65535.f; break; + case r16f: value[0] = from_half(halves[0]); break; + case r32f: value[0] = singles[0]; break; + case depth24: value[0] = float(*reinterpret_cast(p) & 0xffffffu) / 16777215.f; break; + case r16unorm: value[0] = float(halves[0]) / 65535.f; break; + default: break; // Host rejects formats which have no tested decoder. + } +} +__device__ inline void encode(u8* p, Storage fmt, const float* value) { + auto* halves = reinterpret_cast(p); + auto* singles = reinterpret_cast(p); + switch (fmt) { + case rgba16f: for (unsigned i = 0; i < 4; ++i) halves[i] = to_half(value[i]); break; + case rgba32f: for (unsigned i = 0; i < 4; ++i) singles[i] = value[i]; break; + case rgba8: case bgra8: + for (unsigned i = 0; i < 4; ++i) p[i] = u8(quantize(value[i], 255)); + if (fmt == bgra8) { u8 swap = p[0]; p[0] = p[2]; p[2] = swap; } break; + case r11g11b10f: + *reinterpret_cast(p) = half_to_unsigned(to_half(value[0]), 6) | + (half_to_unsigned(to_half(value[1]), 6) << 11) | (half_to_unsigned(to_half(value[2]), 5) << 22); break; + case rgb10a2: + *reinterpret_cast(p) = quantize(value[0], 1023) | (quantize(value[1], 1023) << 10) | + (quantize(value[2], 1023) << 20) | (quantize(value[3], 3) << 30); break; + default: break; // Only color/output formats are encoded in the runtime. + } +} +} + +extern "C" __global__ void d4r_pixel_decode(const u8* source, u64 source_pitch, u8* dest, u64 dest_pitch, + unsigned width, unsigned height, unsigned format, unsigned plane) { + unsigned i = blockIdx.x * 256u + threadIdx.x; + if (u64(i) >= u64(width) * height) return; + unsigned x = i % width, y = i / width; + float value[4]; decode(source + u64(y) * source_pitch + u64(x) * pixel_bytes(Storage(format)), Storage(format), value); + if (plane == depth || plane == exposure) reinterpret_cast(dest + u64(y) * dest_pitch)[x] = value[0]; + else { + unsigned channels = plane == motion ? 2 : 4; + auto* halves = reinterpret_cast(dest + u64(y) * dest_pitch) + x * channels; + for (unsigned c = 0; c < channels; ++c) halves[c] = to_half(value[c]); + } +} + +extern "C" __global__ void d4r_pixel_encode(const u8* source, u64 source_pitch, u8* dest, u64 dest_pitch, + unsigned width, unsigned height, unsigned format, unsigned plane) { + unsigned i = blockIdx.x * 256u + threadIdx.x; + if (u64(i) >= u64(width) * height) return; + unsigned x = i % width, y = i / width; + const auto* halves = reinterpret_cast(source + u64(y) * source_pitch) + x * 4; + float value[4]; for (unsigned c = 0; c < 4; ++c) value[c] = from_half(halves[c]); + encode(dest + u64(y) * dest_pitch + u64(x) * pixel_bytes(Storage(format)), Storage(format), value); +} diff --git a/tools/windows/pixel_format.h b/tools/windows/pixel_format.h new file mode 100644 index 0000000..3acc354 --- /dev/null +++ b/tools/windows/pixel_format.h @@ -0,0 +1,12 @@ +#pragma once + +// Numeric storage descriptions shared by the native HIP kernel and Windows +// host. These are independent of Win32, Wine, KFD and NVIDIA SDK headers. +namespace d4r::pixel { +enum Storage : unsigned { + rgba16f, rgba32f, rgba8, bgra8, r11g11b10f, rgb10a2, + rg16f, rg32f, rg16snorm, rg16unorm, r16f, r32f, + depth24, r16unorm, rgba8srgb, bgra8srgb +}; +enum Plane : unsigned { color, depth, motion, output, exposure }; +}