Merge PR #27: experimental packed and in-place DLSS output
This commit is contained in:
woutermans committed
commit
84e6048066
11 files changed
+1012
-22
No files matched your search
@@ -0,0 +1,53 @@
|
|||||||
|
# Reading the DLSS inputs in place: findings
|
||||||
|
|
||||||
|
The colour, depth and motion inputs are copied from the game's textures into buffers shared with HIP on every frame. This page records an attempt to remove those copies by letting DLSS read the game's textures directly, why it does not work, and what was ruled out. None of the code for it is in the tree; the output-side work it grew out of is described in [performance.md](performance.md).
|
||||||
|
|
||||||
|
All of this was measured in the D3D12 harness on an RX 9070 XT (gfx1201) with ROCm 7.2, Mesa 26.2 (RADV) and Linux 7.2, preset M at 1280×720 → 3840×2160.
|
||||||
|
|
||||||
|
## The fault
|
||||||
|
|
||||||
|
With vkd3d-proton patch 0003 a game texture can be created linear on exportable memory and imported into HIP, exactly as the shim's own buffers are. A pointer to its texels is then available to CUDA. Three ways of reading through that pointer give different answers:
|
||||||
|
|
||||||
|
| Read | Result |
|
||||||
|
|---|---|
|
||||||
|
| `cuMemcpy2D` | current contents, every frame |
|
||||||
|
| A compute kernel using plain loads | current contents, 0 of 7,372,800 bytes differ |
|
||||||
|
| A compute kernel using a pitch-linear texture object | wrong from the first frame |
|
||||||
|
|
||||||
|
DLSS's kernels fetch their inputs through texture objects, so DLSS ran on wrong inputs. A probe kernel showed what the texture fetch returns: the image the texture held at its first upload in 81% of texels, with 32 bytes of unrelated data at every 256-byte step. That is the look of memory that has been freed and reused, not of an older version of the texture.
|
||||||
|
|
||||||
|
Writing through the same kind of mapping is not affected: the output kernel writes the game's output texture with plain stores, and Vulkan sees the result.
|
||||||
|
|
||||||
|
## What was ruled out
|
||||||
|
|
||||||
|
- **NGX and ZLUDA.** The probe is wrong before NGX evaluates anything, and ZLUDA passes texture creation straight to HIP.
|
||||||
|
- **A copy inside ROCm.** ROCm 7.2 creates the texture's image view directly on the imported pointer (`rocmemory.cpp`, `Image::createView`); the workaround path that would allocate a copy is disabled for HIP.
|
||||||
|
- **A wrong descriptor.** The hardware image descriptor stored in a bad texture object is identical to that of a working one apart from its address, and the address is the imported pointer.
|
||||||
|
- **The texture having moved.** The kernel driver's debug files (`amdgpu_gem_info`, `amdgpu_vm_info`) show the texture's memory in VRAM, exported, with an up-to-date mapping from before the import to the end of the run. Vulkan and HIP share one GPU address space in the process.
|
||||||
|
- **Stale handles or mappings.** Recreating the texture object every frame, and releasing and re-importing the memory at a new address every frame, change nothing.
|
||||||
|
- **Memory type and allocation style.** CPU-visible and CPU-invisible memory types behave the same, and a non-dedicated exportable allocation behaves the same.
|
||||||
|
- **A general ROCm, kernel or RADV bug.** `tools/hip_vk_texture_stale_repro.cpp` performs the same sequence natively (Vulkan object on exportable memory, upload, import into HIP, texture object, further uploads) and texture sampling follows Vulkan's writes for a linear image and a buffer, every memory type, one to four uploads before the import, game-like usage flags and zero-initialised memory.
|
||||||
|
|
||||||
|
So the same address in the same address space gives current data to a plain load and other data to a hardware texture fetch, only for textures created by vkd3d-proton inside the Wine process. The cause is not known. The untested differences from the native program are vkd3d-proton's image creation chain (alignment control, compression control, format lists), its allocation call as issued through Wine, and the path its uploads take in the driver.
|
||||||
|
|
||||||
|
A warning for anyone repeating this: copying texels out and back with `cuMemcpy2D` makes exactly the rewritten rows visible to the texture fetch, which means that copy writes into whatever memory the fetch is reading. It was done a few times during this investigation without visible effect, but it should not be repeated.
|
||||||
|
|
||||||
|
## The workaround that works, and why it was dropped
|
||||||
|
|
||||||
|
Since plain loads are correct, the fetches themselves can be replaced. For preset M the inputs are read by three kernels only: enc0 (colour, motion), post (colour, motion) and downsample (colour); depth is passed to enc0 but not fetched in the `mvhi_ldr` variant. All 20 of those fetches are integer texel fetches (`tex.base.2d.v4.f32.s32`), so no filtering is involved. `make_ptx.py` was extended to trace each fetch's texture handle back to the kernel's parameter block (enc0 offsets 80, 96, 136; post 0, 176; downsample 56) and to replace the input fetches with a load through the pointer HIP keeps in the texture object (`__hip_texture`: resource type at byte 96, pointer at 104, channel description at 112, width 136, height 144, pitch 152). An integer fetch outside the image returns zeros, not the edge texel; motion depends on that.
|
||||||
|
|
||||||
|
With colour, depth and motion read in place this way and the output written in place, no frame-sized copy was left and the output of a 48-frame panning scene was byte-identical. It was not faster:
|
||||||
|
|
||||||
|
| Configuration | Frame interval | DLSS kernel time |
|
||||||
|
|---|---|---|
|
||||||
|
| Copies, original kernels | 5.91 ms | 3.09 ms |
|
||||||
|
| Output in place only | 5.76–5.89 ms | 3.15–3.16 ms |
|
||||||
|
| Inputs and output in place | 5.89–5.93 ms | 3.27–3.30 ms |
|
||||||
|
|
||||||
|
The rewritten fetches add about 0.11 ms of kernel time, and about 0.07 ms even when unused because each fetch first checks whether the texture is tagged. The three input copies cost roughly 0.05–0.1 ms. The two cancel, so the code was removed.
|
||||||
|
|
||||||
|
## If this is picked up again
|
||||||
|
|
||||||
|
- The zero-cost fix is a working hardware fetch. The next step is to make the native reproducer fail by adding vkd3d-proton's behaviour to it piece by piece, or by running it through Wine's Vulkan layer.
|
||||||
|
- A cheaper workaround would need the per-fetch check removed (a separate kernel set chosen when the inputs are in place); a plain load plus conversion is still unlikely to beat a hardware fetch.
|
||||||
|
- Real games pass depth in a depth format, which RADV is not expected to offer with linear tiling (not checked here; the harness uses R32F), so depth would probably keep its copy either way.
|
||||||
@@ -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.
|
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
|
## What did not help
|
||||||
|
|
||||||
- **Accumulating in f16 on the WMMA units.** The error grows about 20× and the image degrades.
|
- **Accumulating in f16 on the WMMA units.** The error grows about 20× and the image degrades.
|
||||||
|
|||||||
@@ -338,9 +338,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);
|
extern "C" __attribute__((device)) int __ockl_image_height_2D(tsharp_t* image);
|
||||||
__attribute__((device)) static inline uint32_t redirect_pitch(uint64_t surface)
|
__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];
|
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
|
// D4R_AS1_STORES: redirect stores through address-space-1 pointers (global_store instead of flat_store; the
|
||||||
// redirect target is linear device memory)
|
// redirect target is linear device memory)
|
||||||
@@ -371,7 +385,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 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)));
|
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;
|
return;
|
||||||
}
|
}
|
||||||
@@ -474,8 +493,9 @@ DEV __attribute__((always_inline)) void d4r_sust_b32(uint64_t surface, int32_t x
|
|||||||
tsharp_t* image = (tsharp_t*)surface;
|
tsharp_t* image = (tsharp_t*)surface;
|
||||||
if (redirect_pitch(surface) != 0)
|
if (redirect_pitch(surface) != 0)
|
||||||
{
|
{
|
||||||
// raw bits at byte offset x (RGBA16F rows are 8 bytes per pixel)
|
// raw bits at byte offset x (RGBA16F rows are 8 bytes per pixel); the shim only asks for packed
|
||||||
if (x >= 0 && (x & 3) == 0 && redirect_inside(surface, x >> 3, y))
|
// 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;
|
*D4R_REDIRECT_PTR(uint32_t, redirect_row(surface, y) + x) = data;
|
||||||
return;
|
return;
|
||||||
}
|
}
|
||||||
|
|||||||
@@ -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;
|
||||||
Executable
+11
@@ -0,0 +1,11 @@
|
|||||||
|
#!/usr/bin/env bash
|
||||||
|
set -euo pipefail
|
||||||
|
|
||||||
|
# Builds build/hip_vk_texture_stale_repro: a native Vulkan + HIP check that a HIP texture object on
|
||||||
|
# memory imported from Vulkan follows Vulkan's later writes (see tools/hip_vk_texture_stale_repro.cpp).
|
||||||
|
ROOT="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)"
|
||||||
|
ROCM_ROOT="${ROCM_ROOT:-/opt/rocm}"
|
||||||
|
mkdir -p "$ROOT/build"
|
||||||
|
"$ROCM_ROOT/bin/hipcc" -std=c++17 -O2 "$ROOT/tools/hip_vk_texture_stale_repro.cpp" -lvulkan \
|
||||||
|
-o "$ROOT/build/hip_vk_texture_stale_repro"
|
||||||
|
printf 'Built %s (run with LD_LIBRARY_PATH=%s/lib:%s/lib/llvm/lib)\n' "$ROOT/build/hip_vk_texture_stale_repro" "$ROCM_ROOT" "$ROCM_ROOT"
|
||||||
@@ -504,6 +504,13 @@ int main(int argc, char** argv)
|
|||||||
const bool exposureRGBA32 = std::getenv("D4R_HARNESS_EXPOSURE_RGBA32") != nullptr;
|
const bool exposureRGBA32 = std::getenv("D4R_HARNESS_EXPOSURE_RGBA32") != nullptr;
|
||||||
const char* replayDir = std::getenv("D4R_HARNESS_REPLAY_DIR");
|
const char* replayDir = std::getenv("D4R_HARNESS_REPLAY_DIR");
|
||||||
const bool rgba8 = std::getenv("D4R_HARNESS_RGBA8") != nullptr;
|
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))
|
if (rgba8 && (qualityScene || motionScene || jitterScene || replayDir != nullptr))
|
||||||
{
|
{
|
||||||
std::fprintf(stderr, "RGBA8 format probe requires the static synthetic scene\n");
|
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,
|
ID3D12Resource* exposureTexture = create_texture(1, 1,
|
||||||
exposureRGBA32 ? DXGI_FORMAT_R32G32B32A32_FLOAT : DXGI_FORMAT_R32_FLOAT, D3D12_RESOURCE_FLAG_NONE);
|
exposureRGBA32 ? DXGI_FORMAT_R32G32B32A32_FLOAT : DXGI_FORMAT_R32_FLOAT, D3D12_RESOURCE_FLAG_NONE);
|
||||||
ID3D12Resource* outputTexture = create_texture(outWidth, outHeight,
|
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);
|
D3D12_RESOURCE_FLAG_ALLOW_UNORDERED_ACCESS);
|
||||||
upload(colorTexture, rgba8 ? static_cast<const void*>(color8.data()) : static_cast<const void*>(color.data()),
|
upload(colorTexture, rgba8 ? static_cast<const void*>(color8.data()) : static_cast<const void*>(color.data()),
|
||||||
resourceWidth * (rgba8 ? 4 : 8), srv);
|
resourceWidth * (rgba8 ? 4 : 8), srv);
|
||||||
@@ -636,8 +645,10 @@ int main(int argc, char** argv)
|
|||||||
upload(exposureTexture, exposureRGBA32 ? static_cast<const void*>(exposureRGBA.data())
|
upload(exposureTexture, exposureRGBA32 ? static_cast<const void*>(exposureRGBA.data())
|
||||||
: static_cast<const void*>(&exposureValue),
|
: static_cast<const void*>(&exposureValue),
|
||||||
exposureRGBA32 ? 16 : 4, srv);
|
exposureRGBA32 ? 16 : 4, srv);
|
||||||
upload(outputTexture, rgba8 ? static_cast<const void*>(outputInit8.data()) : static_cast<const void*>(outputInit.data()),
|
const std::vector<uint8_t> outputInit10(rgb10 ? static_cast<size_t>(outWidth) * outHeight * 4 : 0);
|
||||||
outWidth * (rgba8 ? 4 : 8), D3D12_RESOURCE_STATE_UNORDERED_ACCESS);
|
upload(outputTexture, rgba8 ? static_cast<const void*>(outputInit8.data())
|
||||||
|
: rgb10 ? static_cast<const void*>(outputInit10.data()) : static_cast<const void*>(outputInit.data()),
|
||||||
|
outWidth * (rgba8 || rgb10 ? 4 : 8), D3D12_RESOURCE_STATE_UNORDERED_ACCESS);
|
||||||
std::printf("synthetic inputs uploaded\n");
|
std::printf("synthetic inputs uploaded\n");
|
||||||
trace_stage("synthetic inputs uploaded");
|
trace_stage("synthetic inputs uploaded");
|
||||||
|
|
||||||
@@ -970,7 +981,7 @@ int main(int argc, char** argv)
|
|||||||
if (saveFrames)
|
if (saveFrames)
|
||||||
{
|
{
|
||||||
const std::vector<uint8_t> frameOutput =
|
const std::vector<uint8_t> 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);
|
const std::string framePath = std::string(argv[2]) + ".frame" + std::to_string(frame);
|
||||||
if (FILE* frameFile = std::fopen(framePath.c_str(), "wb"))
|
if (FILE* frameFile = std::fopen(framePath.c_str(), "wb"))
|
||||||
{
|
{
|
||||||
@@ -981,7 +992,7 @@ int main(int argc, char** argv)
|
|||||||
}
|
}
|
||||||
|
|
||||||
const std::vector<uint8_t> output = read_back(outputTexture, D3D12_RESOURCE_STATE_UNORDERED_ACCESS,
|
const std::vector<uint8_t> 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");
|
FILE* file = std::fopen(argv[2], "wb");
|
||||||
if (file != nullptr)
|
if (file != nullptr)
|
||||||
{
|
{
|
||||||
@@ -1041,7 +1052,7 @@ int main(int argc, char** argv)
|
|||||||
Sleep(frameWaitMs);
|
Sleep(frameWaitMs);
|
||||||
}
|
}
|
||||||
const std::vector<uint8_t> cycleOutput = read_back(outputTexture, D3D12_RESOURCE_STATE_UNORDERED_ACCESS,
|
const std::vector<uint8_t> cycleOutput = read_back(outputTexture, D3D12_RESOURCE_STATE_UNORDERED_ACCESS,
|
||||||
outWidth * (rgba8 ? 4 : 8));
|
outWidth * (rgba8 || rgb10 ? 4 : 8));
|
||||||
if (verifyRecreation)
|
if (verifyRecreation)
|
||||||
{
|
{
|
||||||
auto& reference = recreationReference[cycle % 2];
|
auto& reference = recreationReference[cycle % 2];
|
||||||
|
|||||||
+183
-8
@@ -34,6 +34,9 @@
|
|||||||
// D4R_SHIM_VRAM_INTEROP=1 keep inputs and output in VRAM (see "VRAM interop")
|
// 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,
|
// D4R_SHIM_SPLIT_FRAME=1 with VRAM interop and the d4r vkd3d-proton patch,
|
||||||
// present each frame's own DLSS result (see "Split frames")
|
// 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
|
#define WIDL_EXPLICIT_AGGREGATE_RETURNS
|
||||||
#include <windows.h>
|
#include <windows.h>
|
||||||
#include <d3d12.h>
|
#include <d3d12.h>
|
||||||
@@ -1611,6 +1614,12 @@ struct FrameParams
|
|||||||
bool vram = false; // inputs and output stay in VRAM (VRAM interop)
|
bool vram = false; // inputs and output stay in VRAM (VRAM interop)
|
||||||
bool split = false; // this frame presents its own result (split frames)
|
bool split = false; // this frame presents its own result (split frames)
|
||||||
unsigned motionDilation = 0; // actually applied, worker only
|
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
|
struct FrameTiming
|
||||||
@@ -1732,6 +1741,9 @@ struct Feature
|
|||||||
// D4R_SHIM_OUTPUT_DIRECT: this frame's result is stored by the native output kernel straight into its
|
// 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.
|
// destination buffer (worker only), so the array -> buffer copy is skipped.
|
||||||
bool outputRedirected = false;
|
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<bool> packedReady{false};
|
||||||
bool outputDirectAllowed = false; // the forced preset's output kernel honours the redirect
|
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
|
// 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.
|
// texture objects (rows padded to 256 bytes), so the buffer -> array copies disappear.
|
||||||
@@ -2140,6 +2152,7 @@ struct VulkanInterop
|
|||||||
ID3D12DXVKInteropDevice1* interop = nullptr;
|
ID3D12DXVKInteropDevice1* interop = nullptr;
|
||||||
ID3D12DXVKInteropDeviceD4R2* lifetime = nullptr;
|
ID3D12DXVKInteropDeviceD4R2* lifetime = nullptr;
|
||||||
ID3D12DXVKInteropDeviceD4R* split = nullptr; // only with the d4r vkd3d-proton patch
|
ID3D12DXVKInteropDeviceD4R* split = nullptr; // only with the d4r vkd3d-proton patch
|
||||||
|
ID3D12DXVKInteropDeviceD4R3* linear = nullptr; // experimental: linear exportable game textures
|
||||||
VkDevice device = VK_NULL_HANDLE;
|
VkDevice device = VK_NULL_HANDLE;
|
||||||
VkPhysicalDeviceMemoryProperties memory = {};
|
VkPhysicalDeviceMemoryProperties memory = {};
|
||||||
VkPhysicalDevice physical = VK_NULL_HANDLE;
|
VkPhysicalDevice physical = VK_NULL_HANDLE;
|
||||||
@@ -2245,6 +2258,9 @@ static void init_vram_interop()
|
|||||||
g_vk.split = nullptr;
|
g_vk.split = nullptr;
|
||||||
if (FAILED(g.device->QueryInterface(__uuidof(ID3D12DXVKInteropDeviceD4R2), reinterpret_cast<void**>(&g_vk.lifetime))))
|
if (FAILED(g.device->QueryInterface(__uuidof(ID3D12DXVKInteropDeviceD4R2), reinterpret_cast<void**>(&g_vk.lifetime))))
|
||||||
g_vk.lifetime = nullptr;
|
g_vk.lifetime = nullptr;
|
||||||
|
if (env_uint("D4R_SHIM_INPLACE", 0) == 0 ||
|
||||||
|
FAILED(g.device->QueryInterface(__uuidof(ID3D12DXVKInteropDeviceD4R3), reinterpret_cast<void**>(&g_vk.linear))))
|
||||||
|
g_vk.linear = nullptr;
|
||||||
g_vk.ready = ok;
|
g_vk.ready = ok;
|
||||||
logf("VRAM interop: %s (VkDevice %p), split frames %s", ok ? "ready" : "missing Vulkan entry points",
|
logf("VRAM interop: %s (VkDevice %p), split frames %s", ok ? "ready" : "missing Vulkan entry points",
|
||||||
static_cast<void*>(g_vk.device), g_vk.split != nullptr ? "available" : "unavailable (stock vkd3d-proton)");
|
static_cast<void*>(g_vk.device), g_vk.split != nullptr ? "available" : "unavailable (stock vkd3d-proton)");
|
||||||
@@ -2470,6 +2486,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<std::pair<uint64_t, CudaDevicePtr>> 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<std::mutex> 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<unsigned long long>(size), imported);
|
||||||
|
return false;
|
||||||
|
}
|
||||||
|
logf("in-place: imported %llu bytes of game texture memory (offset %llu, row pitch %llu)",
|
||||||
|
static_cast<unsigned long long>(size), static_cast<unsigned long long>(offset),
|
||||||
|
static_cast<unsigned long long>(pitch));
|
||||||
|
std::lock_guard<std::mutex> lock(g_inPlaceMutex);
|
||||||
|
g_inPlaceImports.emplace_back(memory, base);
|
||||||
|
}
|
||||||
|
texture.texels = base + offset;
|
||||||
|
texture.rowPitch = static_cast<size_t>(pitch);
|
||||||
|
return true;
|
||||||
|
}
|
||||||
|
|
||||||
// How one input or the output moves between its D3D12 texture and a buffer.
|
// How one input or the output moves between its D3D12 texture and a buffer.
|
||||||
struct VramCopy
|
struct VramCopy
|
||||||
{
|
{
|
||||||
@@ -3858,6 +3920,44 @@ static int synchronize_default_stream(Feature& feature, bool* querySleep = nullp
|
|||||||
return g.cu.ctxSynchronize();
|
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<uint16_t> halves(width * height * 4);
|
||||||
|
std::vector<uint32_t> 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<uint32_t>(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
|
// 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.
|
// 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,
|
static void publish_vram(Feature* feature, uint32_t frame, const FrameParams& params, FrameTiming& timing,
|
||||||
@@ -3873,7 +3973,7 @@ static void publish_vram(Feature* feature, uint32_t frame, const FrameParams& pa
|
|||||||
frame_retired(feature, frame);
|
frame_retired(feature, frame);
|
||||||
return;
|
return;
|
||||||
}
|
}
|
||||||
const VramBuffer& buffer = feature->outputs[target].vram;
|
VramBuffer buffer = feature->outputs[target].vram;
|
||||||
const size_t rowBytes = static_cast<size_t>(feature->output.width) * 8;
|
const size_t rowBytes = static_cast<size_t>(feature->output.width) * 8;
|
||||||
CudaMemcpy2D copy = {};
|
CudaMemcpy2D copy = {};
|
||||||
copy.srcMemoryType = CUDA_MEMORY_ARRAY;
|
copy.srcMemoryType = CUDA_MEMORY_ARRAY;
|
||||||
@@ -3883,7 +3983,18 @@ static void publish_vram(Feature* feature, uint32_t frame, const FrameParams& pa
|
|||||||
copy.dstPitch = rowBytes;
|
copy.dstPitch = rowBytes;
|
||||||
copy.WidthInBytes = rowBytes;
|
copy.WidthInBytes = rowBytes;
|
||||||
copy.Height = feature->output.height;
|
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<size_t>(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)
|
if (result == 0)
|
||||||
{
|
{
|
||||||
const auto syncStart = timing.enabled ? ProfileClock::now() : ProfileClock::time_point{};
|
const auto syncStart = timing.enabled ? ProfileClock::now() : ProfileClock::time_point{};
|
||||||
@@ -4100,10 +4211,20 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram
|
|||||||
(nativeOutput || feature->outputRedirected))
|
(nativeOutput || feature->outputRedirected))
|
||||||
{
|
{
|
||||||
// split frames write their own slot; the staged path downloads from outputLinear
|
// split frames write their own slot; the staged path downloads from outputLinear
|
||||||
const size_t rowBytes = static_cast<size_t>(feature->output.width) * 8, bytes = rowBytes * feature->output.height;
|
size_t rowBytes = static_cast<size_t>(feature->output.width) * 8;
|
||||||
|
const size_t bytes = rowBytes * feature->output.height;
|
||||||
CudaDevicePtr target = 0;
|
CudaDevicePtr target = 0;
|
||||||
if (params.vram && params.split && feature->outputs[frame % kOutputSlots].vram.bytes >= bytes)
|
if (params.packedOutput)
|
||||||
|
rowBytes = static_cast<size_t>(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;
|
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))
|
else if (!params.vram && overlapped && ensure_linear(feature->outputLinear, feature->outputLinearBytes, bytes))
|
||||||
target = feature->outputLinear;
|
target = feature->outputLinear;
|
||||||
const bool redirect = nativeOutput && target != 0 &&
|
const bool redirect = nativeOutput && target != 0 &&
|
||||||
@@ -4138,6 +4259,7 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram
|
|||||||
g.cu.setArrayRedirect(feature->output.array, 0, 0);
|
g.cu.setArrayRedirect(feature->output.array, 0, 0);
|
||||||
feature->outputRedirected = false;
|
feature->outputRedirected = false;
|
||||||
}
|
}
|
||||||
|
feature->packedReady = feature->outputRedirected;
|
||||||
const auto evalReturned = timing.enabled ? ProfileClock::now() : ProfileClock::time_point{};
|
const auto evalReturned = timing.enabled ? ProfileClock::now() : ProfileClock::time_point{};
|
||||||
const int eventEndResult = eventStartResult == 0 ? g.cu.eventRecord(feature->profileEnd, nullptr) : -1;
|
const int eventEndResult = eventStartResult == 0 ? g.cu.eventRecord(feature->profileEnd, nullptr) : -1;
|
||||||
const auto syncStart = timing.enabled ? ProfileClock::now() : ProfileClock::time_point{};
|
const auto syncStart = timing.enabled ? ProfileClock::now() : ProfileClock::time_point{};
|
||||||
@@ -4981,6 +5103,49 @@ static NgxResult evaluate(ID3D12GraphicsCommandList* list, const NgxHandle* hand
|
|||||||
if (!ensure_vram_buffer(*feature, slot.vram[index], geometry.size()))
|
if (!ensure_vram_buffer(*feature, slot.vram[index], geometry.size()))
|
||||||
return NGX_FAIL_PLATFORM_ERROR;
|
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<size_t>(vramOutput.width) * outputTexelBytes)
|
||||||
|
{
|
||||||
|
p.outputInPlace = outputTexture.texels;
|
||||||
|
p.outputInPlacePitch = static_cast<uint32_t>(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<int>(feature->split), feature->outputDirectAllowed,
|
||||||
|
static_cast<int>(outputDesc.Format), vramOutput.convert, SUCCEEDED(linear),
|
||||||
|
static_cast<unsigned long long>(pitch));
|
||||||
|
}
|
||||||
|
}
|
||||||
if (vramInputs[0].convert && !ensure_conversion_image(*feature, feature->colorConversion,
|
if (vramInputs[0].convert && !ensure_conversion_image(*feature, feature->colorConversion,
|
||||||
vramInputs[0].width, vramInputs[0].height,
|
vramInputs[0].width, vramInputs[0].height,
|
||||||
VK_FORMAT_R16G16B16A16_SFLOAT))
|
VK_FORMAT_R16G16B16A16_SFLOAT))
|
||||||
@@ -5091,14 +5256,24 @@ static NgxResult evaluate(ID3D12GraphicsCommandList* list, const NgxHandle* hand
|
|||||||
if (p.vram && feature->split)
|
if (p.vram && feature->split)
|
||||||
p.split = SUCCEEDED(g_vk.split->SplitCommandListForExternalWait(
|
p.split = SUCCEEDED(g_vk.split->SplitCommandListForExternalWait(
|
||||||
list, reinterpret_cast<UINT64>(feature->splitSemaphore), frame));
|
list, reinterpret_cast<UINT64>(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)
|
if (p.split)
|
||||||
{
|
{
|
||||||
// Runs in the second half of the list, once frame's result is in place.
|
// Runs in the second half of the list, once frame's result is in place.
|
||||||
const int target = static_cast<int>(frame % kOutputSlots);
|
const int target = static_cast<int>(frame % kOutputSlots);
|
||||||
transition(list, output, outputState, D3D12_RESOURCE_STATE_COPY_DEST);
|
if (p.outputInPlace == 0)
|
||||||
if (!record_vram_output(list, feature->outputs[target].vram, vramOutput, feature->outputConversion))
|
{
|
||||||
logf("frame %u: BeginVkCommandBufferInterop failed for the output", frame);
|
if (p.packedOutput)
|
||||||
transition(list, output, D3D12_RESOURCE_STATE_COPY_DEST, outputState);
|
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)
|
if (timing.enabled)
|
||||||
timing.presentedFrame = frame;
|
timing.presentedFrame = frame;
|
||||||
}
|
}
|
||||||
|
|||||||
@@ -57,3 +57,13 @@ ID3D12DXVKInteropDeviceD4R2 : public ID3D12DXVKInteropDeviceD4R
|
|||||||
};
|
};
|
||||||
__CRT_UUID_DECL(ID3D12DXVKInteropDeviceD4R2, 0x4f8c9462, 0x8197, 0x4dc1, 0xa5, 0x84, 0xcf, 0x44, 0xd9, 0x70, 0x6d, 0xb2)
|
__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)
|
||||||
@@ -9,7 +9,10 @@ inline void d4r_sleep_us(unsigned int microseconds)
|
|||||||
using Delay = LONG(WINAPI*)(BOOLEAN, LARGE_INTEGER*);
|
using Delay = LONG(WINAPI*)(BOOLEAN, LARGE_INTEGER*);
|
||||||
static const auto delay = reinterpret_cast<Delay>(reinterpret_cast<void*>(
|
static const auto delay = reinterpret_cast<Delay>(reinterpret_cast<void*>(
|
||||||
GetProcAddress(GetModuleHandleW(L"ntdll.dll"), "NtDelayExecution")));
|
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();
|
SwitchToThread();
|
||||||
return;
|
return;
|
||||||
|
|||||||
@@ -0,0 +1,479 @@
|
|||||||
|
// Standalone check (no Wine, no DLSS): does a HIP texture object created on memory imported from Vulkan
|
||||||
|
// see what Vulkan writes to that memory afterwards?
|
||||||
|
//
|
||||||
|
// 1. Vulkan creates a linear-tiled RGBA16F image (or, with --buffer, a plain buffer) on exportable memory
|
||||||
|
// and uploads pattern A.
|
||||||
|
// 2. HIP imports the memory, creates a pitch-linear texture object on it and samples it in a kernel.
|
||||||
|
// 3. Vulkan uploads pattern B. HIP samples again, and also reads the memory with hipMemcpy.
|
||||||
|
// 4. HIP copies the texels out and back (a rewrite with their own bytes) and samples once more.
|
||||||
|
//
|
||||||
|
// The d4r NGX shim sees step 3's sample return pattern A for game textures while hipMemcpy returns B.
|
||||||
|
// usage: hip_vk_texture_stale_repro [--buffer] [--hip-first] [--prewrites N] [--game-usage] [WIDTH HEIGHT]
|
||||||
|
// --buffer export a VkBuffer (written with vkCmdCopyBuffer) instead of a linear VkImage
|
||||||
|
// --hip-first initialise HIP before the Vulkan object exists (default: after pattern A is uploaded)
|
||||||
|
// --prewrites N Vulkan uploads N times (ending on A) before HIP imports the memory (default 1)
|
||||||
|
// --game-usage create the image as vkd3d-proton would a render target: more usages, mutable format
|
||||||
|
#include <hip/hip_runtime.h>
|
||||||
|
#include <vulkan/vulkan.h>
|
||||||
|
|
||||||
|
#include <cstdio>
|
||||||
|
#include <cstdlib>
|
||||||
|
#include <cstring>
|
||||||
|
#include <unistd.h>
|
||||||
|
#include <vector>
|
||||||
|
|
||||||
|
__global__ void sample_texture(hipTextureObject_t texture, float* out, int width, int height)
|
||||||
|
{
|
||||||
|
const int x = blockIdx.x * blockDim.x + threadIdx.x, y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||||
|
if (x >= width || y >= height)
|
||||||
|
return;
|
||||||
|
const float4 texel = tex2D<float4>(texture, x + 0.5f, y + 0.5f);
|
||||||
|
float* pixel = out + (static_cast<size_t>(y) * width + x) * 4;
|
||||||
|
pixel[0] = texel.x, pixel[1] = texel.y, pixel[2] = texel.z, pixel[3] = texel.w;
|
||||||
|
}
|
||||||
|
|
||||||
|
#define VK_CHECK(call) \
|
||||||
|
do \
|
||||||
|
{ \
|
||||||
|
const VkResult vkResult_ = (call); \
|
||||||
|
if (vkResult_ != VK_SUCCESS) \
|
||||||
|
{ \
|
||||||
|
std::fprintf(stderr, "%s failed: %d\n", #call, vkResult_); \
|
||||||
|
std::exit(1); \
|
||||||
|
} \
|
||||||
|
} while (0)
|
||||||
|
#define HIP_CHECK(call) \
|
||||||
|
do \
|
||||||
|
{ \
|
||||||
|
const hipError_t hipResult_ = (call); \
|
||||||
|
if (hipResult_ != hipSuccess) \
|
||||||
|
{ \
|
||||||
|
std::fprintf(stderr, "%s failed: %s\n", #call, hipGetErrorName(hipResult_)); \
|
||||||
|
std::exit(1); \
|
||||||
|
} \
|
||||||
|
} while (0)
|
||||||
|
|
||||||
|
// Values k/256 only: exact as halves.
|
||||||
|
static uint16_t half_of(unsigned k)
|
||||||
|
{
|
||||||
|
if (k == 0)
|
||||||
|
return 0;
|
||||||
|
int exponent = 0;
|
||||||
|
while ((k >> exponent) > 1)
|
||||||
|
++exponent; // k = 1.m * 2^exponent
|
||||||
|
const unsigned mantissa = (k << (10 - exponent)) & 0x3ffu;
|
||||||
|
return static_cast<uint16_t>((exponent - 8 + 15) << 10 | mantissa);
|
||||||
|
}
|
||||||
|
|
||||||
|
// Pattern 0 ("A") and 1 ("B") differ in every texel.
|
||||||
|
static unsigned pattern_value(int pattern, int x, int y, int channel)
|
||||||
|
{
|
||||||
|
const unsigned base = (x * 7 + y * 13 + channel * 29) % 200;
|
||||||
|
return pattern == 0 ? base + 1 : 256 - base;
|
||||||
|
}
|
||||||
|
|
||||||
|
static void fill_pattern(uint16_t* texels, int width, int height, size_t rowPitch, int pattern)
|
||||||
|
{
|
||||||
|
for (int y = 0; y < height; ++y)
|
||||||
|
{
|
||||||
|
uint16_t* row = reinterpret_cast<uint16_t*>(reinterpret_cast<uint8_t*>(texels) + y * rowPitch);
|
||||||
|
for (int x = 0; x < width; ++x)
|
||||||
|
for (int channel = 0; channel < 4; ++channel)
|
||||||
|
row[x * 4 + channel] = half_of(pattern_value(pattern, x, y, channel));
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
// Which pattern the floats hold: 'A', 'B', or '?' with the number of texels matching each.
|
||||||
|
static char classify(const float* values, int width, int height, size_t* matchA, size_t* matchB)
|
||||||
|
{
|
||||||
|
*matchA = *matchB = 0;
|
||||||
|
for (int y = 0; y < height; ++y)
|
||||||
|
for (int x = 0; x < width; ++x)
|
||||||
|
{
|
||||||
|
bool a = true, b = true;
|
||||||
|
for (int channel = 0; channel < 4; ++channel)
|
||||||
|
{
|
||||||
|
const float value = values[(static_cast<size_t>(y) * width + x) * 4 + channel];
|
||||||
|
a &= value == pattern_value(0, x, y, channel) / 256.0f;
|
||||||
|
b &= value == pattern_value(1, x, y, channel) / 256.0f;
|
||||||
|
}
|
||||||
|
*matchA += a, *matchB += b;
|
||||||
|
}
|
||||||
|
const size_t texels = static_cast<size_t>(width) * height;
|
||||||
|
return *matchA == texels ? 'A' : *matchB == texels ? 'B' : '?';
|
||||||
|
}
|
||||||
|
|
||||||
|
static float half_to_float(uint16_t half)
|
||||||
|
{
|
||||||
|
const int exponent = (half >> 10) & 31, mantissa = half & 0x3ff;
|
||||||
|
if (exponent == 0)
|
||||||
|
return mantissa / 1024.0f / 16384.0f;
|
||||||
|
float value = 1.0f + mantissa / 1024.0f;
|
||||||
|
for (int shift = exponent - 15; shift != 0; shift += shift < 0 ? 1 : -1)
|
||||||
|
value = shift < 0 ? value / 2 : value * 2;
|
||||||
|
return value;
|
||||||
|
}
|
||||||
|
|
||||||
|
struct Vulkan
|
||||||
|
{
|
||||||
|
VkInstance instance = VK_NULL_HANDLE;
|
||||||
|
VkPhysicalDevice physical = VK_NULL_HANDLE;
|
||||||
|
VkDevice device = VK_NULL_HANDLE;
|
||||||
|
VkQueue queue = VK_NULL_HANDLE;
|
||||||
|
VkCommandPool pool = VK_NULL_HANDLE;
|
||||||
|
VkCommandBuffer cmd = VK_NULL_HANDLE;
|
||||||
|
VkPhysicalDeviceMemoryProperties memory = {};
|
||||||
|
};
|
||||||
|
|
||||||
|
static uint32_t memory_type(const Vulkan& vk, uint32_t typeBits, VkMemoryPropertyFlags wanted)
|
||||||
|
{
|
||||||
|
for (uint32_t index = 0; index < vk.memory.memoryTypeCount; ++index)
|
||||||
|
if ((typeBits & (1u << index)) && (vk.memory.memoryTypes[index].propertyFlags & wanted) == wanted)
|
||||||
|
return index;
|
||||||
|
std::fprintf(stderr, "no memory type with flags 0x%x\n", wanted);
|
||||||
|
std::exit(1);
|
||||||
|
}
|
||||||
|
|
||||||
|
static void submit(const Vulkan& vk)
|
||||||
|
{
|
||||||
|
VK_CHECK(vkEndCommandBuffer(vk.cmd));
|
||||||
|
VkSubmitInfo info = {VK_STRUCTURE_TYPE_SUBMIT_INFO};
|
||||||
|
info.commandBufferCount = 1;
|
||||||
|
info.pCommandBuffers = &vk.cmd;
|
||||||
|
VK_CHECK(vkQueueSubmit(vk.queue, 1, &info, VK_NULL_HANDLE));
|
||||||
|
VK_CHECK(vkQueueWaitIdle(vk.queue));
|
||||||
|
}
|
||||||
|
|
||||||
|
static void begin(const Vulkan& vk)
|
||||||
|
{
|
||||||
|
VkCommandBufferBeginInfo info = {VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO};
|
||||||
|
info.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT;
|
||||||
|
VK_CHECK(vkResetCommandBuffer(vk.cmd, 0));
|
||||||
|
VK_CHECK(vkBeginCommandBuffer(vk.cmd, &info));
|
||||||
|
}
|
||||||
|
|
||||||
|
int main(int argc, char** argv)
|
||||||
|
{
|
||||||
|
bool useBuffer = false, hipFirst = false, gameUsage = false, zeroInit = false;
|
||||||
|
int prewrites = 1, forcedType = -1;
|
||||||
|
int width = 1280, height = 720, positional = 0;
|
||||||
|
for (int index = 1; index < argc; ++index)
|
||||||
|
{
|
||||||
|
if (!std::strcmp(argv[index], "--buffer"))
|
||||||
|
useBuffer = true;
|
||||||
|
else if (!std::strcmp(argv[index], "--hip-first"))
|
||||||
|
hipFirst = true;
|
||||||
|
else if (!std::strcmp(argv[index], "--zero-init"))
|
||||||
|
zeroInit = true; // allocate with VK_MEMORY_ALLOCATE_ZERO_INITIALIZE_BIT_EXT, as vkd3d-proton does
|
||||||
|
else if (!std::strcmp(argv[index], "--game-usage"))
|
||||||
|
gameUsage = true;
|
||||||
|
else if (!std::strcmp(argv[index], "--memory-type") && index + 1 < argc)
|
||||||
|
forcedType = std::atoi(argv[++index]); // Vulkan memory type index for the shared object
|
||||||
|
else if (!std::strcmp(argv[index], "--prewrites") && index + 1 < argc)
|
||||||
|
prewrites = std::atoi(argv[++index]);
|
||||||
|
else
|
||||||
|
(positional++ == 0 ? width : height) = std::atoi(argv[index]);
|
||||||
|
}
|
||||||
|
if (hipFirst)
|
||||||
|
HIP_CHECK(hipInit(0));
|
||||||
|
|
||||||
|
// --- Vulkan device ---
|
||||||
|
Vulkan vk;
|
||||||
|
VkApplicationInfo app = {VK_STRUCTURE_TYPE_APPLICATION_INFO};
|
||||||
|
app.apiVersion = VK_API_VERSION_1_2;
|
||||||
|
VkInstanceCreateInfo instanceInfo = {VK_STRUCTURE_TYPE_INSTANCE_CREATE_INFO};
|
||||||
|
instanceInfo.pApplicationInfo = &app;
|
||||||
|
VK_CHECK(vkCreateInstance(&instanceInfo, nullptr, &vk.instance));
|
||||||
|
uint32_t count = 0;
|
||||||
|
VK_CHECK(vkEnumeratePhysicalDevices(vk.instance, &count, nullptr));
|
||||||
|
std::vector<VkPhysicalDevice> devices(count);
|
||||||
|
VK_CHECK(vkEnumeratePhysicalDevices(vk.instance, &count, devices.data()));
|
||||||
|
VkPhysicalDeviceProperties properties = {};
|
||||||
|
for (VkPhysicalDevice candidate : devices)
|
||||||
|
{
|
||||||
|
vkGetPhysicalDeviceProperties(candidate, &properties);
|
||||||
|
if (properties.vendorID == 0x1002 && properties.deviceType == VK_PHYSICAL_DEVICE_TYPE_DISCRETE_GPU)
|
||||||
|
{
|
||||||
|
vk.physical = candidate;
|
||||||
|
break;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
if (vk.physical == VK_NULL_HANDLE)
|
||||||
|
{
|
||||||
|
std::fprintf(stderr, "no discrete AMD Vulkan device\n");
|
||||||
|
return 1;
|
||||||
|
}
|
||||||
|
std::printf("Vulkan device: %s (driver version 0x%x)\n", properties.deviceName, properties.driverVersion);
|
||||||
|
vkGetPhysicalDeviceMemoryProperties(vk.physical, &vk.memory);
|
||||||
|
uint32_t families = 0;
|
||||||
|
vkGetPhysicalDeviceQueueFamilyProperties(vk.physical, &families, nullptr);
|
||||||
|
std::vector<VkQueueFamilyProperties> familyProperties(families);
|
||||||
|
vkGetPhysicalDeviceQueueFamilyProperties(vk.physical, &families, familyProperties.data());
|
||||||
|
uint32_t family = 0;
|
||||||
|
while (family < families && !(familyProperties[family].queueFlags & VK_QUEUE_GRAPHICS_BIT))
|
||||||
|
++family;
|
||||||
|
const float priority = 1.0f;
|
||||||
|
VkDeviceQueueCreateInfo queueInfo = {VK_STRUCTURE_TYPE_DEVICE_QUEUE_CREATE_INFO};
|
||||||
|
queueInfo.queueFamilyIndex = family;
|
||||||
|
queueInfo.queueCount = 1;
|
||||||
|
queueInfo.pQueuePriorities = &priority;
|
||||||
|
const char* extensions[] = {VK_KHR_EXTERNAL_MEMORY_FD_EXTENSION_NAME, "VK_EXT_zero_initialize_device_memory"};
|
||||||
|
VkPhysicalDeviceZeroInitializeDeviceMemoryFeaturesEXT zeroFeature = {
|
||||||
|
VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_ZERO_INITIALIZE_DEVICE_MEMORY_FEATURES_EXT};
|
||||||
|
zeroFeature.zeroInitializeDeviceMemory = VK_TRUE;
|
||||||
|
VkDeviceCreateInfo deviceInfo = {VK_STRUCTURE_TYPE_DEVICE_CREATE_INFO, zeroInit ? &zeroFeature : nullptr};
|
||||||
|
deviceInfo.queueCreateInfoCount = 1;
|
||||||
|
deviceInfo.pQueueCreateInfos = &queueInfo;
|
||||||
|
deviceInfo.enabledExtensionCount = zeroInit ? 2 : 1;
|
||||||
|
deviceInfo.ppEnabledExtensionNames = extensions;
|
||||||
|
VK_CHECK(vkCreateDevice(vk.physical, &deviceInfo, nullptr, &vk.device));
|
||||||
|
vkGetDeviceQueue(vk.device, family, 0, &vk.queue);
|
||||||
|
VkCommandPoolCreateInfo poolInfo = {VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO};
|
||||||
|
poolInfo.flags = VK_COMMAND_POOL_CREATE_RESET_COMMAND_BUFFER_BIT;
|
||||||
|
poolInfo.queueFamilyIndex = family;
|
||||||
|
VK_CHECK(vkCreateCommandPool(vk.device, &poolInfo, nullptr, &vk.pool));
|
||||||
|
VkCommandBufferAllocateInfo cmdInfo = {VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO};
|
||||||
|
cmdInfo.commandPool = vk.pool;
|
||||||
|
cmdInfo.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY;
|
||||||
|
cmdInfo.commandBufferCount = 1;
|
||||||
|
VK_CHECK(vkAllocateCommandBuffers(vk.device, &cmdInfo, &vk.cmd));
|
||||||
|
|
||||||
|
// --- the shared object: a linear image or a buffer on exportable device-local memory ---
|
||||||
|
VkExternalMemoryImageCreateInfo externalImage = {VK_STRUCTURE_TYPE_EXTERNAL_MEMORY_IMAGE_CREATE_INFO};
|
||||||
|
externalImage.handleTypes = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT;
|
||||||
|
VkExternalMemoryBufferCreateInfo externalBuffer = {VK_STRUCTURE_TYPE_EXTERNAL_MEMORY_BUFFER_CREATE_INFO};
|
||||||
|
externalBuffer.handleTypes = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT;
|
||||||
|
VkImage image = VK_NULL_HANDLE;
|
||||||
|
VkBuffer shared = VK_NULL_HANDLE;
|
||||||
|
VkMemoryRequirements requirements = {};
|
||||||
|
size_t rowPitch = static_cast<size_t>(width) * 8, texelOffset = 0;
|
||||||
|
VkMemoryDedicatedAllocateInfo dedicated = {VK_STRUCTURE_TYPE_MEMORY_DEDICATED_ALLOCATE_INFO};
|
||||||
|
if (useBuffer)
|
||||||
|
{
|
||||||
|
rowPitch = (rowPitch + 255) & ~static_cast<size_t>(255);
|
||||||
|
VkBufferCreateInfo info = {VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO, &externalBuffer};
|
||||||
|
info.size = rowPitch * height;
|
||||||
|
info.usage = VK_BUFFER_USAGE_TRANSFER_SRC_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT;
|
||||||
|
VK_CHECK(vkCreateBuffer(vk.device, &info, nullptr, &shared));
|
||||||
|
vkGetBufferMemoryRequirements(vk.device, shared, &requirements);
|
||||||
|
dedicated.buffer = shared;
|
||||||
|
}
|
||||||
|
else
|
||||||
|
{
|
||||||
|
VkImageCreateInfo info = {VK_STRUCTURE_TYPE_IMAGE_CREATE_INFO, &externalImage};
|
||||||
|
info.imageType = VK_IMAGE_TYPE_2D;
|
||||||
|
info.format = VK_FORMAT_R16G16B16A16_SFLOAT;
|
||||||
|
info.extent = {static_cast<uint32_t>(width), static_cast<uint32_t>(height), 1};
|
||||||
|
info.mipLevels = info.arrayLayers = 1;
|
||||||
|
info.samples = VK_SAMPLE_COUNT_1_BIT;
|
||||||
|
info.tiling = VK_IMAGE_TILING_LINEAR;
|
||||||
|
info.usage = VK_IMAGE_USAGE_TRANSFER_SRC_BIT | VK_IMAGE_USAGE_TRANSFER_DST_BIT | VK_IMAGE_USAGE_SAMPLED_BIT;
|
||||||
|
if (gameUsage)
|
||||||
|
{
|
||||||
|
info.usage |= VK_IMAGE_USAGE_COLOR_ATTACHMENT_BIT | VK_IMAGE_USAGE_STORAGE_BIT;
|
||||||
|
info.flags |= VK_IMAGE_CREATE_MUTABLE_FORMAT_BIT;
|
||||||
|
}
|
||||||
|
info.initialLayout = VK_IMAGE_LAYOUT_UNDEFINED;
|
||||||
|
VK_CHECK(vkCreateImage(vk.device, &info, nullptr, &image));
|
||||||
|
vkGetImageMemoryRequirements(vk.device, image, &requirements);
|
||||||
|
dedicated.image = image;
|
||||||
|
}
|
||||||
|
VkExportMemoryAllocateInfo exportInfo = {VK_STRUCTURE_TYPE_EXPORT_MEMORY_ALLOCATE_INFO, &dedicated};
|
||||||
|
exportInfo.handleTypes = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT;
|
||||||
|
VkMemoryAllocateFlagsInfo allocateFlags = {VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_FLAGS_INFO, &exportInfo};
|
||||||
|
allocateFlags.flags = VK_MEMORY_ALLOCATE_ZERO_INITIALIZE_BIT_EXT;
|
||||||
|
VkMemoryAllocateInfo allocation = {VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO,
|
||||||
|
zeroInit ? static_cast<const void*>(&allocateFlags) : &exportInfo};
|
||||||
|
allocation.allocationSize = requirements.size;
|
||||||
|
allocation.memoryTypeIndex = memory_type(vk, requirements.memoryTypeBits, VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT);
|
||||||
|
if (forcedType >= 0)
|
||||||
|
allocation.memoryTypeIndex = static_cast<uint32_t>(forcedType);
|
||||||
|
VkDeviceMemory sharedMemory = VK_NULL_HANDLE;
|
||||||
|
VK_CHECK(vkAllocateMemory(vk.device, &allocation, nullptr, &sharedMemory));
|
||||||
|
if (useBuffer)
|
||||||
|
VK_CHECK(vkBindBufferMemory(vk.device, shared, sharedMemory, 0));
|
||||||
|
else
|
||||||
|
{
|
||||||
|
VK_CHECK(vkBindImageMemory(vk.device, image, sharedMemory, 0));
|
||||||
|
const VkImageSubresource subresource = {VK_IMAGE_ASPECT_COLOR_BIT, 0, 0};
|
||||||
|
VkSubresourceLayout layout = {};
|
||||||
|
vkGetImageSubresourceLayout(vk.device, image, &subresource, &layout);
|
||||||
|
rowPitch = layout.rowPitch;
|
||||||
|
texelOffset = layout.offset;
|
||||||
|
}
|
||||||
|
std::printf("shared %s: %dx%d RGBA16F, row pitch %zu, offset %zu, allocation %llu bytes, memory type %u\n",
|
||||||
|
useBuffer ? "VkBuffer" : "linear VkImage", width, height, rowPitch, texelOffset,
|
||||||
|
static_cast<unsigned long long>(requirements.size), allocation.memoryTypeIndex);
|
||||||
|
|
||||||
|
// --- host-visible staging buffer for uploads and Vulkan-side readback ---
|
||||||
|
const size_t stagingBytes = rowPitch * height;
|
||||||
|
VkBufferCreateInfo stagingInfo = {VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO};
|
||||||
|
stagingInfo.size = stagingBytes;
|
||||||
|
stagingInfo.usage = VK_BUFFER_USAGE_TRANSFER_SRC_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT;
|
||||||
|
VkBuffer staging = VK_NULL_HANDLE;
|
||||||
|
VK_CHECK(vkCreateBuffer(vk.device, &stagingInfo, nullptr, &staging));
|
||||||
|
VkMemoryRequirements stagingRequirements = {};
|
||||||
|
vkGetBufferMemoryRequirements(vk.device, staging, &stagingRequirements);
|
||||||
|
VkMemoryAllocateInfo stagingAllocation = {VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO};
|
||||||
|
stagingAllocation.allocationSize = stagingRequirements.size;
|
||||||
|
stagingAllocation.memoryTypeIndex = memory_type(
|
||||||
|
vk, stagingRequirements.memoryTypeBits, VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | VK_MEMORY_PROPERTY_HOST_COHERENT_BIT);
|
||||||
|
VkDeviceMemory stagingMemory = VK_NULL_HANDLE;
|
||||||
|
VK_CHECK(vkAllocateMemory(vk.device, &stagingAllocation, nullptr, &stagingMemory));
|
||||||
|
VK_CHECK(vkBindBufferMemory(vk.device, staging, stagingMemory, 0));
|
||||||
|
void* stagingMap = nullptr;
|
||||||
|
VK_CHECK(vkMapMemory(vk.device, stagingMemory, 0, VK_WHOLE_SIZE, 0, &stagingMap));
|
||||||
|
|
||||||
|
VkBufferImageCopy region = {};
|
||||||
|
region.bufferRowLength = static_cast<uint32_t>(rowPitch / 8);
|
||||||
|
region.imageSubresource = {VK_IMAGE_ASPECT_COLOR_BIT, 0, 0, 1};
|
||||||
|
region.imageExtent = {static_cast<uint32_t>(width), static_cast<uint32_t>(height), 1};
|
||||||
|
const VkBufferCopy bufferRegion = {0, 0, stagingBytes};
|
||||||
|
bool imageInitialised = false;
|
||||||
|
const auto vulkan_upload = [&](int pattern) {
|
||||||
|
std::memset(stagingMap, 0, stagingBytes);
|
||||||
|
fill_pattern(static_cast<uint16_t*>(stagingMap), width, height, rowPitch, pattern);
|
||||||
|
begin(vk);
|
||||||
|
if (useBuffer)
|
||||||
|
vkCmdCopyBuffer(vk.cmd, staging, shared, 1, &bufferRegion);
|
||||||
|
else
|
||||||
|
{
|
||||||
|
if (!imageInitialised)
|
||||||
|
{
|
||||||
|
VkImageMemoryBarrier barrier = {VK_STRUCTURE_TYPE_IMAGE_MEMORY_BARRIER};
|
||||||
|
barrier.dstAccessMask = VK_ACCESS_TRANSFER_WRITE_BIT | VK_ACCESS_TRANSFER_READ_BIT;
|
||||||
|
barrier.oldLayout = VK_IMAGE_LAYOUT_UNDEFINED;
|
||||||
|
barrier.newLayout = VK_IMAGE_LAYOUT_GENERAL;
|
||||||
|
barrier.srcQueueFamilyIndex = barrier.dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED;
|
||||||
|
barrier.image = image;
|
||||||
|
barrier.subresourceRange = {VK_IMAGE_ASPECT_COLOR_BIT, 0, 1, 0, 1};
|
||||||
|
vkCmdPipelineBarrier(vk.cmd, VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT, VK_PIPELINE_STAGE_TRANSFER_BIT, 0, 0,
|
||||||
|
nullptr, 0, nullptr, 1, &barrier);
|
||||||
|
imageInitialised = true;
|
||||||
|
}
|
||||||
|
vkCmdCopyBufferToImage(vk.cmd, staging, image, VK_IMAGE_LAYOUT_GENERAL, 1, ®ion);
|
||||||
|
}
|
||||||
|
submit(vk);
|
||||||
|
};
|
||||||
|
std::vector<float> floats(static_cast<size_t>(width) * height * 4);
|
||||||
|
size_t matchA = 0, matchB = 0;
|
||||||
|
const auto report = [&](const char* what, char expected, const char* source) {
|
||||||
|
const char seen = classify(floats.data(), width, height, &matchA, &matchB);
|
||||||
|
std::printf(" %-34s %-22s pattern %c (texels matching A: %zu, B: %zu) %s\n", what, source, seen, matchA, matchB,
|
||||||
|
seen == expected ? "ok" : "MISMATCH");
|
||||||
|
return seen == expected;
|
||||||
|
};
|
||||||
|
const auto halves_to_floats = [&](const uint16_t* texels) {
|
||||||
|
for (int y = 0; y < height; ++y)
|
||||||
|
{
|
||||||
|
const uint16_t* row = reinterpret_cast<const uint16_t*>(reinterpret_cast<const uint8_t*>(texels) + y * rowPitch);
|
||||||
|
for (int x = 0; x < width * 4; ++x)
|
||||||
|
floats[static_cast<size_t>(y) * width * 4 + x] = half_to_float(row[x]);
|
||||||
|
}
|
||||||
|
};
|
||||||
|
const auto vulkan_read = [&](const char* what, char expected) {
|
||||||
|
std::memset(stagingMap, 0, stagingBytes);
|
||||||
|
begin(vk);
|
||||||
|
if (useBuffer)
|
||||||
|
vkCmdCopyBuffer(vk.cmd, shared, staging, 1, &bufferRegion);
|
||||||
|
else
|
||||||
|
vkCmdCopyImageToBuffer(vk.cmd, image, VK_IMAGE_LAYOUT_GENERAL, staging, 1, ®ion);
|
||||||
|
submit(vk);
|
||||||
|
halves_to_floats(static_cast<const uint16_t*>(stagingMap));
|
||||||
|
return report(what, expected, "Vulkan copy");
|
||||||
|
};
|
||||||
|
|
||||||
|
for (int write = prewrites - 1; write >= 0; --write)
|
||||||
|
vulkan_upload(write & 1); // alternating, ending on A
|
||||||
|
std::printf("Vulkan uploaded %d time(s) before the import, ending on pattern A\n", prewrites);
|
||||||
|
|
||||||
|
// --- HIP: import, texture object, sampling kernel ---
|
||||||
|
if (!hipFirst)
|
||||||
|
HIP_CHECK(hipInit(0));
|
||||||
|
int fd = -1;
|
||||||
|
VkMemoryGetFdInfoKHR fdInfo = {VK_STRUCTURE_TYPE_MEMORY_GET_FD_INFO_KHR};
|
||||||
|
fdInfo.memory = sharedMemory;
|
||||||
|
fdInfo.handleType = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT;
|
||||||
|
const auto getMemoryFd = reinterpret_cast<PFN_vkGetMemoryFdKHR>(vkGetDeviceProcAddr(vk.device, "vkGetMemoryFdKHR"));
|
||||||
|
VK_CHECK(getMemoryFd(vk.device, &fdInfo, &fd));
|
||||||
|
hipExternalMemoryHandleDesc handle = {};
|
||||||
|
handle.type = hipExternalMemoryHandleTypeOpaqueFd;
|
||||||
|
handle.handle.fd = fd;
|
||||||
|
handle.size = requirements.size;
|
||||||
|
handle.flags = hipExternalMemoryDedicated;
|
||||||
|
hipExternalMemory_t external = nullptr;
|
||||||
|
HIP_CHECK(hipImportExternalMemory(&external, &handle));
|
||||||
|
close(fd);
|
||||||
|
hipExternalMemoryBufferDesc mapping = {};
|
||||||
|
mapping.size = requirements.size;
|
||||||
|
void* base = nullptr;
|
||||||
|
HIP_CHECK(hipExternalMemoryGetMappedBuffer(&base, external, &mapping));
|
||||||
|
uint8_t* texels = static_cast<uint8_t*>(base) + texelOffset;
|
||||||
|
std::printf("HIP imported the memory at %p\n", base);
|
||||||
|
|
||||||
|
hipResourceDesc resource = {};
|
||||||
|
resource.resType = hipResourceTypePitch2D;
|
||||||
|
resource.res.pitch2D.devPtr = texels;
|
||||||
|
resource.res.pitch2D.desc = hipCreateChannelDesc(16, 16, 16, 16, hipChannelFormatKindFloat);
|
||||||
|
resource.res.pitch2D.width = width;
|
||||||
|
resource.res.pitch2D.height = height;
|
||||||
|
resource.res.pitch2D.pitchInBytes = rowPitch;
|
||||||
|
hipTextureDesc sampler = {};
|
||||||
|
sampler.addressMode[0] = sampler.addressMode[1] = hipAddressModeClamp;
|
||||||
|
sampler.filterMode = hipFilterModePoint;
|
||||||
|
sampler.readMode = hipReadModeElementType;
|
||||||
|
sampler.normalizedCoords = 0;
|
||||||
|
hipTextureObject_t texture = nullptr;
|
||||||
|
HIP_CHECK(hipCreateTextureObject(&texture, &resource, &sampler, nullptr));
|
||||||
|
float* deviceOut = nullptr;
|
||||||
|
HIP_CHECK(hipMalloc(&deviceOut, floats.size() * sizeof(float)));
|
||||||
|
const auto hip_sample = [&](const char* what, char expected) {
|
||||||
|
HIP_CHECK(hipMemset(deviceOut, 0, floats.size() * sizeof(float)));
|
||||||
|
const dim3 block(16, 16), grid((width + 15) / 16, (height + 15) / 16);
|
||||||
|
hipLaunchKernelGGL(sample_texture, grid, block, 0, 0, texture, deviceOut, width, height);
|
||||||
|
HIP_CHECK(hipDeviceSynchronize());
|
||||||
|
HIP_CHECK(hipMemcpy(floats.data(), deviceOut, floats.size() * sizeof(float), hipMemcpyDeviceToHost));
|
||||||
|
return report(what, expected, "HIP texture sample");
|
||||||
|
};
|
||||||
|
std::vector<uint16_t> hostTexels(stagingBytes / 2);
|
||||||
|
const auto hip_memcpy = [&](const char* what, char expected) {
|
||||||
|
HIP_CHECK(hipMemcpy(hostTexels.data(), texels, stagingBytes, hipMemcpyDeviceToHost));
|
||||||
|
halves_to_floats(hostTexels.data());
|
||||||
|
return report(what, expected, "hipMemcpy to host");
|
||||||
|
};
|
||||||
|
|
||||||
|
bool ok = true;
|
||||||
|
std::printf("after import (Vulkan last wrote A):\n");
|
||||||
|
ok &= hip_sample("texture created after A", 'A');
|
||||||
|
ok &= hip_memcpy("", 'A');
|
||||||
|
|
||||||
|
vulkan_upload(1);
|
||||||
|
std::printf("Vulkan uploaded pattern B:\n");
|
||||||
|
ok &= vulkan_read("", 'B');
|
||||||
|
ok &= hip_memcpy("", 'B');
|
||||||
|
ok &= hip_sample("same texture object", 'B');
|
||||||
|
hipTextureObject_t second = nullptr;
|
||||||
|
HIP_CHECK(hipCreateTextureObject(&second, &resource, &sampler, nullptr));
|
||||||
|
std::swap(texture, second);
|
||||||
|
ok &= hip_sample("new texture object", 'B');
|
||||||
|
std::swap(texture, second);
|
||||||
|
|
||||||
|
// HIP rewrites the texels with their own bytes.
|
||||||
|
void* scratch = nullptr;
|
||||||
|
HIP_CHECK(hipMalloc(&scratch, stagingBytes));
|
||||||
|
HIP_CHECK(hipMemcpy(scratch, texels, stagingBytes, hipMemcpyDeviceToDevice));
|
||||||
|
HIP_CHECK(hipMemcpy(texels, scratch, stagingBytes, hipMemcpyDeviceToDevice));
|
||||||
|
HIP_CHECK(hipDeviceSynchronize());
|
||||||
|
std::printf("HIP copied the texels out and back:\n");
|
||||||
|
ok &= hip_sample("same texture object", 'B');
|
||||||
|
ok &= vulkan_read("", 'B');
|
||||||
|
|
||||||
|
vulkan_upload(0);
|
||||||
|
std::printf("Vulkan uploaded pattern A again:\n");
|
||||||
|
ok &= hip_memcpy("", 'A');
|
||||||
|
ok &= hip_sample("same texture object", 'A');
|
||||||
|
|
||||||
|
std::printf("%s\n", ok ? "RESULT: consistent - HIP texture sampling follows Vulkan's writes"
|
||||||
|
: "RESULT: INCONSISTENT - see MISMATCH lines");
|
||||||
|
return ok ? 0 : 3;
|
||||||
|
}
|
||||||
@@ -1399,8 +1399,9 @@ static CUresult write_redirect_tail(CUsurfObject object, CUdeviceptr pointer, ui
|
|||||||
static unsigned char ring[256][16];
|
static unsigned char ring[256][16];
|
||||||
static unsigned int next;
|
static unsigned int next;
|
||||||
unsigned char* tail = ring[next++ % 256];
|
unsigned char* tail = ring[next++ % 256];
|
||||||
// dword 21: 0x5232 ('R2') << 16 | pitch / 8 (0 = off); dwords 22-23: pointer
|
// dword 21: 0x5232 ('R2', RGBA16F texels) << 16 | pitch / 8 (0 = off); dwords 22-23: pointer.
|
||||||
const uint32_t word = pitch != 0 ? 0x52320000u | ((pitch >> 3) & 0xffffu) : 0u;
|
// 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);
|
memset(tail, 0, 16);
|
||||||
memcpy(tail, &word, 4);
|
memcpy(tail, &word, 4);
|
||||||
memcpy(tail + 4, &pointer, 8);
|
memcpy(tail + 4, &pointer, 8);
|
||||||
|
|||||||
Reference in new issue
Block a user