ayo
ac4ee6c225
Add a native D3D12 FSR Ray Regeneration route for gfx1200/gfx1201, using the game's device and queue with embedded HLSL guide preparation and output composition. Preserve the vkd3d compatibility route for RDNA3 and expose explicit route overrides and diagnostics. Improve Windows runtime compatibility: - Retry rejected application-ID initialization with a project identity. - Defer expensive RR backend creation until feature creation. - Recover from transient submission failures without disabling later frames. - Add shared-buffer export fallbacks through explicit heaps and a private device on the same adapter. - Safely evict old RR descriptor sets when input views change. - Treat requested RR presets as hints when a configured model is available. - Expand SR/RR input formats and handle diffuse/specular hit distances, packed normals, depth-stencil resources, and unused FSR guides. Reduce SR overhead by batching output downloads and format conversion under one completion wait. Optimize strict M enc2/tube attention with bounded FP32 accumulation and the original FP64 fallback for ambiguous rounding cases. Extend RDNA4 tooling with gfx1200/gfx1201 native FP8 RR model builds, capability checks, architecture-aware deployment, and experimental K/M kernel paths. Require matching model and runtime arithmetic families. Add optional embedded FSR shader optimizations keyed by shader hash and exact original SPIR-V contents. Preserve reflected metadata and record optimizer provenance, artifact hashes, and cache identity. Simplify Windows installation with automatic backend dependency downloads, download progress and ETA, backed-up absolute OptiScaler library paths, and recovery from interrupted RR model builds. Add build-kernels.bat and a kernels-only build matrix for all eleven registered RDNA targets, with code-object target validation and SHA256 manifests. Expand conversion probes, numerical stress tests, temporal benchmark tooling, installer/deployment tests, and setup documentation. Document synthetic regression and timing results, experimental build options, and remaining RDNA4 hardware validation requirements.
272 lines
25 KiB
C++
272 lines
25 KiB
C++
// 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_R32G8X24_TYPELESS, DXGI_FORMAT_R32_FLOAT_X8X24_TYPELESS, 1, "depth32-stencil8-plane0"},
|
|
{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"},
|
|
{DXGI_FORMAT_R8G8B8A8_UNORM, DXGI_FORMAT_R8G8B8A8_UNORM, 4, "exposure-rgba8"},
|
|
{DXGI_FORMAT_B8G8R8A8_UNORM, DXGI_FORMAT_B8G8R8A8_UNORM, 4, "exposure-bgra8"},
|
|
{DXGI_FORMAT_R9G9B9E5_SHAREDEXP, DXGI_FORMAT_R9G9B9E5_SHAREDEXP, 0, "rgb9e5"},
|
|
{DXGI_FORMAT_R16_FLOAT, DXGI_FORMAT_R16_FLOAT, 1, "depth-r16f"},
|
|
{DXGI_FORMAT_R8G8B8A8_UNORM, DXGI_FORMAT_R8G8B8A8_UNORM, 2, "motion-rgba8"},
|
|
};
|
|
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;
|
|
if (fmt == DXGI_FORMAT_R32G8X24_TYPELESS) desc.Flags = D3D12_RESOURCE_FLAG_ALLOW_DEPTH_STENCIL;
|
|
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;
|
|
}
|
|
ResourceAccess shader_access() {
|
|
ResourceAccess result; result.enhanced = true; result.layout = D3D12_BARRIER_LAYOUT_DIRECT_QUEUE_SHADER_RESOURCE;
|
|
result.access = D3D12_BARRIER_ACCESS_SHADER_RESOURCE; return result;
|
|
}
|
|
void native_layout(ID3D12GraphicsCommandList* list, ID3D12Resource* texture, D3D12_BARRIER_LAYOUT before, D3D12_BARRIER_LAYOUT after,
|
|
D3D12_BARRIER_ACCESS accessBefore, D3D12_BARRIER_ACCESS accessAfter) {
|
|
ComPtr<ID3D12GraphicsCommandList7> newer; dx(list->QueryInterface(IID_PPV_ARGS(newer.GetAddressOf())), "Probe enhanced list");
|
|
D3D12_TEXTURE_BARRIER barrier{}; barrier.pResource = texture; barrier.LayoutBefore = before; barrier.LayoutAfter = after;
|
|
barrier.AccessBefore = accessBefore; barrier.AccessAfter = accessAfter; barrier.SyncBefore = barrier.SyncAfter = D3D12_BARRIER_SYNC_ALL;
|
|
D3D12_BARRIER_GROUP group{}; group.Type = D3D12_BARRIER_TYPE_TEXTURE; group.NumBarriers = 1; group.pTextureBarriers = &barrier; newer->Barrier(1,&group);
|
|
}
|
|
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;
|
|
case d4r::pixel::rgb9e5: *reinterpret_cast<uint32_t*>(p) = i * 2654435761u; break; // every exponent and mantissa
|
|
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_architecture(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; }
|
|
ComPtr<ID3D12Device> device; dx(D3D12CreateDevice(adapter.Get(), D3D_FEATURE_LEVEL_12_0, IID_PPV_ARGS(device.GetAddressOf())), "Pixel D3D12");
|
|
const bool enhanced = args.interop_mode == "enhanced";
|
|
if (enhanced) { D3D12_FEATURE_DATA_D3D12_OPTIONS12 support{}; dx(device->CheckFeatureSupport(D3D12_FEATURE_D3D12_OPTIONS12, &support, sizeof(support)), "Enhanced barrier support"); if (!support.EnhancedBarriersSupported) throw std::runtime_error("Enhanced barriers unsupported"); }
|
|
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()));
|
|
std::printf("PIXEL_FOOTPRINT format=%s plane0_format=%u pitch=%u storage_bytes=%u\n", item.name, shared.footprint.Footprint.Format, shared.footprint.Footprint.RowPitch, spec.bytes);
|
|
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, enhanced ? D3D12_RESOURCE_STATE_COMMON : D3D12_RESOURCE_STATE_NON_PIXEL_SHADER_RESOURCE);
|
|
if (enhanced) native_layout(list.list.Get(),input.Get(),D3D12_BARRIER_LAYOUT_COMMON,D3D12_BARRIER_LAYOUT_DIRECT_QUEUE_SHADER_RESOURCE,D3D12_BARRIER_ACCESS_COMMON,D3D12_BARRIER_ACCESS_SHADER_RESOURCE);
|
|
shared.copy_input(list.list.Get(), input.Get(), enhanced ? shader_access() : ResourceAccess(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));
|
|
// D3D12 has no typed UAV store for R9G9B9E5, so its encoder has no hardware reference here.
|
|
if (item.plane == 0 && spec.storage != d4r::pixel::rgb9e5) {
|
|
// 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);
|
|
ResourceAccess state; state.enhanced = enhanced; shared.copy_output(list.list.Get(), actualTexture.Get(), state);
|
|
if (enhanced) native_layout(list.list.Get(),actualTexture.Get(),D3D12_BARRIER_LAYOUT_COMMON,D3D12_BARRIER_LAYOUT_DIRECT_QUEUE_SHADER_RESOURCE,D3D12_BARRIER_ACCESS_COMMON,D3D12_BARRIER_ACCESS_SHADER_RESOURCE);
|
|
else 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=%s cases=%u iterations=%u width=%u height=%u enhanced=%u typed_srv_reference=1 typed_uav_reference=1\n", d4r::diag::HipApi::target_arch(), verified,args.iterations,width,height,enhanced); return 0;
|
|
} catch (const std::exception& error) { std::fprintf(stderr,"FAIL D3D12_GPU_FORMATS %s\n",error.what()); loaded_modules(); return 1; }
|
|
}
|