From ff3e2695a7b68c335449a33641644f92ed0980fd Mon Sep 17 00:00:00 2001 From: James Newling Date: Wed, 15 Jul 2026 09:39:04 -0700 Subject: [PATCH] fix(rocjitsu): avoid D16 memory completion self-races --- .../rocjitsu/vm/amdgpu/memory_pipeline.cpp | 10 ++- .../tests/race-detector/CMakeLists.txt | 3 + .../race-detector/hip_race_gfx1151_test.hip | 83 +++++++++++++++++++ 3 files changed, 94 insertions(+), 2 deletions(-) diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/memory_pipeline.cpp b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/memory_pipeline.cpp index 1b883e73ed94..3db45a6c5879 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/memory_pipeline.cpp +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/memory_pipeline.cpp @@ -121,6 +121,12 @@ MemoryAccessCompletion vector_complete(VectorMemState &d, Wavefront &wf, Compute if (!d.is_load) return MemoryAccessCompletion::Complete; + // D16 memory completion merges a loaded half into the destination dword. The + // preserved half is storage state, not an instruction-visible source operand. + auto read_dst_lane_storage_for_d16_merge = [&](uint32_t reg, uint32_t lane) { + return reinterpret_cast(cu.raw_vgpr_data(reg))[lane]; + }; + // Buffer load with LDS bit: scatter loaded data into LDS instead of VGPRs. // Each lane writes num_elems * elem_size bytes to LDS at lds_base + lane_offset. if (d.lds_dst) @@ -156,7 +162,7 @@ MemoryAccessCompletion vector_complete(VectorMemState &d, Wavefront &wf, Compute for (uint32_t i = 0; i < vgpr_count; ++i) { uint32_t val = 0; if (!cu.sram_ecc() && d.elem_size <= 2 && (d.d16_hi || d.d16_lo)) { - const uint32_t old = cu.read_vgpr(d.dst_reg_base + i, lane); + const uint32_t old = read_dst_lane_storage_for_d16_merge(d.dst_reg_base + i, lane); val = d.d16_hi ? (old & 0x0000FFFFu) : (old & 0xFFFF0000u); } cu.write_vgpr(d.dst_reg_base + i, lane, val); @@ -185,7 +191,7 @@ MemoryAccessCompletion vector_complete(VectorMemState &d, Wavefront &wf, Compute else val = val & 0xFFFF; } else { - uint32_t old = cu.read_vgpr(d.dst_reg_base + i, lane); + uint32_t old = read_dst_lane_storage_for_d16_merge(d.dst_reg_base + i, lane); if (d.d16_hi) val = (old & 0xFFFF) | (val << 16); else diff --git a/emulation/rocjitsu/tests/race-detector/CMakeLists.txt b/emulation/rocjitsu/tests/race-detector/CMakeLists.txt index 6947a9a95199..408ba34cbff4 100644 --- a/emulation/rocjitsu/tests/race-detector/CMakeLists.txt +++ b/emulation/rocjitsu/tests/race-detector/CMakeLists.txt @@ -91,5 +91,8 @@ rj_add_gfx950_race_test_case(f64_vgpr_race) rj_add_gfx1151_race_test_case(vgpr_waitcnt) rj_add_gfx1151_race_test_case(vgpr_waitcnt_race) +rj_add_gfx1151_race_test_case(d16_global_load_merge_no_race) +rj_add_gfx1151_race_test_case(d16_global_load_hi_merge_no_race) +rj_add_gfx1151_race_test_case(d16_global_load_missing_waitcnt_race) rj_add_gfx1151_race_test_case(lds_wave32_cross_wave) rj_add_gfx1151_race_test_case(lds_wave32_cross_wave_race) diff --git a/emulation/rocjitsu/tests/race-detector/hip_race_gfx1151_test.hip b/emulation/rocjitsu/tests/race-detector/hip_race_gfx1151_test.hip index 6454a7dd91d9..b2b388dff1cb 100644 --- a/emulation/rocjitsu/tests/race-detector/hip_race_gfx1151_test.hip +++ b/emulation/rocjitsu/tests/race-detector/hip_race_gfx1151_test.hip @@ -8,6 +8,7 @@ #include #include +#include #include #include #include @@ -111,6 +112,50 @@ __global__ void vgpr_race_kernel(const float *src, float *dst) { dst[tid] = val; } +// D16 loads write only one half of the destination dword on RDNA. This kernel +// initializes the high half, performs a low-half global_load_d16_u8, waits for +// vmcnt, and then stores the merged dword. The test below checks two things: +// rocjitsu preserves the untouched high half, and the internal merge read used +// during memory completion is not reported as an instruction-visible VGPR read. +__global__ void d16_global_load_merge_kernel(const uint8_t *src, uint32_t *dst) { + int tid = threadIdx.x; + uint32_t val = 0xABCD0000u; + asm volatile("global_load_d16_u8 %0, %1, off\n" + "s_waitcnt vmcnt(0)\n" + : "+v"(val) + : "v"(&src[tid]) + : "memory"); + dst[tid] = val; +} + +// High-half D16 loads have the same internal merge requirement, but preserve the +// low half and write the loaded byte into bits [31:16]. +__global__ void d16_global_load_hi_merge_kernel(const uint8_t *src, uint32_t *dst) { + int tid = threadIdx.x; + uint32_t val = 0x0000ABCDu; + asm volatile("global_load_d16_hi_u8 %0, %1, off\n" + "s_waitcnt vmcnt(0)\n" + : "+v"(val) + : "v"(&src[tid]) + : "memory"); + dst[tid] = val; +} + +// Same D16 load shape as d16_global_load_merge_kernel, but this intentionally +// omits s_waitcnt before storing the VGPR. The store consumes the load +// destination while the low-half load is still outstanding, so the race detector +// should report a real VGPR race after the memory-completion self-race is fixed. +__global__ void d16_global_load_missing_waitcnt_kernel(const uint8_t *src, uint32_t *dst) { + int tid = threadIdx.x; + uint32_t val = 0xABCD0000u; + asm volatile("global_load_d16_u8 %0, %1, off\n" + "s_nop 0\n" + : "+v"(val) + : "v"(&src[tid]) + : "memory"); + dst[tid] = val; +} + __attribute__((amdgpu_flat_work_group_size(64, 64))) __global__ void lds_wave32_cross_wave_safe_kernel(int *out) { __shared__ int lds[64]; @@ -147,6 +192,44 @@ TEST_F(Gfx1151RaceTest, vgpr_waitcnt_race) { EXPECT_EQ(records[0].type, "VGPR"); } +TEST_F(Gfx1151RaceTest, d16_global_load_merge_no_race) { + auto *src = allocWithData(32); + auto *dst = alloc(32); + d16_global_load_merge_kernel<<<1, 32>>>(src, dst); + sync(); + + std::vector h(32); + (void)hipMemcpy(h.data(), dst, h.size() * sizeof(uint32_t), hipMemcpyDeviceToHost); + for (uint32_t i = 0; i < h.size(); ++i) + EXPECT_EQ(h[i], 0xABCD0000u | i); + ExpectNoRace(); +} + +TEST_F(Gfx1151RaceTest, d16_global_load_hi_merge_no_race) { + auto *src = allocWithData(32); + auto *dst = alloc(32); + d16_global_load_hi_merge_kernel<<<1, 32>>>(src, dst); + sync(); + + std::vector h(32); + (void)hipMemcpy(h.data(), dst, h.size() * sizeof(uint32_t), hipMemcpyDeviceToHost); + for (uint32_t i = 0; i < h.size(); ++i) + EXPECT_EQ(h[i], (i << 16) | 0xABCDu); + ExpectNoRace(); +} + +TEST_F(Gfx1151RaceTest, d16_global_load_missing_waitcnt_race) { + auto *src = allocWithData(32); + auto *dst = alloc(32); + d16_global_load_missing_waitcnt_kernel<<<1, 32>>>(src, dst); + sync(); + + auto records = ExpectRace(); + ASSERT_FALSE(records.empty()); + EXPECT_EQ(records[0].type, "VGPR"); + EXPECT_NE(records[0].message.find("global_load_d16_u8"), std::string::npos); +} + TEST_F(Gfx1151RaceTest, lds_wave32_cross_wave) { lds_wave32_cross_wave_safe_kernel<<<1, 64>>>(alloc(64)); sync();