From 87365077e084dc5e42abaca7b1c7d6739e38b7ff Mon Sep 17 00:00:00 2001 From: woutermans Date: Wed, 7 Oct 2026 10:59:53 +0200 Subject: [PATCH] Add experimental packed and in-place DLSS output - Tag redirected surfaces as R10G10B10A2 ('R3') as well as RGBA16F ('R2'), and let the native output kernel store the game's R10G10B10A2 texels itself. - Import linear, exportable game textures (vkd3d-proton patch 0003) into HIP so the output kernel writes the game's output texture with no game-side copy. - Add ID3D12DXVKInteropDeviceD4R3 with GetVulkanLinearImageInfo. - Gate both paths behind D4R_SHIM_OUTPUT_PACKED and D4R_SHIM_INPLACE, and use them only for split frames of preset M whose output kernel honours the redirect. - Convert on the host when a recorded packed frame finds the wrong kernel. - Give the harness an R10G10B10A2 output, and add D4R_SHIM_SPIN_POLL so frame timing is not quantised to the poll sleep. - Document the measured removal of the output copies in performance.md. --- docs/performance.md | 23 ++ kernels/tex/tex_common.h | 30 ++- .../0003-linear-exportable-textures.patch | 204 ++++++++++++++++++ tools/d3d12_dlss_harness.cpp | 23 +- tools/d4r_nvngx_shim.cpp | 191 +++++++++++++++- tools/d4r_vkd3d_interop.h | 10 + tools/d4r_win32_wait.h | 5 +- tools/wine_nvcuda_bridge.c | 5 +- 8 files changed, 469 insertions(+), 22 deletions(-) create mode 100644 patches/vkd3d-proton/0003-linear-exportable-textures.patch diff --git a/docs/performance.md b/docs/performance.md index 8d45382..2a43e21 100644 --- a/docs/performance.md +++ b/docs/performance.md @@ -58,6 +58,29 @@ In SILENT HILL Townfall (static scene, K Quality 1706×960 → 2560×1440, Mango These were measured in the work copy. They have not been re-measured from this repository since the port. +## Experimental: removing the output copies + +After DLSS, the result normally goes from the shared buffer into the game's output texture with one copy, or with a copy and a blit when the output is not RGBA16F. Two opt-in settings remove those passes: + +- `D4R_SHIM_OUTPUT_PACKED=1` (preset M): the native output kernel stores R10G10B10A2 texels itself, rounding as the RGBA16F store followed by the blit does, so the conversion image and blit are not needed. +- `D4R_SHIM_INPLACE=1` with `VKD3D_D4R_LINEAR_TEXTURES=1` (vkd3d-proton patch 0003): committed single-level RGBA16F and R10G10B10A2 textures are created linear on exportable memory, the shim imports the output texture into HIP, and the output kernel writes it directly. + +Measured on an RX 9070 XT (gfx1201) in the D3D12 harness, preset M at 1280×720 → 3840×2160, median frame interval of 4 alternating runs of 400 frames. The harness runs DLSS back to back with no game rendering. + +| Output texture | Passes after DLSS | Frame interval | +|---|---|---| +| RGBA16F, copy | 1 | 5.76 ms | +| RGBA16F, in place | 0 | 5.52 ms | +| R10G10B10A2, copy and blit | 2 | 5.97 ms | +| R10G10B10A2, packed store | 1 | 5.76 ms | +| R10G10B10A2, packed store in place | 0 | 5.65 ms | +| R10G10B10A2, accuracy kernels, copy and blit | 2 | 6.35 ms | +| R10G10B10A2, accuracy kernels, packed store in place | 0 | 5.91 ms | + +Every variant was byte-identical to the path it replaces on moving scenes (`D4R_HARNESS_OUTPUT_RGB10A2=1` gives the harness an R10G10B10A2 output). In Ghost of Tsushima, which passes an R10G10B10A2 output, both settings engaged from the fourth frame. + +Limits: the harness does not render into its textures, so the cost of linear render targets to a game is not measured, and every committed single-level texture of those two formats becomes linear, not only the DLSS output. The packed store is only built into M's downsample kernel. A frame recorded for packed output whose output kernel turns out not to be the native one is converted on the CPU; that path has not been exercised. The input copies remain; see [in-place-inputs.md](in-place-inputs.md). + ## What did not help - **Accumulating in f16 on the WMMA units.** The error grows about 20× and the image degrades. diff --git a/kernels/tex/tex_common.h b/kernels/tex/tex_common.h index 3910696..b680cd7 100644 --- a/kernels/tex/tex_common.h +++ b/kernels/tex/tex_common.h @@ -299,9 +299,23 @@ extern "C" __attribute__((device)) int __ockl_image_width_2D(tsharp_t* image); extern "C" __attribute__((device)) int __ockl_image_height_2D(tsharp_t* image); __attribute__((device)) static inline uint32_t redirect_pitch(uint64_t surface) { - // 'R2' tag in the high half (other handles may have data there), row pitch / 8 in the low half + // 'R2' (RGBA16F texels) or 'R3' (R10G10B10A2_UNORM texels) tag in the high half (other handles may have + // data there), row pitch / 8 in the low half const uint32_t word = ((const __attribute__((address_space(4))) uint32_t*)surface)[21]; - return (word >> 16) == 0x5232u ? (word & 0xffffu) << 3 : 0u; + return (word >> 16) == 0x5232u || (word >> 16) == 0x5233u ? (word & 0xffffu) << 3 : 0u; +} +__attribute__((device)) static inline bool redirect_rgb10a2(uint64_t surface) +{ + return (((const __attribute__((address_space(4))) uint32_t*)surface)[21] >> 16) == 0x5233u; +} +// One channel as the game-side path produces it: the f32 value rounded toward zero to a half (the RGBA16F +// store), then converted to UNORM by a blit (NaN and negatives 0, clamped at 1, rounded to nearest). +__attribute__((device)) static inline uint32_t redirect_unorm(uint16_t half, float scale) +{ + float value = (float)__builtin_bit_cast(_Float16, half); + value = value > 0.0f ? value : 0.0f; + value = value < 1.0f ? value : 1.0f; + return (uint32_t)(value * scale + 0.5f); } // D4R_AS1_STORES: redirect stores through address-space-1 pointers (global_store instead of flat_store; the // redirect target is linear device memory) @@ -332,7 +346,12 @@ DEV __attribute__((always_inline)) void d4r_sust_p_v4b32(uint64_t surface, int32 { const uint32_t lo = __builtin_bit_cast(uint32_t, __builtin_amdgcn_cvt_pkrtz(__builtin_bit_cast(float, a), __builtin_bit_cast(float, b))); const uint32_t hi = __builtin_bit_cast(uint32_t, __builtin_amdgcn_cvt_pkrtz(__builtin_bit_cast(float, c), __builtin_bit_cast(float, d))); - *D4R_REDIRECT_PTR(uint2_t, redirect_row(surface, y) + 8 * x) = (uint2_t){lo, hi}; + if (redirect_rgb10a2(surface)) + *D4R_REDIRECT_PTR(uint32_t, redirect_row(surface, y) + 4 * x) = + redirect_unorm((uint16_t)lo, 1023.0f) | redirect_unorm((uint16_t)(lo >> 16), 1023.0f) << 10 | + redirect_unorm((uint16_t)hi, 1023.0f) << 20 | redirect_unorm((uint16_t)(hi >> 16), 3.0f) << 30; + else + *D4R_REDIRECT_PTR(uint2_t, redirect_row(surface, y) + 8 * x) = (uint2_t){lo, hi}; } return; } @@ -435,8 +454,9 @@ DEV __attribute__((always_inline)) void d4r_sust_b32(uint64_t surface, int32_t x tsharp_t* image = (tsharp_t*)surface; if (redirect_pitch(surface) != 0) { - // raw bits at byte offset x (RGBA16F rows are 8 bytes per pixel) - if (x >= 0 && (x & 3) == 0 && redirect_inside(surface, x >> 3, y)) + // raw bits at byte offset x (RGBA16F rows are 8 bytes per pixel); the shim only asks for packed + // RGB10A2 texels from kernels without raw stores + if (!redirect_rgb10a2(surface) && x >= 0 && (x & 3) == 0 && redirect_inside(surface, x >> 3, y)) *D4R_REDIRECT_PTR(uint32_t, redirect_row(surface, y) + x) = data; return; } diff --git a/patches/vkd3d-proton/0003-linear-exportable-textures.patch b/patches/vkd3d-proton/0003-linear-exportable-textures.patch new file mode 100644 index 0000000..074232a --- /dev/null +++ b/patches/vkd3d-proton/0003-linear-exportable-textures.patch @@ -0,0 +1,204 @@ +diff --git a/include/vkd3d_device_vkd3d_ext.idl b/include/vkd3d_device_vkd3d_ext.idl +index 860640d5..fc91e283 100644 +--- a/include/vkd3d_device_vkd3d_ext.idl ++++ b/include/vkd3d_device_vkd3d_ext.idl +@@ -132,6 +132,21 @@ interface ID3D12DXVKInteropDeviceD4R2 : ID3D12DXVKInteropDeviceD4R + HRESULT RetainExternalResources(ID3D12CommandList *list, IUnknown *resources); + } + ++[ ++ uuid(7b1e5d0a-93c4-4e8f-b6a2-51c08f2e4d73), ++ object, ++ local, ++ pointer_default(unique) ++] ++interface ID3D12DXVKInteropDeviceD4R3 : ID3D12DXVKInteropDeviceD4R2 ++{ ++ /* Textures created linear and exportable (VKD3D_D4R_LINEAR_TEXTURES=1): the ++ * dedicated VkDeviceMemory behind the image and the layout of its only ++ * subresource, so an external API can read or write the texels in place. */ ++ HRESULT GetVulkanLinearImageInfo(ID3D12Resource *resource, UINT64 *vk_memory, UINT64 *memory_size, ++ UINT64 *offset, UINT64 *row_pitch); ++} ++ + [ + uuid(f3112584-41f9-348d-a59b-00b7e1d285d6), + object, +diff --git a/libs/vkd3d/device.c b/libs/vkd3d/device.c +index fd753830..3ba6b1b0 100644 +--- a/libs/vkd3d/device.c ++++ b/libs/vkd3d/device.c +@@ -4659,7 +4659,8 @@ HRESULT STDMETHODCALLTYPE d3d12_device_QueryInterface(d3d12_device_iface *iface, + || IsEqualGUID(riid, &IID_ID3D12DXVKInteropDevice2) + || IsEqualGUID(riid, &IID_ID3D12DXVKInteropDevice3) + || IsEqualGUID(riid, &IID_ID3D12DXVKInteropDeviceD4R) +- || IsEqualGUID(riid, &IID_ID3D12DXVKInteropDeviceD4R2)) ++ || IsEqualGUID(riid, &IID_ID3D12DXVKInteropDeviceD4R2) ++ || IsEqualGUID(riid, &IID_ID3D12DXVKInteropDeviceD4R3)) + { + d3d12_dxvk_interop_device_AddRef(&device->ID3D12DXVKInteropDevice_iface); + *object = &device->ID3D12DXVKInteropDevice_iface; +@@ -11381,7 +11382,7 @@ static void d3d12_device_replace_vtable(struct d3d12_device *device) + } + + extern CONST_VTBL struct ID3D12DeviceExt5Vtbl d3d12_device_vkd3d_ext_vtbl; +-extern CONST_VTBL struct ID3D12DXVKInteropDeviceD4R2Vtbl d3d12_dxvk_interop_device_vtbl; ++extern CONST_VTBL struct ID3D12DXVKInteropDeviceD4R3Vtbl d3d12_dxvk_interop_device_vtbl; + extern CONST_VTBL struct ID3DLowLatencyDeviceVtbl d3d_low_latency_device_vtbl; + extern CONST_VTBL struct IAmdExtAntiLagApiVtbl d3d_amd_ext_anti_lag_vtbl; + +diff --git a/libs/vkd3d/device_vkd3d_ext.c b/libs/vkd3d/device_vkd3d_ext.c +index fc5290da..f18e3c1a 100644 +--- a/libs/vkd3d/device_vkd3d_ext.c ++++ b/libs/vkd3d/device_vkd3d_ext.c +@@ -1145,6 +1145,35 @@ static HRESULT STDMETHODCALLTYPE d3d12_dxvk_interop_device_GetVulkanHeapInfo(d3d + return S_OK; + } + ++static HRESULT STDMETHODCALLTYPE d3d12_dxvk_interop_device_GetVulkanLinearImageInfo(d3d12_dxvk_interop_device_iface *iface, ++ ID3D12Resource *resource, UINT64 *vk_memory, UINT64 *memory_size, UINT64 *offset, UINT64 *row_pitch) ++{ ++ struct d3d12_device *device = d3d12_device_from_ID3D12DXVKInteropDevice(iface); ++ const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs; ++ struct d3d12_resource *resource_impl; ++ VkImageSubresource subresource; ++ VkSubresourceLayout layout; ++ ++ if (!resource || !vk_memory || !memory_size || !offset || !row_pitch) ++ return E_INVALIDARG; ++ ++ resource_impl = impl_from_ID3D12Resource(resource); ++ if (!(resource_impl->flags & VKD3D_RESOURCE_D4R_LINEAR) || ++ resource_impl->mem.device_allocation.vk_memory == VK_NULL_HANDLE) ++ return E_FAIL; ++ ++ subresource.aspectMask = VK_IMAGE_ASPECT_COLOR_BIT; ++ subresource.mipLevel = 0; ++ subresource.arrayLayer = 0; ++ VK_CALL(vkGetImageSubresourceLayout(device->vk_device, resource_impl->res.vk_image, &subresource, &layout)); ++ ++ *vk_memory = (UINT64)resource_impl->mem.device_allocation.vk_memory; ++ *memory_size = (UINT64)resource_impl->mem.device_allocation.size; ++ *offset = (UINT64)resource_impl->mem.offset + layout.offset; ++ *row_pitch = (UINT64)layout.rowPitch; ++ return S_OK; ++} ++ + static HRESULT STDMETHODCALLTYPE d3d12_dxvk_interop_device_SplitCommandListForExternalWait( + d3d12_dxvk_interop_device_iface *iface, ID3D12CommandList *list, UINT64 vk_semaphore, UINT64 value) + { +@@ -1166,7 +1195,7 @@ static HRESULT STDMETHODCALLTYPE d3d12_dxvk_interop_device_RetainExternalResourc + return d3d12_command_list_retain_external_resources(cmd_list, resources); + } + +-CONST_VTBL struct ID3D12DXVKInteropDeviceD4R2Vtbl d3d12_dxvk_interop_device_vtbl = ++CONST_VTBL struct ID3D12DXVKInteropDeviceD4R3Vtbl d3d12_dxvk_interop_device_vtbl = + { + /* IUnknown methods */ + d3d12_dxvk_interop_device_QueryInterface, +@@ -1204,6 +1233,9 @@ CONST_VTBL struct ID3D12DXVKInteropDeviceD4R2Vtbl d3d12_dxvk_interop_device_vtbl + + /* ID3D12DXVKInteropDeviceD4R2 methods */ + d3d12_dxvk_interop_device_RetainExternalResources, ++ ++ /* ID3D12DXVKInteropDeviceD4R3 methods */ ++ d3d12_dxvk_interop_device_GetVulkanLinearImageInfo, + }; + + static inline struct d3d12_device *d3d12_device_from_ID3DLowLatencyDevice(d3d_low_latency_device_iface *iface) +diff --git a/libs/vkd3d/resource.c b/libs/vkd3d/resource.c +index e521ec69..0683c75e 100644 +--- a/libs/vkd3d/resource.c ++++ b/libs/vkd3d/resource.c +@@ -701,6 +701,28 @@ static bool d3d12_device_supports_universal_color_ds_copy(struct d3d12_device *d + return true; + } + ++/* d4r experiment (VKD3D_D4R_LINEAR_TEXTURES=1): committed single-level 2D textures in the formats the ++ * d4r shim can write a DLSS result in are created with linear tiling on dedicated exportable memory. */ ++static bool d4r_linear_texture_candidate(const D3D12_HEAP_PROPERTIES *heap_properties, ++ const D3D12_RESOURCE_DESC1 *desc, const struct d3d12_resource *resource, const struct vkd3d_format *format) ++{ ++ /* Read from the live process environment until it is set: the d4r shim applies its d4r.ini after ++ * this library has loaded. */ ++ static bool enabled; ++ char value[8]; ++ if (!enabled && vkd3d_get_env_var("VKD3D_D4R_LINEAR_TEXTURES", value, sizeof(value)) && value[0] == '1') ++ enabled = true; ++ if (!enabled || !heap_properties || heap_properties->Type != D3D12_HEAP_TYPE_DEFAULT) ++ return false; ++ if (!(resource->flags & VKD3D_RESOURCE_COMMITTED) || (resource->flags & VKD3D_RESOURCE_LINEAR_STAGING_COPY)) ++ return false; ++ if (desc->Dimension != D3D12_RESOURCE_DIMENSION_TEXTURE2D || desc->MipLevels != 1 || ++ desc->DepthOrArraySize != 1 || desc->SampleDesc.Count != 1 || ++ (desc->Flags & D3D12_RESOURCE_FLAG_ALLOW_DEPTH_STENCIL)) ++ return false; ++ return format->vk_format == VK_FORMAT_R16G16B16A16_SFLOAT || format->vk_format == VK_FORMAT_A2B10G10R10_UNORM_PACK32; ++} ++ + static HRESULT vkd3d_get_image_create_info(struct d3d12_device *device, + const D3D12_HEAP_PROPERTIES *heap_properties, + const D3D12_RESOURCE_DESC1 *desc, struct d3d12_resource *resource, +@@ -749,6 +771,14 @@ static HRESULT vkd3d_get_image_create_info(struct d3d12_device *device, + external_info->handleTypes = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT; + vk_prepend_struct(image_info, external_info); + } ++ else if (resource && d4r_linear_texture_candidate(heap_properties, desc, resource, format)) ++ { ++ /* d4r: linear, exportable texels an external API can address in place. */ ++ resource->flags |= VKD3D_RESOURCE_D4R_LINEAR; ++ external_info->sType = VK_STRUCTURE_TYPE_EXTERNAL_MEMORY_IMAGE_CREATE_INFO; ++ external_info->handleTypes = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT; ++ vk_prepend_struct(image_info, external_info); ++ } + + disable_compression = false; + +@@ -905,6 +935,8 @@ static HRESULT vkd3d_get_image_create_info(struct d3d12_device *device, + image_info->mipLevels = min(desc->MipLevels, max_miplevel_count(desc)); + image_info->samples = vk_samples_from_dxgi_sample_desc(&desc->SampleDesc); + image_info->tiling = format->vk_image_tiling; ++ if (resource && (resource->flags & VKD3D_RESOURCE_D4R_LINEAR)) ++ image_info->tiling = VK_IMAGE_TILING_LINEAR; + image_info->initialLayout = VK_IMAGE_LAYOUT_UNDEFINED; + + if (resource && (resource->flags & VKD3D_RESOURCE_COMMITTED) && +@@ -4430,6 +4462,15 @@ HRESULT d3d12_resource_create_committed(struct d3d12_device *device, const D3D12 + #endif + } + ++ if (object->flags & VKD3D_RESOURCE_D4R_LINEAR) ++ { ++ dedicated_requirements.prefersDedicatedAllocation = VK_TRUE; ++ export_info.sType = VK_STRUCTURE_TYPE_EXPORT_MEMORY_ALLOCATE_INFO; ++ export_info.pNext = allocate_info.pNext; ++ export_info.handleTypes = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT; ++ allocate_info.pNext = &export_info; ++ } ++ + /* Requires implies that prefers is also set by spec. */ + if (dedicated_requirements.prefersDedicatedAllocation) + { +diff --git a/libs/vkd3d/vkd3d_private.h b/libs/vkd3d/vkd3d_private.h +index b2f66434..c60ffe00 100644 +--- a/libs/vkd3d/vkd3d_private.h ++++ b/libs/vkd3d/vkd3d_private.h +@@ -1081,7 +1081,8 @@ enum vkd3d_resource_flag + VKD3D_RESOURCE_ZERO_INITIALIZED = (1u << 8), + VKD3D_RESOURCE_RETAINED_GPU_REFERENCE = (1u << 9), + VKD3D_RESOURCE_COPY_QUEUE_COMPATIBLE = (1u << 10), +- VKD3D_RESOURCE_INPUT_ATTACHMENT = (1u << 11) ++ VKD3D_RESOURCE_INPUT_ATTACHMENT = (1u << 11), ++ VKD3D_RESOURCE_D4R_LINEAR = (1u << 12) + }; + + #define VKD3D_INVALID_TILE_INDEX (~0u) +@@ -5480,7 +5481,7 @@ HRESULT d3d12_command_list_split_for_external_wait(struct d3d12_command_list *li + + HRESULT d3d12_command_list_retain_external_resources(struct d3d12_command_list *list, IUnknown *resources); + +-typedef ID3D12DXVKInteropDeviceD4R2 d3d12_dxvk_interop_device_iface; ++typedef ID3D12DXVKInteropDeviceD4R3 d3d12_dxvk_interop_device_iface; + + /* ID3DLowLatencyDevice */ + typedef ID3DLowLatencyDevice d3d_low_latency_device_iface; diff --git a/tools/d3d12_dlss_harness.cpp b/tools/d3d12_dlss_harness.cpp index 586133c..8fa0702 100644 --- a/tools/d3d12_dlss_harness.cpp +++ b/tools/d3d12_dlss_harness.cpp @@ -504,6 +504,13 @@ int main(int argc, char** argv) const bool exposureRGBA32 = std::getenv("D4R_HARNESS_EXPOSURE_RGBA32") != nullptr; const char* replayDir = std::getenv("D4R_HARNESS_REPLAY_DIR"); const bool rgba8 = std::getenv("D4R_HARNESS_RGBA8") != nullptr; + // D4R_HARNESS_OUTPUT_RGB10A2=1: an R10G10B10A2_UNORM output texture (inputs unchanged), as some games pass + const bool rgb10 = std::getenv("D4R_HARNESS_OUTPUT_RGB10A2") != nullptr; + if (rgb10 && (rgba8 || std::getenv("D4R_HARNESS_QUALITY_SCENE") != nullptr)) + { + std::fprintf(stderr, "RGB10A2 output cannot be combined with RGBA8 or the quality scene\n"); + return 2; + } if (rgba8 && (qualityScene || motionScene || jitterScene || replayDir != nullptr)) { std::fprintf(stderr, "RGBA8 format probe requires the static synthetic scene\n"); @@ -627,7 +634,9 @@ int main(int argc, char** argv) ID3D12Resource* exposureTexture = create_texture(1, 1, exposureRGBA32 ? DXGI_FORMAT_R32G32B32A32_FLOAT : DXGI_FORMAT_R32_FLOAT, D3D12_RESOURCE_FLAG_NONE); ID3D12Resource* outputTexture = create_texture(outWidth, outHeight, - rgba8 ? DXGI_FORMAT_R8G8B8A8_UNORM : DXGI_FORMAT_R16G16B16A16_FLOAT, + rgba8 ? DXGI_FORMAT_R8G8B8A8_UNORM + : rgb10 ? DXGI_FORMAT_R10G10B10A2_UNORM + : DXGI_FORMAT_R16G16B16A16_FLOAT, D3D12_RESOURCE_FLAG_ALLOW_UNORDERED_ACCESS); upload(colorTexture, rgba8 ? static_cast(color8.data()) : static_cast(color.data()), resourceWidth * (rgba8 ? 4 : 8), srv); @@ -636,8 +645,10 @@ int main(int argc, char** argv) upload(exposureTexture, exposureRGBA32 ? static_cast(exposureRGBA.data()) : static_cast(&exposureValue), exposureRGBA32 ? 16 : 4, srv); - upload(outputTexture, rgba8 ? static_cast(outputInit8.data()) : static_cast(outputInit.data()), - outWidth * (rgba8 ? 4 : 8), D3D12_RESOURCE_STATE_UNORDERED_ACCESS); + const std::vector outputInit10(rgb10 ? static_cast(outWidth) * outHeight * 4 : 0); + upload(outputTexture, rgba8 ? static_cast(outputInit8.data()) + : rgb10 ? static_cast(outputInit10.data()) : static_cast(outputInit.data()), + outWidth * (rgba8 || rgb10 ? 4 : 8), D3D12_RESOURCE_STATE_UNORDERED_ACCESS); std::printf("synthetic inputs uploaded\n"); trace_stage("synthetic inputs uploaded"); @@ -970,7 +981,7 @@ int main(int argc, char** argv) if (saveFrames) { const std::vector frameOutput = - read_back(outputTexture, D3D12_RESOURCE_STATE_UNORDERED_ACCESS, outWidth * (rgba8 ? 4 : 8)); + read_back(outputTexture, D3D12_RESOURCE_STATE_UNORDERED_ACCESS, outWidth * (rgba8 || rgb10 ? 4 : 8)); const std::string framePath = std::string(argv[2]) + ".frame" + std::to_string(frame); if (FILE* frameFile = std::fopen(framePath.c_str(), "wb")) { @@ -981,7 +992,7 @@ int main(int argc, char** argv) } const std::vector output = read_back(outputTexture, D3D12_RESOURCE_STATE_UNORDERED_ACCESS, - outWidth * (rgba8 ? 4 : 8)); + outWidth * (rgba8 || rgb10 ? 4 : 8)); FILE* file = std::fopen(argv[2], "wb"); if (file != nullptr) { @@ -1041,7 +1052,7 @@ int main(int argc, char** argv) Sleep(frameWaitMs); } const std::vector cycleOutput = read_back(outputTexture, D3D12_RESOURCE_STATE_UNORDERED_ACCESS, - outWidth * (rgba8 ? 4 : 8)); + outWidth * (rgba8 || rgb10 ? 4 : 8)); if (verifyRecreation) { auto& reference = recreationReference[cycle % 2]; diff --git a/tools/d4r_nvngx_shim.cpp b/tools/d4r_nvngx_shim.cpp index 8868f0b..6668beb 100644 --- a/tools/d4r_nvngx_shim.cpp +++ b/tools/d4r_nvngx_shim.cpp @@ -34,6 +34,9 @@ // D4R_SHIM_VRAM_INTEROP=1 keep inputs and output in VRAM (see "VRAM interop") // D4R_SHIM_SPLIT_FRAME=1 with VRAM interop and the d4r vkd3d-proton patch, // present each frame's own DLSS result (see "Split frames") +// D4R_SHIM_INPLACE=1 experimental, with VKD3D_D4R_LINEAR_TEXTURES=1 in the patched vkd3d-proton: the output +// kernel writes the game's output texture itself (no game-side output copy). +// D4R_SHIM_OUTPUT_PACKED=1 lets M's output kernel store R10G10B10A2 texels itself. #define WIDL_EXPLICIT_AGGREGATE_RETURNS #include #include @@ -1603,6 +1606,12 @@ struct FrameParams bool vram = false; // inputs and output stay in VRAM (VRAM interop) bool split = false; // this frame presents its own result (split frames) unsigned motionDilation = 0; // actually applied, worker only + // D4R_SHIM_INPLACE: the game's linear output texture, written without a game-side copy (0 = none) + CudaDevicePtr outputInPlace = 0; + uint32_t outputInPlacePitch = 0; + // D4R_SHIM_OUTPUT_PACKED: the output kernel stores the game's R10G10B10A2 texels itself, so the game-side + // copy (or the in-place texture) takes them without the RGBA16F conversion image and blit + bool packedOutput = false; }; struct FrameTiming @@ -1719,6 +1728,9 @@ struct Feature // D4R_SHIM_OUTPUT_DIRECT: this frame's result is stored by the native output kernel straight into its // destination buffer (worker only), so the array -> buffer copy is skipped. bool outputRedirected = false; + // The previous frame's output kernel was the native one and honoured the redirect (worker writes, game + // thread reads): only then is a frame recorded for packed output. + std::atomic packedReady{false}; bool outputDirectAllowed = false; // the forced preset's output kernel honours the redirect // D4R_SHIM_LINEAR_INPUTS: NGX samples the interop buffers directly through pitch-linear // texture objects (rows padded to 256 bytes), so the buffer -> array copies disappear. @@ -2127,6 +2139,7 @@ struct VulkanInterop ID3D12DXVKInteropDevice1* interop = nullptr; ID3D12DXVKInteropDeviceD4R2* lifetime = nullptr; ID3D12DXVKInteropDeviceD4R* split = nullptr; // only with the d4r vkd3d-proton patch + ID3D12DXVKInteropDeviceD4R3* linear = nullptr; // experimental: linear exportable game textures VkDevice device = VK_NULL_HANDLE; VkPhysicalDeviceMemoryProperties memory = {}; VkPhysicalDevice physical = VK_NULL_HANDLE; @@ -2232,6 +2245,9 @@ static void init_vram_interop() g_vk.split = nullptr; if (FAILED(g.device->QueryInterface(__uuidof(ID3D12DXVKInteropDeviceD4R2), reinterpret_cast(&g_vk.lifetime)))) g_vk.lifetime = nullptr; + if (env_uint("D4R_SHIM_INPLACE", 0) == 0 || + FAILED(g.device->QueryInterface(__uuidof(ID3D12DXVKInteropDeviceD4R3), reinterpret_cast(&g_vk.linear)))) + g_vk.linear = nullptr; g_vk.ready = ok; logf("VRAM interop: %s (VkDevice %p), split frames %s", ok ? "ready" : "missing Vulkan entry points", static_cast(g_vk.device), g_vk.split != nullptr ? "available" : "unavailable (stock vkd3d-proton)"); @@ -2453,6 +2469,52 @@ static size_t vk_texel_bytes(VkFormat format) } } +// D4R_SHIM_INPLACE (experimental): a game texture that vkd3d-proton created linear and exportable is imported +// into CUDA once, so DLSS reads or writes its texels without a game-side copy. Imports are kept for the +// process lifetime (ROCm does not free mapped external memory, see g_vramPool). +struct InPlaceTexture +{ + CudaDevicePtr texels = 0; + size_t rowPitch = 0; +}; +static std::mutex g_inPlaceMutex; +static std::vector> g_inPlaceImports; // VkDeviceMemory -> CUDA mapping + +// Game thread. +static bool describe_in_place(ID3D12Resource* resource, InPlaceTexture& texture) +{ + UINT64 memory = 0, size = 0, offset = 0, pitch = 0; + if (g_vk.linear == nullptr || FAILED(g_vk.linear->GetVulkanLinearImageInfo(resource, &memory, &size, &offset, &pitch))) + return false; + CudaDevicePtr base = 0; + { + std::lock_guard lock(g_inPlaceMutex); + for (const auto& entry : g_inPlaceImports) + if (entry.first == memory) + base = entry.second; + } + if (base == 0) + { + void* external = nullptr; + const int imported = g.worker.call([&] { + return g_vk.import(g_vk.device, memory, size, &base, &external); + }); + if (imported != 0 || base == 0) + { + logf("in-place: importing %llu bytes of texture memory failed (%d)", static_cast(size), imported); + return false; + } + logf("in-place: imported %llu bytes of game texture memory (offset %llu, row pitch %llu)", + static_cast(size), static_cast(offset), + static_cast(pitch)); + std::lock_guard lock(g_inPlaceMutex); + g_inPlaceImports.emplace_back(memory, base); + } + texture.texels = base + offset; + texture.rowPitch = static_cast(pitch); + return true; +} + // How one input or the output moves between its D3D12 texture and a buffer. struct VramCopy { @@ -3781,6 +3843,44 @@ static int synchronize_default_stream(Feature& feature, bool* querySleep = nullp return g.cu.ctxSynchronize(); } +// A frame recorded for packed output whose output kernel then did not pack it (it was not the native one +// after all): the RGBA16F result is converted on the host, rounding as the kernel store and the game-side +// blit do. Slow, and only expected around a kernel change. +static int pack_output_on_host(Feature& feature, CudaMemcpy2D copy, size_t packedPitch, uint32_t frame) +{ + const size_t width = feature.output.width, height = feature.output.height; + std::vector halves(width * height * 4); + std::vector packed(width * height); + const CudaDevicePtr target = copy.dstDevice; + copy.dstMemoryType = CUDA_MEMORY_HOST; + copy.dstHost = halves.data(); + copy.dstPitch = width * 8; + int result = g.cu.memcpy2D(©); + if (result != 0) + return result; + const auto unorm = [](uint16_t half, float scale) { + float value = half_to_float(half); + value = value > 0.0f ? value : 0.0f; + value = value < 1.0f ? value : 1.0f; + return static_cast(value * scale + 0.5f); + }; + for (size_t texel = 0; texel < packed.size(); ++texel) + packed[texel] = unorm(halves[texel * 4], 1023.0f) | unorm(halves[texel * 4 + 1], 1023.0f) << 10 | + unorm(halves[texel * 4 + 2], 1023.0f) << 20 | unorm(halves[texel * 4 + 3], 3.0f) << 30; + CudaMemcpy2D upload = {}; + upload.srcMemoryType = CUDA_MEMORY_HOST; + upload.srcHost = packed.data(); + upload.srcPitch = width * 4; + upload.dstMemoryType = CUDA_MEMORY_DEVICE; + upload.dstDevice = target; + upload.dstPitch = packedPitch; + upload.WidthInBytes = width * 4; + upload.Height = height; + result = g.cu.memcpy2D(&upload); + logf("frame %u: output kernel did not pack its result; converted on the host (%d)", frame, result); + return result; +} + // VRAM interop, worker thread: copies the finished result into a free output // slot's imported buffer and publishes it; evaluate() copies it on to the game. static void publish_vram(Feature* feature, uint32_t frame, const FrameParams& params, FrameTiming& timing, @@ -3796,7 +3896,7 @@ static void publish_vram(Feature* feature, uint32_t frame, const FrameParams& pa frame_retired(feature, frame); return; } - const VramBuffer& buffer = feature->outputs[target].vram; + VramBuffer buffer = feature->outputs[target].vram; const size_t rowBytes = static_cast(feature->output.width) * 8; CudaMemcpy2D copy = {}; copy.srcMemoryType = CUDA_MEMORY_ARRAY; @@ -3806,7 +3906,18 @@ static void publish_vram(Feature* feature, uint32_t frame, const FrameParams& pa copy.dstPitch = rowBytes; copy.WidthInBytes = rowBytes; copy.Height = feature->output.height; - int result = rowBytes * feature->output.height > buffer.bytes ? -1 : feature->outputRedirected ? 0 : copy_2d(copy); + if (params.outputInPlace != 0) + { + // no game-side copy was recorded: a result the output kernel did not redirect goes to the texture here + buffer.device = copy.dstDevice = params.outputInPlace; + copy.dstPitch = params.outputInPlacePitch; + buffer.bytes = static_cast(params.outputInPlacePitch) * feature->output.height; + } + int result = rowBytes / (params.packedOutput ? 2 : 1) * feature->output.height > buffer.bytes ? -1 + : feature->outputRedirected ? 0 + : params.packedOutput ? pack_output_on_host(*feature, copy, params.outputInPlace != 0 ? copy.dstPitch + : rowBytes / 2, frame) + : copy_2d(copy); if (result == 0) { const auto syncStart = timing.enabled ? ProfileClock::now() : ProfileClock::time_point{}; @@ -4021,10 +4132,20 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram (nativeOutput || feature->outputRedirected)) { // split frames write their own slot; the staged path downloads from outputLinear - const size_t rowBytes = static_cast(feature->output.width) * 8, bytes = rowBytes * feature->output.height; + size_t rowBytes = static_cast(feature->output.width) * 8; + const size_t bytes = rowBytes * feature->output.height; CudaDevicePtr target = 0; - if (params.vram && params.split && feature->outputs[frame % kOutputSlots].vram.bytes >= bytes) + if (params.packedOutput) + rowBytes = static_cast(feature->output.width) * 4; + if (params.vram && params.split && params.outputInPlace != 0) + { + target = params.outputInPlace; + rowBytes = params.outputInPlacePitch; + } + else if (params.vram && params.split && feature->outputs[frame % kOutputSlots].vram.bytes >= bytes) target = feature->outputs[frame % kOutputSlots].vram.device; + if (params.packedOutput) + rowBytes |= 1; // the bridge's flag for R10G10B10A2 texels (row pitches are multiples of 8) else if (!params.vram && overlapped && ensure_linear(feature->outputLinear, feature->outputLinearBytes, bytes)) target = feature->outputLinear; const bool redirect = nativeOutput && target != 0 && @@ -4051,6 +4172,7 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram g.cu.setArrayRedirect(feature->output.array, 0, 0); feature->outputRedirected = false; } + feature->packedReady = feature->outputRedirected; const auto evalReturned = timing.enabled ? ProfileClock::now() : ProfileClock::time_point{}; const int eventEndResult = eventStartResult == 0 ? g.cu.eventRecord(feature->profileEnd, nullptr) : -1; const auto syncStart = timing.enabled ? ProfileClock::now() : ProfileClock::time_point{}; @@ -4864,6 +4986,49 @@ static NgxResult evaluate(ID3D12GraphicsCommandList* list, const NgxHandle* hand if (!ensure_vram_buffer(*feature, slot.vram[index], geometry.size())) return NGX_FAIL_PLATFORM_ERROR; } + // M's downsample kernel has only formatted stores, which the native store can pack (kernels/tex/tex_common.h). + p.packedOutput = feature->split && feature->outputDirectAllowed && feature->preset == 13 && vramOutput.convert && + outputDesc.Format == DXGI_FORMAT_R10G10B10A2_UNORM && vramOutput.width % 2 == 0 && + env_uint("D4R_SHIM_OUTPUT_DIRECT", 0) != 0 && env_uint("D4R_SHIM_OUTPUT_PACKED", 0) != 0 && + feature->packedReady.load(); + static bool packedLogged = false; + if (p.packedOutput && !packedLogged) + { + packedLogged = true; + logf("frame %u: packed output (the output kernel stores R10G10B10A2 texels; no conversion blit)", frame); + } + InPlaceTexture outputTexture; + const size_t outputTexelBytes = p.packedOutput ? 4 : 8; + if (feature->split && feature->outputDirectAllowed && (!vramOutput.convert || p.packedOutput) && + env_uint("D4R_SHIM_OUTPUT_DIRECT", 0) != 0 && + describe_in_place(output, outputTexture) && outputTexture.rowPitch % 8 == 0 && + outputTexture.rowPitch / 8 <= 0xffffu && + outputTexture.rowPitch >= static_cast(vramOutput.width) * outputTexelBytes) + { + p.outputInPlace = outputTexture.texels; + p.outputInPlacePitch = static_cast(outputTexture.rowPitch); + } + static bool outputInPlaceLogged = false; + if (p.outputInPlace != 0 && !outputInPlaceLogged) + { + outputInPlaceLogged = true; + logf("frame %u: output in place (the output kernel writes the game's %s texture; no game-side copy)", frame, + p.packedOutput ? "R10G10B10A2" : "RGBA16F"); + } + static bool inPlaceLogged = false; + if (g_vk.linear != nullptr && !inPlaceLogged) + { + inPlaceLogged = true; + if (p.outputInPlace == 0) + { + UINT64 memory = 0, size = 0, offset = 0, pitch = 0; + const HRESULT linear = g_vk.linear->GetVulkanLinearImageInfo(output, &memory, &size, &offset, &pitch); + logf("in-place: output copied on the first frame (split %d, direct output allowed %d, format %d, needs conversion %d, " + "linear exportable texture %d, row pitch %llu)", static_cast(feature->split), feature->outputDirectAllowed, + static_cast(outputDesc.Format), vramOutput.convert, SUCCEEDED(linear), + static_cast(pitch)); + } + } if (vramInputs[0].convert && !ensure_conversion_image(*feature, feature->colorConversion, vramInputs[0].width, vramInputs[0].height, VK_FORMAT_R16G16B16A16_SFLOAT)) @@ -4970,14 +5135,24 @@ static NgxResult evaluate(ID3D12GraphicsCommandList* list, const NgxHandle* hand if (p.vram && feature->split) p.split = SUCCEEDED(g_vk.split->SplitCommandListForExternalWait( list, reinterpret_cast(feature->splitSemaphore), frame)); + if (!p.split) + { + p.outputInPlace = 0; // only a split frame waits for the result before the game reads its texture + p.packedOutput = false; + } if (p.split) { // Runs in the second half of the list, once frame's result is in place. const int target = static_cast(frame % kOutputSlots); - transition(list, output, outputState, D3D12_RESOURCE_STATE_COPY_DEST); - if (!record_vram_output(list, feature->outputs[target].vram, vramOutput, feature->outputConversion)) - logf("frame %u: BeginVkCommandBufferInterop failed for the output", frame); - transition(list, output, D3D12_RESOURCE_STATE_COPY_DEST, outputState); + if (p.outputInPlace == 0) + { + if (p.packedOutput) + vramOutput.convert = false; // the buffer already holds the texture's own texels + transition(list, output, outputState, D3D12_RESOURCE_STATE_COPY_DEST); + if (!record_vram_output(list, feature->outputs[target].vram, vramOutput, feature->outputConversion)) + logf("frame %u: BeginVkCommandBufferInterop failed for the output", frame); + transition(list, output, D3D12_RESOURCE_STATE_COPY_DEST, outputState); + } if (timing.enabled) timing.presentedFrame = frame; } diff --git a/tools/d4r_vkd3d_interop.h b/tools/d4r_vkd3d_interop.h index e636b78..9087d26 100644 --- a/tools/d4r_vkd3d_interop.h +++ b/tools/d4r_vkd3d_interop.h @@ -57,3 +57,13 @@ ID3D12DXVKInteropDeviceD4R2 : public ID3D12DXVKInteropDeviceD4R }; __CRT_UUID_DECL(ID3D12DXVKInteropDeviceD4R2, 0x4f8c9462, 0x8197, 0x4dc1, 0xa5, 0x84, 0xcf, 0x44, 0xd9, 0x70, 0x6d, 0xb2) + +// Experimental (VKD3D_D4R_LINEAR_TEXTURES=1 in the patched vkd3d-proton): the dedicated, exportable memory +// behind a linear-tiled texture and the layout of its only subresource. +MIDL_INTERFACE("7b1e5d0a-93c4-4e8f-b6a2-51c08f2e4d73") +ID3D12DXVKInteropDeviceD4R3 : public ID3D12DXVKInteropDeviceD4R2 +{ + virtual HRESULT STDMETHODCALLTYPE GetVulkanLinearImageInfo(ID3D12Resource* resource, UINT64* memory, UINT64* memorySize, + UINT64* offset, UINT64* rowPitch) = 0; +}; +__CRT_UUID_DECL(ID3D12DXVKInteropDeviceD4R3, 0x7b1e5d0a, 0x93c4, 0x4e8f, 0xb6, 0xa2, 0x51, 0xc0, 0x8f, 0x2e, 0x4d, 0x73) diff --git a/tools/d4r_win32_wait.h b/tools/d4r_win32_wait.h index 56a28d3..d4d86c9 100644 --- a/tools/d4r_win32_wait.h +++ b/tools/d4r_win32_wait.h @@ -9,7 +9,10 @@ inline void d4r_sleep_us(unsigned int microseconds) using Delay = LONG(WINAPI*)(BOOLEAN, LARGE_INTEGER*); static const auto delay = reinterpret_cast(reinterpret_cast( GetProcAddress(GetModuleHandleW(L"ntdll.dll"), "NtDelayExecution"))); - if (microseconds == 0) + // D4R_SHIM_SPIN_POLL=1 (measurement only): every poll yields instead of sleeping, so frame timings are + // not quantised to the ~250 us a 200 us sleep takes. Costs a CPU core per waiting thread. + static const bool spin = GetEnvironmentVariableA("D4R_SHIM_SPIN_POLL", nullptr, 0) > 1; + if (microseconds == 0 || spin) { SwitchToThread(); return; diff --git a/tools/wine_nvcuda_bridge.c b/tools/wine_nvcuda_bridge.c index db9b33e..f48bde3 100644 --- a/tools/wine_nvcuda_bridge.c +++ b/tools/wine_nvcuda_bridge.c @@ -1201,8 +1201,9 @@ static CUresult write_redirect_tail(CUsurfObject object, CUdeviceptr pointer, ui static unsigned char ring[256][16]; static unsigned int next; unsigned char* tail = ring[next++ % 256]; - // dword 21: 0x5232 ('R2') << 16 | pitch / 8 (0 = off); dwords 22-23: pointer - const uint32_t word = pitch != 0 ? 0x52320000u | ((pitch >> 3) & 0xffffu) : 0u; + // dword 21: 0x5232 ('R2', RGBA16F texels) << 16 | pitch / 8 (0 = off); dwords 22-23: pointer. + // Row pitches are multiples of 8, so bit 0 of `pitch` selects 'R3': R10G10B10A2_UNORM texels. + const uint32_t word = pitch != 0 ? ((pitch & 1u) != 0 ? 0x52330000u : 0x52320000u) | ((pitch >> 3) & 0xffffu) : 0u; memset(tail, 0, 16); memcpy(tail, &word, 4); memcpy(tail + 4, &pointer, 8);