[rocjitsu] Fix DS transpose reads ignoring AccVGPR acc bit

emulation/rocjitsu/lib/python/amdisa/codegen/_generator.py

@@ -2043,11 +2043,12 @@ def _gen_ds_read_tr(
             'ds_read_tr_b16': (4, 2, 4),  # elem_size=4, num_elems=2, transpose=4
         }
         esz, ne, tr_kind = tr_map.get(sem.semantic_class, (4, 2, 4))
+        acc = self._acc_vgpr_expr
         L = []
         L.append(
             '  auto d = std::make_unique<amdgpu::VectorMemState>(amdgpu::LOCAL_MEM);'
         )
-        L.append(f'  d->dst_reg_base = wf.vgpr_alloc().base + inst_.vdst;')
+        L.append(f'  d->dst_reg_base = wf.vgpr_alloc().base + {acc} + inst_.vdst;')
         L.append(f'  d->elem_size = {esz};')
         L.append(f'  d->num_elems = {ne};')
         L.append('  d->is_load = true;')

emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/cdna4/ds.cpp

@@ -4280,7 +4280,7 @@ void DsReadB64TrB4Ds::execute_impl(amdgpu::Wavefront &wf) {
   if (inst_.gds)
     throw util::UnimplementedInst(mnemonic());
   auto d = std::make_unique<amdgpu::VectorMemState>(amdgpu::LOCAL_MEM);
-  d->dst_reg_base = wf.vgpr_alloc().base + inst_.vdst;
+  d->dst_reg_base = wf.vgpr_alloc().base + (inst_.acc ? 256u : 0u) + inst_.vdst;
   d->elem_size = 4;
   d->num_elems = 2;
   d->is_load = true;
@@ -4305,7 +4305,7 @@ void DsReadB96TrB6Ds::execute_impl(amdgpu::Wavefront &wf) {
   if (inst_.gds)
     throw util::UnimplementedInst(mnemonic());
   auto d = std::make_unique<amdgpu::VectorMemState>(amdgpu::LOCAL_MEM);
-  d->dst_reg_base = wf.vgpr_alloc().base + inst_.vdst;
+  d->dst_reg_base = wf.vgpr_alloc().base + (inst_.acc ? 256u : 0u) + inst_.vdst;
   d->elem_size = 4;
   d->num_elems = 3;
   d->is_load = true;
@@ -4330,7 +4330,7 @@ void DsReadB64TrB8Ds::execute_impl(amdgpu::Wavefront &wf) {
   if (inst_.gds)
     throw util::UnimplementedInst(mnemonic());
   auto d = std::make_unique<amdgpu::VectorMemState>(amdgpu::LOCAL_MEM);
-  d->dst_reg_base = wf.vgpr_alloc().base + inst_.vdst;
+  d->dst_reg_base = wf.vgpr_alloc().base + (inst_.acc ? 256u : 0u) + inst_.vdst;
   d->elem_size = 4;
   d->num_elems = 2;
   d->is_load = true;
@@ -4355,7 +4355,7 @@ void DsReadB64TrB16Ds::execute_impl(amdgpu::Wavefront &wf) {
   if (inst_.gds)
     throw util::UnimplementedInst(mnemonic());
   auto d = std::make_unique<amdgpu::VectorMemState>(amdgpu::LOCAL_MEM);
-  d->dst_reg_base = wf.vgpr_alloc().base + inst_.vdst;
+  d->dst_reg_base = wf.vgpr_alloc().base + (inst_.acc ? 256u : 0u) + inst_.vdst;
   d->elem_size = 4;
   d->num_elems = 2;
   d->is_load = true;

emulation/rocjitsu/tests/CMakeLists.txt

@@ -473,6 +473,7 @@ if(HIPCC_EXECUTABLE AND HSA_RUNTIME64)
     rj_add_hip_test_case(hip_memcpy_test "HipMemcpyTest.DeviceToDevice")
     rj_add_hip_test(hip_vector_add_test)
     rj_add_hip_test_case(hip_vector_add_test "HipVectorAddTest.CorrectResult")
+    rj_add_hip_test(hip_ds_transpose_acc_test)
 
     # --- RCCL collective test ---
     find_library(

emulation/rocjitsu/tests/amdgpu_vm_test.cpp

@@ -457,8 +457,9 @@ constexpr uint32_t v_cmp_eq_f32(uint32_t s0, uint32_t vs1) { return vopc(66, s0,
 // DS: 64-bit instruction.
 // dword0: offset0[7:0], offset1[15:8], gds[16], op[24:17], acc[25], encoding[31:26]=0x36
 // dword1: addr[7:0], data0[15:8], data1[23:16], vdst[31:24]
-constexpr uint32_t ds_lo(uint32_t op, uint8_t offset0 = 0, uint8_t offset1 = 0) {
-  return (0x36u << 26) | (op << 17) | (static_cast<uint32_t>(offset1) << 8) | offset0;
+constexpr uint32_t ds_lo(uint32_t op, uint8_t offset0 = 0, uint8_t offset1 = 0, uint8_t acc = 0) {
+  return (0x36u << 26) | (static_cast<uint32_t>(acc) << 25) | (op << 17) |
+         (static_cast<uint32_t>(offset1) << 8) | offset0;
 }
 constexpr uint32_t ds_hi(uint32_t vdst, uint32_t data0, uint32_t addr, uint32_t data1 = 0) {
   return (vdst << 24) | (data1 << 16) | (data0 << 8) | addr;
@@ -1129,4 +1130,57 @@ TEST(AtomicStressTest, GlobalAtomicAdd_MultiWorkgroup) {
   EXPECT_EQ(final_val, 1256u) << "1000 + 256 global atomic adds = 1256";
 }
 
+// Verify that ds_read_b64_tr_b16 with acc=1 writes to AccVGPR (vb+256+vdst),
+// not to VGPR (vb+vdst).
+TEST(DsTransposeTest, ReadB64TrB16_AccBit) {
+  VmFixture f("cdna4", 1, 10);
+
+  constexpr uint32_t VDST = 4;
+  constexpr uint32_t ADDR_REG = 0;
+  constexpr uint32_t DS_OP = 227; // ds_read_b64_tr_b16
+
+  // Kernel:
+  //   v_mov_b32 v0, 0          ; addr = LDS offset 0
+  //   v_mov_b32 v4, 0x42       ; sentinel in VGPR v4
+  //   v_mov_b32 v5, 0x42       ; sentinel in VGPR v5
+  //   ds_read_b64_tr_b16 a[4:5], v0  ; acc=1: write to AccVGPR
+  //   s_waitcnt lgkmcnt(0)
+  //   s_endpgm
+  using namespace enc;
+  const uint32_t code[] = {
+      v_mov_b32(ADDR_REG, INLINE_CONST(0)),
+      v_mov_b32(VDST, INLINE_CONST(42)),
+      v_mov_b32(VDST + 1, INLINE_CONST(42)),
+      ds_lo(DS_OP, /*offset0=*/0, /*offset1=*/0, /*acc=*/1),
+      ds_hi(VDST, /*data0=*/0, ADDR_REG),
+      S_WAITCNT_0,
+      S_ENDPGM,
+  };
+  uint64_t ko = f.write_kernel(0x1000, code, sizeof(code));
+
+  auto *cu = f.cu();
+
+  // Write a known non-zero pattern to LDS.
+  for (uint32_t i = 0; i < 256; ++i)
+    cu->lds().write32(i * 4, 0xDEADBEEF);
+
+  test::AqlQueue queue(f.mem(), f.cp());
+  queue.dispatch(ko, 64);
+  f.engine->run();
+
+  auto *wf = cu->wf(0);
+  ASSERT_NE(wf, nullptr);
+  uint32_t vb = wf->vgpr_alloc().base;
+
+  // VGPR v4 should still hold the sentinel (42), not overwritten by ds_read.
+  uint32_t vgpr_val = cu->read_vgpr(vb + VDST, 0);
+  EXPECT_EQ(vgpr_val, 42u) << "VGPR v" << VDST << " should NOT have been written when acc=1";
+
+  // AccVGPR a4 should have been written with LDS data (not 42, not 0).
+  uint32_t acc_val = cu->read_vgpr(vb + 256 + VDST, 0);
+  EXPECT_NE(acc_val, 0u) << "AccVGPR a" << VDST << " should have been written by ds_read";
+  EXPECT_NE(acc_val, 42u) << "AccVGPR a" << VDST
+                          << " should contain LDS data, not the VGPR sentinel";
+}
+
 } // namespace

emulation/rocjitsu/tests/hip_ds_transpose_acc_test.cpp

@@ -0,0 +1,88 @@
+// Copyright (c) 2026 Advanced Micro Devices, Inc.
+// SPDX-License-Identifier: MIT
+
+/// @file hip_ds_transpose_acc_test.cpp
+/// @brief Verify ds_read_b64_tr_b16 with acc=1 writes to AccVGPR, not VGPR.
+
+#include <cstdint>
+#include <hip/hip_runtime.h>
+#include <vector>
+
+#include <gtest/gtest.h>
+
+int main(int argc, char **argv) {
+  ::testing::InitGoogleTest(&argc, argv);
+  int rc = RUN_ALL_TESTS();
+  (void)hipDeviceReset();
+  return rc;
+}
+
+#define HIP_ASSERT(call)                                                                           \
+  do {                                                                                             \
+    hipError_t err = (call);                                                                       \
+    ASSERT_EQ(err, hipSuccess) << "HIP error: " << hipGetErrorString(err);                         \
+  } while (0)
+
+// out[0..63]   = v2 per lane (should be 0xC0DE if ds_read went to AccVGPR)
+// out[64..127] = a2 per lane (should differ from 0xF00D if ds_read wrote here)
+__global__ void ds_transpose_acc_kernel(uint32_t *out) {
+  __shared__ uint32_t lds[128];
+  int tid = threadIdx.x;
+
+  lds[tid] = 0xCAFE0000u | tid;
+  if (tid + 64 < 128)
+    lds[tid + 64] = 0xBEEF0000u | (tid + 64);
+  __syncthreads();
+
+  int lds_addr = static_cast<int>(reinterpret_cast<uintptr_t>(lds) & 0xFFFF);
+
+  uint32_t v2_val, a2_val;
+  uint32_t v2_init = 0xC0DEu;
+  uint32_t a2_init = 0xF00Du;
+
+  asm volatile("v_accvgpr_write_b32 a2, %2\n"
+               "v_accvgpr_write_b32 a3, %2\n"
+               "v_mov_b32 v2, %3\n"
+               "v_mov_b32 v3, %3\n"
+               "ds_read_b64_tr_b16 a[2:3], %4\n"
+               "s_waitcnt lgkmcnt(0)\n"
+               "v_mov_b32 %0, v2\n"
+               "v_accvgpr_read_b32 %1, a2\n"
+               : "=v"(v2_val), "=v"(a2_val)
+               : "v"(a2_init), "v"(v2_init), "v"(lds_addr)
+               : "v2", "v3", "a2", "a3");
+
+  out[tid] = v2_val;
+  out[64 + tid] = a2_val;
+}
+
+TEST(DsTransposeAccTest, AccBitWritesToAccVGPR) {
+  constexpr int N = 128;
+  uint32_t *d_out = nullptr;
+  HIP_ASSERT(hipMalloc(&d_out, N * sizeof(uint32_t)));
+  HIP_ASSERT(hipMemset(d_out, 0, N * sizeof(uint32_t)));
+
+  ds_transpose_acc_kernel<<<1, 64>>>(d_out);
+  HIP_ASSERT(hipDeviceSynchronize());
+
+  std::vector<uint32_t> h(N);
+  HIP_ASSERT(hipMemcpy(h.data(), d_out, N * sizeof(uint32_t), hipMemcpyDeviceToHost));
+  (void)hipFree(d_out);
+
+  // v2 should be 0xC0DE for all lanes — the ds_read with acc=1 should not
+  // have touched VGPR v2.
+  int v2_ok = 0;
+  for (int i = 0; i < 64; ++i)
+    if (h[i] == 0xC0DEu)
+      v2_ok++;
+  EXPECT_EQ(v2_ok, 64) << "VGPR v2 was clobbered by ds_read — acc bit not respected";
+
+  // a2 should have changed from the initial 0xF00D for at least some lanes
+  // (the transpose shuffle may produce 0xF00D for a few lanes by coincidence,
+  // but not for all 64).
+  int a2_unchanged = 0;
+  for (int i = 0; i < 64; ++i)
+    if (h[64 + i] == 0xF00Du)
+      a2_unchanged++;
+  EXPECT_LT(a2_unchanged, 64) << "AccVGPR a2 was never written by ds_read";
+}