fix(rocjitsu): avoid plugin false-firing on D16 load completion

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<const uint32_t *>(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

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)

emulation/rocjitsu/tests/race-detector/hip_race_gfx1151_test.hip

@@ -8,6 +8,7 @@
 #include <hip/hip_runtime.h>
 
 #include <cstdlib>
+#include <cstdint>
 #include <fstream>
 #include <string>
 #include <vector>
@@ -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<uint8_t>(32);
+  auto *dst = alloc<uint32_t>(32);
+  d16_global_load_merge_kernel<<<1, 32>>>(src, dst);
+  sync();
+
+  std::vector<uint32_t> 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<uint8_t>(32);
+  auto *dst = alloc<uint32_t>(32);
+  d16_global_load_hi_merge_kernel<<<1, 32>>>(src, dst);
+  sync();
+
+  std::vector<uint32_t> 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<uint8_t>(32);
+  auto *dst = alloc<uint32_t>(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<int>(64));
   sync();