Convert native Windows D3D12 formats on gfx1201 with GPU validation

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 "$<TARGET_FILE:d4r_nvngx_windows>" --interop-mode indirect --iterations 32)
 set_tests_properties(d3d12_command_backend PROPERTIES TIMEOUT 120 LABELS "hardware;d3d12;command_lists;lifetime")

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
+**Windows / RDNA4 development:** `windows-rdna4` runs native K and M networks
-gfx1201 HIP/ZLUDA diagnostics, validated D3D12/HIP interop, an NGX init probe,
+in a standalone Windows D3D12 harness on an RX 9070 XT (gfx1201). All K layers
-and a CPU-checked gfx12 WMMA backend. K `enc1` and `enc2` execute on
+and all spatial blocks in forty temporal M launches pass NumPy/PTX validation.
-gfx1201 and match the existing numpy model on synthetic nonzero fixtures.
+The Windows shim imports D3D12 VRAM through HIP, executes DLSS at command-list
-This branch incorporates upstream RDNA4 commit `dbef4b2`; native Windows CMake
+submission, and returns output in the same frame without CPU image copies.
-builds all 11 K and 5 M modules with the upstream gfx12 layout. Other layers,
+RGBA16/32F, R11G11B10, RGBA/BGRA8, depth, exposure and motion formats have
-full K/M image validation and the Windows game runtime are still in progress.
+independent D3D12 typed SRV/UAV conversion tests. Four-frame K/M runs with packed
-Build commands, hardware results and remaining
+output match the FP16 baseline exactly after resource-format quantisation.
-porting gates are in [docs/windows-rdna4-port.md](docs/windows-rdna4-port.md).
+The build incorporates upstream RDNA4 commit `dbef4b2`. OptiScaler/game
-This branch does not yet provide a Windows DLSS runtime.
+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.
 

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

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

tools/windows/d3d12_dlss_probe.cpp

@@ -3,6 +3,7 @@
 #include "ngx_cuda_evaluate.h"
 #include <dxgi1_6.h>
 #include <memory>
+#include <array>
 
 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<uint16_t> pixels(size_t(outWidth) * outHeight * 4);
+            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<uint8_t*>(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<uint16_t, 256> 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<uint8_t*>(raw) + y * textures[3].footprint.Footprint.RowPitch;
+                    native.write(reinterpret_cast<const char*>(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) {

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 <dxgi1_6.h>
+#include <d3dcompiler.h>
+#include <array>
+#include <cmath>
+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<ID3D12CommandAllocator> alloc;
+    ComPtr<ID3D12GraphicsCommandList> 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<ID3D12Resource> 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<ID3D12Resource> 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<ID3D12Resource> 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<ID3D12Resource> 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<uint16_t*>(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<float*>(p)[c] = float(int((i * 23 + c * 13) % 4097) - 2048) / 4096.f; break;
+            case d4r::pixel::r11g11b10f:
+                *reinterpret_cast<uint32_t*>(p) = (i % 1984) | (((i * 13) % 1984) << 11) | (((i * 7) % 992) << 22); break;
+            case d4r::pixel::rgb10a2:
+                *reinterpret_cast<uint32_t*>(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<uint16_t*>(p)[c] = uint16_t(i < 4 ? i * 32767 : i * 193 + c * 37); break;
+            case d4r::pixel::depth24: *reinterpret_cast<uint32_t*>(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<double> values = [] {
+        std::vector<double> 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<ID3D12RootSignature> root;
+    ComPtr<ID3D12PipelineState> 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<decltype(&D3DCompile)>("D3DCompile");
+        const char* source = R"(
+Texture2D<float4> Input : register(t0);
+StructuredBuffer<uint4> Canonical : register(t1);
+RWStructuredBuffer<uint4> Output : register(u0);
+RWTexture2D<float4> 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<ID3DBlob> 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<char*>(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<uint32_t> 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<uint32_t> 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<IDXGIFactory4> factory; dx(CreateDXGIFactory1(IID_PPV_ARGS(factory.GetAddressOf())), "Pixel DXGI");
+        ComPtr<IDXGIAdapter1> 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<ID3D12Device> device; dx(D3D12CreateDevice(adapter.Get(), D3D_FEATURE_LEVEL_12_0, IID_PPV_ARGS(device.GetAddressOf())), "Pixel D3D12");
+        ComPtr<ID3D12CommandQueue> 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<uint8_t*>(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<ID3D12DescriptorHeap> 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<uint8_t> 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<uint8_t> 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<uint32_t*>(data); std::vector<uint16_t> 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<uint8_t> 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; }
+}

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;

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)

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<LPCWSTR>(&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<SharedTimeline> timeline;
+    std::unique_ptr<PixelProgram> 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<SharedPlane> shared;
         std::unique_ptr<cuda::Image> image;
+        std::unique_ptr<PixelAllocation> 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 :
+            if (desc.Width > UINT32_MAX) throw std::runtime_error("Texture width exceeds GPU conversion ABI");
-                channels == 2 ? DXGI_FORMAT_R16G16_FLOAT : DXGI_FORMAT_R32_FLOAT;
+            const PixelSpec spec = pixel_spec(i, desc.Format);
-            if (desc.Format != format && !(i == 1 && (desc.Format == DXGI_FORMAT_R32_TYPELESS || desc.Format == DXGI_FORMAT_D32_FLOAT)))
+            if (!pixel_supported(i, spec.storage)) throw std::runtime_error(std::string("Unsupported ") + names[i] + " DXGI format=" + std::to_string(desc.Format));
-                throw std::runtime_error(std::string("GPU conversion not yet implemented for ") + names[i] + " 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<SharedPlane>(rt_->external, rt_->device.Get(), desc);
+                if (!spec.direct) {
+                    if (!rt_->pixels) rt_->pixels = std::make_unique<PixelProgram>(rt_->hip, pixel_module_path());
+                    plane.canonical = std::make_unique<PixelAllocation>(rt_->hip, unsigned(desc.Width), desc.Height, channels * (half ? 2 : 4));
+                }
                 plane.image = std::make_unique<cuda::Image>(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<uintptr_t>(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<uintptr_t>(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<uintptr_t>(planes_[3].shared->mapped), planes_[3].shared->footprint.Footprint.RowPitch);
+        auto& output = planes_[3];
+        output.image->download_device(reinterpret_cast<uintptr_t>(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<double, std::milli>(std::chrono::steady_clock::now() - begin).count());

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

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<const u16*>(p);
+    const auto* singles = reinterpret_cast<const float*>(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<const u32*>(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<const u32*>(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<const short*>(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<const u32*>(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<u16*>(p);
+    auto* singles = reinterpret_cast<float*>(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<u32*>(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<u32*>(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<float*>(dest + u64(y) * dest_pitch)[x] = value[0];
+    else {
+        unsigned channels = plane == motion ? 2 : 4;
+        auto* halves = reinterpret_cast<u16*>(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<const u16*>(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);
+}

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