Validate native Windows DLSS output on GPU and record actual M gameplay

tools/windows/d3d12_pixel_probe.cpp

@@ -172,6 +172,14 @@ int main(int argc, char** argv) {
     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));
+        {
+            PixelAllocation diagnostic(hip,3,1,8);
+            const uint16_t bad[] = {0x7e00,0x7c00,0xfc00,0x3c00,0,1,0x3ff,0x7bff,0x8000,0x8001,0xbc00,0x400};
+            hip.check(hip.hipMemcpy(diagnostic.data,bad,sizeof(bad),hipMemcpyHostToDevice), "Upload nonfinite validation fixture");
+            const auto counts = program.validate(diagnostic.data,diagnostic.pitch,3,1);
+            if (counts[0] != 1 || counts[1] != 2) throw std::runtime_error("GPU NaN/Inf validation counters mismatch");
+            std::printf("PASS OUTPUT_VALIDATION_FIXTURE nan=1 inf=2 finite=9\n");
+        }
         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; }

tools/windows/ngx_windows_runtime.h

@@ -322,6 +322,13 @@ private:
         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 (std::getenv("D4R_VALIDATE_OUTPUT")) {
+            if (!rt_->pixels) rt_->pixels = std::make_unique<PixelProgram>(rt_->hip, pixel_module_path());
+            const auto counts = rt_->pixels->validate(output.canonical ? output.canonical->data : output.shared->mapped,
+                output.canonical ? output.canonical->pitch : output.shared->footprint.Footprint.RowPitch, unsigned(output.desc.Width), output.desc.Height);
+            std::printf("D4R_OUTPUT_VALIDATION elements=%llu nan=%u inf=%u diagnostics_cpu_bytes=8\n", output.desc.Width * output.desc.Height * 4, counts[0], counts[1]);
+            if (counts[0] || counts[1]) throw std::runtime_error("DLSS output contains NaN/Inf");
+        }
         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);

tools/windows/pixel_conversion.h

@@ -1,6 +1,7 @@
 #pragma once
 #include "d3d12_external.h"
 #include "pixel_format.h"
+#include <array>
 
 namespace d4r::win {
 struct PixelSpec {
@@ -52,6 +53,8 @@ class PixelProgram {
     diag::HipApi& hip_;
     hipModule_t module_ = nullptr;
     hipFunction_t decode_ = nullptr, encode_ = nullptr;
+    hipFunction_t validate_ = nullptr;
+    void* counters_ = nullptr;
 public:
     PixelProgram(diag::HipApi& hip, const std::filesystem::path& module) : hip_(hip) {
         try {
@@ -60,8 +63,20 @@ public:
             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() { if (counters_) (void)hip_.hipFree(counters_); if (module_) (void)hip_.hipModuleUnload(module_); }
     PixelProgram(const PixelProgram&) = delete;
+    std::array<unsigned,2> validate(const void* source, uint64_t pitch, unsigned width, unsigned height) {
+        if (!width || !height || uint64_t(width) * height > UINT32_MAX / 4) throw std::runtime_error("Invalid output validation dimensions");
+        if (!validate_) hip_.check(hip_.hipModuleGetFunction(&validate_, module_, "d4r_pixel_validate"), "Resolve output validation kernel");
+        if (!counters_) hip_.check(hip_.hipMalloc(&counters_, 2 * sizeof(unsigned)), "Allocate validation counters");
+        std::array<unsigned,2> result{};
+        hip_.check(hip_.hipMemcpy(counters_, result.data(), sizeof(result), hipMemcpyHostToDevice), "Clear validation counters");
+        void* args[] = {&source, &pitch, &width, &height, &counters_};
+        hip_.check(hip_.hipModuleLaunchKernel(validate_, unsigned((uint64_t(width) * height + 255) / 256),1,1,256,1,1,0,nullptr,args,nullptr), "Validate output on GPU");
+        hip_.check(hip_.hipDeviceSynchronize(), "Output validation completion");
+        hip_.check(hip_.hipMemcpy(result.data(), counters_, sizeof(result), hipMemcpyDeviceToHost), "Read validation counters");
+        return result;
+    }
     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");

tools/windows/pixel_convert_gfx1201.hip

@@ -131,3 +131,17 @@ extern "C" __global__ void d4r_pixel_encode(const u8* source, u64 source_pitch,
     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);
 }
+
+// Diagnostic scan of the canonical RGBA16F output; no image readback.
+extern "C" __global__ void d4r_pixel_validate(const u8* source, u64 pitch,
+    unsigned width, unsigned height, unsigned* counters) {
+    unsigned i = blockIdx.x * 256u + threadIdx.x;
+    if (u64(i) >= u64(width) * height) return;
+    const auto* values = reinterpret_cast<const u16*>(source + u64(i / width) * pitch) + (i % width) * 4;
+    unsigned nans = 0, infinities = 0;
+    for (unsigned c = 0; c < 4; ++c) if ((values[c] & 0x7c00u) == 0x7c00u) {
+        if (values[c] & 0x3ffu) ++nans; else ++infinities;
+    }
+    if (nans) __atomic_fetch_add(counters, nans, __ATOMIC_RELAXED);
+    if (infinities) __atomic_fetch_add(counters + 1, infinities, __ATOMIC_RELAXED);
+}