From 283738d626f38716a5e9c49a30b43022b275a79e Mon Sep 17 00:00:00 2001 From: woutermans Date: Wed, 7 Oct 2026 10:59:56 +0200 Subject: [PATCH] Document the linear-texture input investigation - Add docs/in-place-inputs.md: reading the DLSS inputs in place does not work, because a hardware texture fetch on a HIP texture object over memory imported from vkd3d-proton returns stale data while plain loads and cuMemcpy2D are correct; records the measurements, the rule-outs and the reworked-fetch workaround that turned out not to pay for itself. - Add tools/hip_vk_texture_stale_repro.cpp and its build script, the native Vulkan + HIP reproducer that rules out a general ROCm or RADV bug. --- docs/in-place-inputs.md | 53 +++ scripts/build_hip_vk_texture_stale_repro.sh | 11 + tools/hip_vk_texture_stale_repro.cpp | 479 ++++++++++++++++++++ 3 files changed, 543 insertions(+) create mode 100644 docs/in-place-inputs.md create mode 100755 scripts/build_hip_vk_texture_stale_repro.sh create mode 100644 tools/hip_vk_texture_stale_repro.cpp diff --git a/docs/in-place-inputs.md b/docs/in-place-inputs.md new file mode 100644 index 0000000..382209c --- /dev/null +++ b/docs/in-place-inputs.md @@ -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. diff --git a/scripts/build_hip_vk_texture_stale_repro.sh b/scripts/build_hip_vk_texture_stale_repro.sh new file mode 100755 index 0000000..6938868 --- /dev/null +++ b/scripts/build_hip_vk_texture_stale_repro.sh @@ -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" diff --git a/tools/hip_vk_texture_stale_repro.cpp b/tools/hip_vk_texture_stale_repro.cpp new file mode 100644 index 0000000..676e14b --- /dev/null +++ b/tools/hip_vk_texture_stale_repro.cpp @@ -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 +#include + +#include +#include +#include +#include +#include + +__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(texture, x + 0.5f, y + 0.5f); + float* pixel = out + (static_cast(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((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(reinterpret_cast(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(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(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 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 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(width) * 8, texelOffset = 0; + VkMemoryDedicatedAllocateInfo dedicated = {VK_STRUCTURE_TYPE_MEMORY_DEDICATED_ALLOCATE_INFO}; + if (useBuffer) + { + rowPitch = (rowPitch + 255) & ~static_cast(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(width), static_cast(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(&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(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(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(rowPitch / 8); + region.imageSubresource = {VK_IMAGE_ASPECT_COLOR_BIT, 0, 0, 1}; + region.imageExtent = {static_cast(width), static_cast(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(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 floats(static_cast(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(reinterpret_cast(texels) + y * rowPitch); + for (int x = 0; x < width * 4; ++x) + floats[static_cast(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(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(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(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 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; +}