From 5a8eaea095c5904511a093f946748e250ef9dfb4 Mon Sep 17 00:00:00 2001 From: Alan Li Date: Mon, 6 Jul 2026 20:17:24 -0700 Subject: [PATCH] fix(rocjitsu): detect VGPR WAW races Add an instruction-write VGPR hook and keep simulator setup/memory writeback on raw VGPR paths so race-detector WAW checks only see architectural instruction writes. Refresh race-detector docs and add focused lifecycle, profiling, builder, and HIP coverage for the new behavior. --- emulation/rocjitsu/docs/race-detector.md | 16 +- .../lib/python/amdisa/codegen/_generator.py | 29 +- .../tests/test_generator_profile_gates.py | 38 +++ .../isa/arch/amdgpu/cdna1/encodings.h | 12 + .../isa/arch/amdgpu/cdna2/encodings.h | 12 + .../isa/arch/amdgpu/cdna3/encodings.h | 12 + .../isa/arch/amdgpu/cdna4/encodings.h | 12 + .../isa/arch/amdgpu/gfx1250/encodings.h | 27 ++ .../isa/arch/amdgpu/rdna1/encodings.h | 12 + .../isa/arch/amdgpu/rdna2/encodings.h | 12 + .../isa/arch/amdgpu/rdna3/encodings.h | 27 ++ .../isa/arch/amdgpu/rdna3_5/encodings.h | 27 ++ .../isa/arch/amdgpu/rdna4/encodings.h | 27 ++ .../isa/arch/amdgpu/shared/dpp_sdwa_ops.h | 39 +++ .../isa/arch/amdgpu/shared/mma_exec.h | 35 +-- .../isa/arch/amdgpu/shared/simd_glue.h | 48 ++- .../rocjitsu/src/rocjitsu/isa/instruction.h | 21 ++ .../src/rocjitsu/isa/isa_operand_simd_inl.h | 2 +- .../rocjitsu/vm/amdgpu/command_processor.cpp | 8 +- .../src/rocjitsu/vm/amdgpu/compute_unit.cpp | 123 +++++++- .../src/rocjitsu/vm/amdgpu/compute_unit.h | 89 +++++- .../rocjitsu/vm/amdgpu/memory_pipeline.cpp | 10 +- .../rocjitsu/vm/plugins/execution_plugin.h | 8 + .../vm/plugins/execution_plugin_group.h | 6 + .../plugins/profiled_execution_plugin_group.h | 24 +- .../race_detector/core/race_detector.h | 3 +- .../race_detector/core/wave_race_state.cpp | 11 + .../race_detector/core/wave_race_state.h | 3 + .../vm/plugins/race_detector/plugin.cpp | 16 +- .../vm/plugins/race_detector/plugin.h | 3 + .../rocjitsu/tests/execution_plugin_test.cpp | 289 +++++++++++++++++- .../tests/race-detector/CMakeLists.txt | 1 + .../race-detector/hip_race_gfx950_test.hip | 51 ++-- .../race-detector/race_detector_tests.cpp | 43 ++- .../tests/race-detector/race_test_builder.h | 4 + .../rocjitsu/tests/shared_infra_test.cpp | 11 + 36 files changed, 1008 insertions(+), 103 deletions(-) diff --git a/emulation/rocjitsu/docs/race-detector.md b/emulation/rocjitsu/docs/race-detector.md index 10b131336cf3..a25cd0e9be92 100644 --- a/emulation/rocjitsu/docs/race-detector.md +++ b/emulation/rocjitsu/docs/race-detector.md @@ -76,7 +76,7 @@ RJ_RACE=1 build/tools/rocjitsu/rocjitsu --config configs/gfx950_cdna4.json -- /t You should see output: ``` -RACE type=LDS reg=508 wave=0 lane=0 wg=0,0,0 conflict=unknown +RACE type=LDS access=read reg=508 wave=0 lane=0 wg=0,0,0 conflict=unknown Race on LDS byte 508 [workgroup (0, 0, 0), wave 0, lane 0] ==> ds_write_b32 v0, v1 ; <-- wave 1 v_sub_u32_e32 v1, 0, v0 @@ -127,7 +127,8 @@ races. Some examples: ## What this plugin detects - **VGPR races**: a vector register is read before a pending global or LDS load - has completed (`s_waitcnt vmcnt` / `s_waitcnt lgkmcnt` insufficient). + has completed (`s_waitcnt vmcnt` / `s_waitcnt lgkmcnt` insufficient), or an + instruction writes a VGPR that still has an active pending load. - **SGPR races**: a scalar register is read before a pending scalar load has completed (`s_waitcnt lgkmcnt` insufficient). - **LDS races**: an LDS byte is read or written by one wave while another wave @@ -155,7 +156,9 @@ operations are in flight. When an instruction in the emulator accesses an LDS byte, there is a check to see what memory events are still in flight that read/write that byte, from the perspective of the accessing thread. In this way, RAW (read-after-write) and WAR (write-after-read) hazards can be detected. -Similar logic applies for VGPR and SGPR accesses. +Similar logic applies for VGPR and SGPR accesses. VGPR writes are also checked +against outstanding loads to catch write-after-write hazards where a later ALU +write may be clobbered by a still-pending memory result. **LDS race detection** uses coarse-grained counters (one per 16-byte chunk) for fast-path checks, with interval-based overlap scanning as a fallback. Live @@ -266,10 +269,9 @@ ctest --test-dir build -R "RaceTest" (missing `s_waitcnt` and `s_barrier`). It does not detect inter-workgroup races, races between dispatches, or host-device synchronization issues. -- **No WAW detection**: write-after-write hazards are not currently flagged. - This includes both LDS WAW (two waves writing to the same LDS byte without a - barrier) and VGPR WAW (an ALU instruction overwriting a register that has a - pending global load). +- **No LDS WAW detection**: LDS write-after-write hazards are not currently + flagged. VGPR write-after-load hazards are checked, but two waves writing the + same LDS byte without a barrier are still not reported. - **Kernel name resolution**: kernel names in race reports may show as `"?"` if symbol information is not available in the code object. diff --git a/emulation/rocjitsu/lib/python/amdisa/codegen/_generator.py b/emulation/rocjitsu/lib/python/amdisa/codegen/_generator.py index fc8debb1af93..8380fb5ad709 100644 --- a/emulation/rocjitsu/lib/python/amdisa/codegen/_generator.py +++ b/emulation/rocjitsu/lib/python/amdisa/codegen/_generator.py @@ -2029,6 +2029,33 @@ def gen_encodings(self) -> None: public_members.append(func_decl) class_func_impls.append(func_body) + _enc_upper = inst_enc.enc_name.upper() + _dpp_struct, _dpp8_struct = self._vop_dpp_struct_names(_enc_upper) + if _enc_upper in ('ENC_VOP1', 'ENC_VOP2'): + public_members.append( + cgen.Line( + 'std::optional vgpr_write_hook_filter(' + 'uint32_t wf_size, uint64_t exec) const override {' + ' return amdgpu::vop_write_hook_filter(' + 'inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_,' + ' dpp_bank_mask_, dpp_bound_ctrl_, sdwa_dst_sel_,' + ' sdwa_dst_unused_); }' + ) + ) + elif _dpp_struct and _enc_upper in ( + 'ENC_VOP3', + 'ENC_VOP3P', + 'VOP3_SDST_ENC', + ): + public_members.append( + cgen.Line( + 'std::optional vgpr_write_hook_filter(' + 'uint32_t wf_size, uint64_t exec) const override {' + ' return amdgpu::dpp_write_hook_filter(' + 'inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_,' + ' dpp_bank_mask_, dpp_bound_ctrl_); }' + ) + ) class_members.extend(public_members) class_members.append( cgen.Statement(f'using OpEncoding = {inst_enc.fmt_enc_name}MachineInst') @@ -2041,8 +2068,6 @@ def gen_encodings(self) -> None: class_members.append(cgen.Statement('std::string owned_mnemonic_')) # VOP encoding bases store DPP control fields. # apply_dpp() is a free function in dpp_sdwa_ops.h. - _enc_upper = inst_enc.enc_name.upper() - _dpp_struct, _dpp8_struct = self._vop_dpp_struct_names(_enc_upper) if ( _dpp_struct or _dpp8_struct diff --git a/emulation/rocjitsu/lib/python/amdisa/tests/test_generator_profile_gates.py b/emulation/rocjitsu/lib/python/amdisa/tests/test_generator_profile_gates.py index 93a80a629ba6..c71ecaa591bb 100644 --- a/emulation/rocjitsu/lib/python/amdisa/tests/test_generator_profile_gates.py +++ b/emulation/rocjitsu/lib/python/amdisa/tests/test_generator_profile_gates.py @@ -1280,6 +1280,44 @@ def test_generated_dpp_cleanup_uses_full_write_mask_for_dpp16( assert 'dpp_bound_ctrl_, dpp_fi_' in body +def test_vop_encoding_bases_publish_architectural_write_hook_filter( + amdgpu_generated_root: Path, +): + arches = ( + 'cdna1', + 'cdna2', + 'cdna3', + 'cdna4', + 'rdna1', + 'rdna2', + 'rdna3', + 'rdna3_5', + 'rdna4', + 'gfx1250', + ) + for arch in arches: + encodings = (amdgpu_generated_root / arch / 'encodings.h').read_text() + for class_name, next_class in (('Vop1', 'Vopc'), ('Vop2', 'Vop3p')): + start = encodings.index(f'class {class_name} :') + end = encodings.index(f'class {next_class} :', start) + body = encodings[start:end] + assert 'vgpr_write_hook_filter' in body, arch + assert 'amdgpu::vop_write_hook_filter(' in body, arch + assert 'sdwa_dst_unused_' in body, arch + + for arch in ('rdna3', 'rdna3_5', 'rdna4', 'gfx1250'): + encodings = (amdgpu_generated_root / arch / 'encodings.h').read_text() + classes = ('Vop3', 'Vop3p', 'Vop3SdstEnc') + for class_name, next_class in zip(classes, classes[1:] + ('}',)): + start = encodings.index(f'class {class_name} :') + end = encodings.index( + f'class {next_class} :' if next_class != '}' else '\n};', start + ) + body = encodings[start:end] + assert 'vgpr_write_hook_filter' in body, (arch, class_name) + assert 'amdgpu::dpp_write_hook_filter(' in body, (arch, class_name) + + def test_rdna1_2_generated_vopc_dpp_is_explicitly_unsupported( amdgpu_generated_root: Path, ): diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/cdna1/encodings.h b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/cdna1/encodings.h index 94d11639205b..4d5d754f5a05 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/cdna1/encodings.h +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/cdna1/encodings.h @@ -476,6 +476,12 @@ class Vop1 : public IsaInstruction { public: Vop1(std::string_view mnemonic, const Vop1MachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::vop_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_, sdwa_dst_sel_, + sdwa_dst_unused_); + } bool default_encoding(); bool has_lit(); bool has_dpp(); @@ -533,6 +539,12 @@ class Vop2 : public IsaInstruction { public: Vop2(std::string_view mnemonic, const Vop2MachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::vop_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_, sdwa_dst_sel_, + sdwa_dst_unused_); + } bool default_encoding(); bool has_lit(); bool has_dpp(); diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/cdna2/encodings.h b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/cdna2/encodings.h index 4ef380e9b919..dba1d1391eb1 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/cdna2/encodings.h +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/cdna2/encodings.h @@ -460,6 +460,12 @@ class Vop1 : public IsaInstruction { public: Vop1(std::string_view mnemonic, const Vop1MachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::vop_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_, sdwa_dst_sel_, + sdwa_dst_unused_); + } bool default_encoding(); bool has_lit(); bool has_dpp(); @@ -517,6 +523,12 @@ class Vop2 : public IsaInstruction { public: Vop2(std::string_view mnemonic, const Vop2MachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::vop_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_, sdwa_dst_sel_, + sdwa_dst_unused_); + } bool default_encoding(); bool has_lit(); bool has_dpp(); diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/cdna3/encodings.h b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/cdna3/encodings.h index 6e8c1ea1bd07..74a6a496356f 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/cdna3/encodings.h +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/cdna3/encodings.h @@ -448,6 +448,12 @@ class Vop1 : public IsaInstruction { public: Vop1(std::string_view mnemonic, const Vop1MachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::vop_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_, sdwa_dst_sel_, + sdwa_dst_unused_); + } bool default_encoding(); bool has_lit(); bool has_dpp(); @@ -505,6 +511,12 @@ class Vop2 : public IsaInstruction { public: Vop2(std::string_view mnemonic, const Vop2MachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::vop_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_, sdwa_dst_sel_, + sdwa_dst_unused_); + } bool default_encoding(); bool has_lit(); bool has_dpp(); diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/cdna4/encodings.h b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/cdna4/encodings.h index 8cdfabf9d78b..89c458f2116e 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/cdna4/encodings.h +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/cdna4/encodings.h @@ -450,6 +450,12 @@ class Vop1 : public IsaInstruction { public: Vop1(std::string_view mnemonic, const Vop1MachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::vop_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_, sdwa_dst_sel_, + sdwa_dst_unused_); + } bool default_encoding(); bool has_lit(); bool has_dpp(); @@ -507,6 +513,12 @@ class Vop2 : public IsaInstruction { public: Vop2(std::string_view mnemonic, const Vop2MachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::vop_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_, sdwa_dst_sel_, + sdwa_dst_unused_); + } bool default_encoding(); bool has_lit(); bool has_dpp(); diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/gfx1250/encodings.h b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/gfx1250/encodings.h index 97115146fc3d..89d7eee13bb2 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/gfx1250/encodings.h +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/gfx1250/encodings.h @@ -434,6 +434,12 @@ class Vop1 : public IsaInstruction { public: Vop1(std::string_view mnemonic, const Vop1MachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::vop_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_, sdwa_dst_sel_, + sdwa_dst_unused_); + } bool default_encoding(); bool has_lit(); bool has_lit64(); @@ -496,6 +502,12 @@ class Vop2 : public IsaInstruction { public: Vop2(std::string_view mnemonic, const Vop2MachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::vop_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_, sdwa_dst_sel_, + sdwa_dst_unused_); + } bool default_encoding(); bool has_lit(); bool has_lit64(); @@ -530,6 +542,11 @@ class Vop3 : public IsaInstruction { public: Vop3(std::string_view mnemonic, const Vop3MachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::dpp_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_); + } bool has_lit_0(); bool has_lit_1(); bool has_lit_0_and_has_lit_1(); @@ -555,6 +572,11 @@ class Vop3p : public IsaInstruction { public: Vop3p(std::string_view mnemonic, const Vop3pMachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::dpp_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_); + } bool has_lit_0(); bool has_lit_1(); bool has_lit_0_and_has_lit_1(); @@ -626,6 +648,11 @@ class Vop3SdstEnc : public IsaInstruction { public: Vop3SdstEnc(std::string_view mnemonic, const Vop3SdstEncMachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::dpp_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_); + } bool has_lit_0(); bool has_lit_1(); bool has_lit_0_and_has_lit_1(); diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/rdna1/encodings.h b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/rdna1/encodings.h index ce63bb2dce27..c313ba35f12f 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/rdna1/encodings.h +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/rdna1/encodings.h @@ -425,6 +425,12 @@ class Vop1 : public IsaInstruction { public: Vop1(std::string_view mnemonic, const Vop1MachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::vop_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_, sdwa_dst_sel_, + sdwa_dst_unused_); + } bool default_encoding(); bool has_lit(); using OpEncoding = Vop1MachineInst; @@ -480,6 +486,12 @@ class Vop2 : public IsaInstruction { public: Vop2(std::string_view mnemonic, const Vop2MachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::vop_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_, sdwa_dst_sel_, + sdwa_dst_unused_); + } bool default_encoding(); bool has_lit(); bool hasImpliedLiteral(); diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/rdna2/encodings.h b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/rdna2/encodings.h index 2b12636e6541..c46f5b91fabf 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/rdna2/encodings.h +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/rdna2/encodings.h @@ -421,6 +421,12 @@ class Vop1 : public IsaInstruction { public: Vop1(std::string_view mnemonic, const Vop1MachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::vop_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_, sdwa_dst_sel_, + sdwa_dst_unused_); + } bool default_encoding(); bool has_lit(); using OpEncoding = Vop1MachineInst; @@ -476,6 +482,12 @@ class Vop2 : public IsaInstruction { public: Vop2(std::string_view mnemonic, const Vop2MachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::vop_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_, sdwa_dst_sel_, + sdwa_dst_unused_); + } bool default_encoding(); bool has_lit(); bool hasImpliedLiteral(); diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/rdna3/encodings.h b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/rdna3/encodings.h index a1bd57699551..f72b4b4656ce 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/rdna3/encodings.h +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/rdna3/encodings.h @@ -399,6 +399,12 @@ class Vop1 : public IsaInstruction { public: Vop1(std::string_view mnemonic, const Vop1MachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::vop_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_, sdwa_dst_sel_, + sdwa_dst_unused_); + } bool default_encoding(); bool has_lit(); using OpEncoding = Vop1MachineInst; @@ -455,6 +461,12 @@ class Vop2 : public IsaInstruction { public: Vop2(std::string_view mnemonic, const Vop2MachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::vop_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_, sdwa_dst_sel_, + sdwa_dst_unused_); + } bool default_encoding(); bool has_lit(); bool hasImpliedLiteral(); @@ -486,6 +498,11 @@ class Vop3 : public IsaInstruction { public: Vop3(std::string_view mnemonic, const Vop3MachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::dpp_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_); + } bool has_lit_0(); bool has_lit_1(); bool has_lit_0_has_lit_1(); @@ -509,6 +526,11 @@ class Vop3p : public IsaInstruction { public: Vop3p(std::string_view mnemonic, const Vop3pMachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::dpp_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_); + } bool has_lit_0(); bool has_lit_1(); bool has_lit_0_has_lit_1(); @@ -595,6 +617,11 @@ class Vop3SdstEnc : public IsaInstruction { public: Vop3SdstEnc(std::string_view mnemonic, const Vop3SdstEncMachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::dpp_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_); + } bool has_lit_0(); bool has_lit_1(); bool has_lit_0_has_lit_1(); diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/rdna3_5/encodings.h b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/rdna3_5/encodings.h index 65034b26cf47..14497b489259 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/rdna3_5/encodings.h +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/rdna3_5/encodings.h @@ -414,6 +414,12 @@ class Vop1 : public IsaInstruction { public: Vop1(std::string_view mnemonic, const Vop1MachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::vop_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_, sdwa_dst_sel_, + sdwa_dst_unused_); + } bool default_encoding(); bool has_lit(); using OpEncoding = Vop1MachineInst; @@ -470,6 +476,12 @@ class Vop2 : public IsaInstruction { public: Vop2(std::string_view mnemonic, const Vop2MachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::vop_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_, sdwa_dst_sel_, + sdwa_dst_unused_); + } bool default_encoding(); bool has_lit(); bool hasImpliedLiteral(); @@ -501,6 +513,11 @@ class Vop3 : public IsaInstruction { public: Vop3(std::string_view mnemonic, const Vop3MachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::dpp_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_); + } bool has_lit_0(); bool has_lit_1(); bool has_lit_0_has_lit_1(); @@ -524,6 +541,11 @@ class Vop3p : public IsaInstruction { public: Vop3p(std::string_view mnemonic, const Vop3pMachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::dpp_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_); + } bool has_lit_0(); bool has_lit_1(); bool has_lit_0_has_lit_1(); @@ -610,6 +632,11 @@ class Vop3SdstEnc : public IsaInstruction { public: Vop3SdstEnc(std::string_view mnemonic, const Vop3SdstEncMachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::dpp_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_); + } bool has_lit_0(); bool has_lit_1(); bool has_lit_0_has_lit_1(); diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/rdna4/encodings.h b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/rdna4/encodings.h index 6fad3d9bc192..8de017d63bcb 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/rdna4/encodings.h +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/rdna4/encodings.h @@ -431,6 +431,12 @@ class Vop1 : public IsaInstruction { public: Vop1(std::string_view mnemonic, const Vop1MachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::vop_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_, sdwa_dst_sel_, + sdwa_dst_unused_); + } bool default_encoding(); bool has_lit(); using OpEncoding = Vop1MachineInst; @@ -487,6 +493,12 @@ class Vop2 : public IsaInstruction { public: Vop2(std::string_view mnemonic, const Vop2MachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::vop_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_, sdwa_dst_sel_, + sdwa_dst_unused_); + } bool default_encoding(); bool has_lit(); bool hasImpliedLiteral(); @@ -518,6 +530,11 @@ class Vop3 : public IsaInstruction { public: Vop3(std::string_view mnemonic, const Vop3MachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::dpp_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_); + } bool has_lit_0(); bool has_lit_1(); bool has_lit_0_has_lit_1(); @@ -541,6 +558,11 @@ class Vop3p : public IsaInstruction { public: Vop3p(std::string_view mnemonic, const Vop3pMachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::dpp_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_); + } bool has_lit_0(); bool has_lit_1(); bool has_lit_0_has_lit_1(); @@ -639,6 +661,11 @@ class Vop3SdstEnc : public IsaInstruction { public: Vop3SdstEnc(std::string_view mnemonic, const Vop3SdstEncMachineInst *inst, ExecuteFn exec_fn); void implicit_uses(RegisterSet &uses) const override; + std::optional vgpr_write_hook_filter(uint32_t wf_size, + uint64_t exec) const override { + return amdgpu::dpp_write_hook_filter(inst_.src0, wf_size, exec, dpp_ctrl_, dpp_row_mask_, + dpp_bank_mask_, dpp_bound_ctrl_); + } bool has_lit_0(); bool has_lit_1(); bool has_lit_0_has_lit_1(); diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/shared/dpp_sdwa_ops.h b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/shared/dpp_sdwa_ops.h index 20e334d60a27..da34670b9719 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/shared/dpp_sdwa_ops.h +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/shared/dpp_sdwa_ops.h @@ -389,6 +389,19 @@ enum SdwaUnused : uint32_t { UNUSED_PRESERVE = 2, ///< Keep the original destination register value. }; +/// @brief Return the destination bytes architecturally written by an SDWA op. +/// +/// PAD and SEXT define every non-selected byte, so they write the full dword. +/// PRESERVE leaves those bytes untouched and reports only the selected byte or +/// word. DWORD always writes the full destination. +inline uint8_t sdwa_dst_write_mask(uint32_t dst_sel, uint32_t dst_unused) { + if (dst_sel == DWORD || dst_unused != UNUSED_PRESERVE) + return 0xF; + if (dst_sel <= BYTE_3) + return static_cast(1u << dst_sel); + return dst_sel == WORD_0 ? 0x3 : 0xC; +} + /// @brief Extract a sub-dword from a source value per SDWA sel. /// /// @param val The full 32-bit source value. @@ -469,6 +482,32 @@ inline uint32_t sdwa_clamp_f32(uint32_t result) { } } // namespace sdwa + +/// @brief Build the architectural VGPR hook filter for VOP1/VOP2 suffixes. +inline std::optional +dpp_write_hook_filter(uint32_t src0, uint32_t wf_size, uint64_t exec, uint32_t dpp_ctrl, + uint32_t dpp_row_mask, uint32_t dpp_bank_mask, uint32_t dpp_bound_ctrl) { + if (src0 == SRC_DPP) + return VgprWriteHookFilter{ + exec & dpp::dpp_write_mask(wf_size, dpp_ctrl, dpp_row_mask, dpp_bank_mask, dpp_bound_ctrl), + // DPP changes destination lanes, but not the instruction's architectural byte mask. + 0, true}; + return std::nullopt; +} + +inline std::optional +vop_write_hook_filter(uint32_t src0, uint32_t wf_size, uint64_t exec, uint32_t dpp_ctrl, + uint32_t dpp_row_mask, uint32_t dpp_bank_mask, uint32_t dpp_bound_ctrl, + uint32_t sdwa_dst_sel, uint32_t sdwa_dst_unused) { + if (auto filter = dpp_write_hook_filter(src0, wf_size, exec, dpp_ctrl, dpp_row_mask, + dpp_bank_mask, dpp_bound_ctrl)) + return filter; + if (src0 == SRC_SDWA) + return VgprWriteHookFilter{exec, sdwa::sdwa_dst_write_mask(sdwa_dst_sel, sdwa_dst_unused), + sdwa_dst_sel != sdwa::DWORD}; + return std::nullopt; +} + } // namespace amdgpu } // namespace rocjitsu diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/shared/mma_exec.h b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/shared/mma_exec.h index 32f28b115267..33aeabe679f6 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/shared/mma_exec.h +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/shared/mma_exec.h @@ -1648,7 +1648,7 @@ inline void exec_wmma_f32_16x16x32_f16(amdgpu::ComputeUnitCore &cu, uint32_t dst const_acc, c_modifier); return; } else { - if (util::force_scalar() || util::native::size() != 16) { + if (!cu.plugin_group().empty() || util::force_scalar() || util::native::size() != 16) { exec_wmma_f32(cu, M, N, K, in_bits, dst, s0, s1, s2, amdgpu::extract_f16, amdgpu::extract_f16, const_acc, c_modifier); return; @@ -1724,7 +1724,7 @@ inline void exec_wmma_f32_16x16x32_bf16(amdgpu::ComputeUnitCore &cu, uint32_t ds const_acc, c_modifier); return; } else { - if (util::force_scalar() || util::native::size() != 16) { + if (!cu.plugin_group().empty() || util::force_scalar() || util::native::size() != 16) { exec_wmma_f32(cu, M, N, K, in_bits, dst, s0, s1, s2, amdgpu::extract_bf16, amdgpu::extract_bf16, const_acc, c_modifier); return; @@ -1838,7 +1838,7 @@ void exec_wmma_f32_f8_spec(amdgpu::ComputeUnitCore &cu, uint32_t dst, uint32_t s fallback(); return; } else { - if (util::force_scalar() || util::native::size() != 16) { + if (!cu.plugin_group().empty() || util::force_scalar() || util::native::size() != 16) { fallback(); return; } @@ -1914,7 +1914,7 @@ void exec_wmma_f32_f32_spec(amdgpu::ComputeUnitCore &cu, uint32_t dst, uint32_t const_acc, c_modifier); return; } else { - if (util::force_scalar() || util::native::size() != 16) { + if (!cu.plugin_group().empty() || util::force_scalar() || util::native::size() != 16) { exec_wmma_f32(cu, M, N, K, in_bits, dst, s0, s1, s2, amdgpu::extract_f32, amdgpu::extract_f32, const_acc, c_modifier); return; @@ -2354,7 +2354,7 @@ void exec_wmma_f16_spec(amdgpu::ComputeUnitCore &cu, uint32_t dst, uint32_t s0, fallback(); return; } else { - if (util::force_scalar() || util::native::size() != 16) { + if (!cu.plugin_group().empty() || util::force_scalar() || util::native::size() != 16) { fallback(); return; } @@ -2478,7 +2478,7 @@ void exec_wmma_bf16_spec(amdgpu::ComputeUnitCore &cu, uint32_t dst, uint32_t s0, fallback(); return; } else { - if (util::force_scalar() || util::native::size() != 16) { + if (!cu.plugin_group().empty() || util::force_scalar() || util::native::size() != 16) { fallback(); return; } @@ -2576,7 +2576,7 @@ void exec_wmma_f16_f8_spec(amdgpu::ComputeUnitCore &cu, uint32_t dst, uint32_t s fallback(); return; } else { - if (util::force_scalar() || util::native::size() != 16) { + if (!cu.plugin_group().empty() || util::force_scalar() || util::native::size() != 16) { fallback(); return; } @@ -3121,7 +3121,7 @@ inline void exec_wmma_i32_16x16x64_iu8(amdgpu::ComputeUnitCore &cu, uint32_t dst fallback(); return; } else { - if (util::force_scalar() || util::native::size() != 16) { + if (!cu.plugin_group().empty() || util::force_scalar() || util::native::size() != 16) { fallback(); return; } @@ -3817,8 +3817,8 @@ void exec_f32_mfma_f32_spec(amdgpu::ComputeUnitCore &cu, uint32_t dst, uint32_t const_acc, cbsz, abid, blgp); return; } else { - if (util::force_scalar() || cbsz != 0 || blgp != 0 || util::native::size() != 16 || - cu.wf_size() != 64) { + if (!cu.plugin_group().empty() || util::force_scalar() || cbsz != 0 || blgp != 0 || + util::native::size() != 16 || cu.wf_size() != 64) { exec_f32(cu, M, N, K, BATCH, in_bits, dst, s0, s1, s2, amdgpu::extract_f32, amdgpu::extract_f32, const_acc, cbsz, abid, blgp); return; @@ -3892,8 +3892,8 @@ void exec_f32_mfma_f16_spec(amdgpu::ComputeUnitCore &cu, uint32_t dst, uint32_t const_acc, cbsz, abid, blgp); return; } else { - if (util::force_scalar() || cbsz != 0 || blgp != 0 || util::native::size() != 16 || - cu.wf_size() != 64) { + if (!cu.plugin_group().empty() || util::force_scalar() || cbsz != 0 || blgp != 0 || + util::native::size() != 16 || cu.wf_size() != 64) { exec_f32(cu, M, N, K, B, in_bits, dst, s0, s1, s2, amdgpu::extract_f16, amdgpu::extract_f16, const_acc, cbsz, abid, blgp); return; @@ -3981,8 +3981,8 @@ void exec_f32_mfma_bf16_spec(amdgpu::ComputeUnitCore &cu, uint32_t dst, uint32_t const_acc, cbsz, abid, blgp); return; } else { - if (util::force_scalar() || cbsz != 0 || blgp != 0 || util::native::size() != 16 || - cu.wf_size() != 64) { + if (!cu.plugin_group().empty() || util::force_scalar() || cbsz != 0 || blgp != 0 || + util::native::size() != 16 || cu.wf_size() != 64) { exec_f32(cu, M, N, K, B, in_bits, dst, s0, s1, s2, amdgpu::extract_bf16, amdgpu::extract_bf16, const_acc, cbsz, abid, blgp); return; @@ -4071,8 +4071,8 @@ void exec_f32_mfma_f8_spec(amdgpu::ComputeUnitCore &cu, uint32_t dst, uint32_t s exec_f32(cu, M, N, K, B, in_bits, dst, s0, s1, s2, ea, eb, const_acc, cbsz, abid, blgp); return; } else { - if (util::force_scalar() || cbsz != 0 || blgp != 0 || util::native::size() != 16 || - cu.wf_size() != 64) { + if (!cu.plugin_group().empty() || util::force_scalar() || cbsz != 0 || blgp != 0 || + util::native::size() != 16 || cu.wf_size() != 64) { exec_f32(cu, M, N, K, B, in_bits, dst, s0, s1, s2, ea, eb, const_acc, cbsz, abid, blgp); return; } @@ -4159,7 +4159,8 @@ void exec_i32_mfma_i8_spec(amdgpu::ComputeUnitCore &cu, uint32_t dst, uint32_t s exec_i32_i8(cu, M, N, K, B, dst, s0, s1, s2, const_acc); return; } else { - if (util::force_scalar() || util::native::size() != 16 || cu.wf_size() != 64) { + if (!cu.plugin_group().empty() || util::force_scalar() || util::native::size() != 16 || + cu.wf_size() != 64) { exec_i32_i8(cu, M, N, K, B, dst, s0, s1, s2, const_acc); return; } diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/shared/simd_glue.h b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/shared/simd_glue.h index 144591d3c07c..11a8fb9ccc18 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/shared/simd_glue.h +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/shared/simd_glue.h @@ -308,11 +308,12 @@ inline void write_vop3_true16_dst(const Operand &dst, Wavefront &wf, uint32_t la uint32_t off = reg->index + (wf.vgpr_msb_for_role(dst.vgpr_msb_role()) << 8); uint32_t voff = wf.gpr_idx_en() ? apply_gpr_idx(wf, off, true) : off; uint32_t idx = wf.vgpr_alloc().base + voff; - uint32_t old_dst = wf.cu().read_vgpr(idx, lane); + uint32_t old_dst = low_dst_zeroes_high ? 0 : wf.cu().read_vgpr_raw(idx, lane); uint32_t merged = dst_hi ? ((old_dst & 0x0000ffffu) | (src_half << 16)) : (low_dst_zeroes_high ? src_half : ((old_dst & 0xffff0000u) | src_half)); - wf.cu().write_vgpr(idx, lane, merged); + const uint8_t byte_mask = low_dst_zeroes_high ? 0xF : (dst_hi ? 0xC : 0x3); + wf.cu().write_vgpr_masked(idx, lane, merged, byte_mask); return; } dst.write_lane(wf, lane, dst_hi ? (src_half << 16) : src_half); @@ -494,7 +495,7 @@ inline void write_simd(const Op &op, Wavefront &wf, uint32_t lane_base, util::na uint64_t mask) { static_assert(sizeof(T) == sizeof(uint32_t)); constexpr std::size_t W = util::native_width_v; - if (VgprStorage *r = SimdAccess::vgpr_storage_mut(op, wf)) { + if (VgprStorage *r = SimdAccess::vgpr_storage_mut(op, wf); r && wf.cu().plugin_group().empty()) { r->template simd_store(lane_base, v, mask); return; } @@ -516,7 +517,7 @@ inline void write_simd_at(VgprStorage *rd, const Op &op, Wavefront &wf, uint32_t util::native v, uint64_t mask) { static_assert(sizeof(T) == sizeof(uint32_t)); constexpr std::size_t W = util::native_width_v; - if (rd) { + if (rd && wf.cu().plugin_group().empty()) { rd->template simd_store(base, v, mask); return; } @@ -537,7 +538,30 @@ inline void write_vop3_true16_simd_at(VgprStorage *rd, const Op &op, Wavefront & : ((old & util::broadcast(0xffff0000u)) | src_half); if (!(opsel & 0x8u) && cdna_low_dst_zeroes_high && cdna_vop3_low_dst_zeroes_high(wf)) merged = src_half; - write_simd_at(rd, op, wf, base, merged, mask); + if (wf.cu().plugin_group().empty()) { + write_simd_at(rd, op, wf, base, merged, mask); + return; + } + + auto reg = op.to_register_ref(); + if (!reg || reg->cls != RegClass::VGPR) { + write_simd_at(rd, op, wf, base, merged, mask); + return; + } + + constexpr std::size_t W = util::native_width_v; + alignas(util::native) uint32_t buf[W]; + util::blit_to_buffer(buf, merged); + uint32_t off = reg->index + (wf.vgpr_msb_for_role(op.vgpr_msb_role()) << 8); + uint32_t voff = wf.gpr_idx_en() ? apply_gpr_idx(wf, off, true) : off; + uint32_t idx = wf.vgpr_alloc().base + voff; + const bool dst_hi = (opsel & 0x8u) != 0; + const bool low_dst_zeroes_high = + !dst_hi && cdna_low_dst_zeroes_high && cdna_vop3_low_dst_zeroes_high(wf); + const uint8_t byte_mask = low_dst_zeroes_high ? 0xF : (dst_hi ? 0xC : 0x3); + for (uint32_t i = 0; i < W; ++i) + if (mask & (1ULL << i)) + wf.cu().write_vgpr_masked(idx, base + i, buf[i], byte_mask); } /// Narrow (native_width64-wide) counterpart of write_simd_at, for the f64->32-bit @@ -550,7 +574,7 @@ inline void write_simd_narrow_at(VgprStorage *rd, const Op &op, Wavefront &wf, u util::narrow32 v, uint64_t mask) { static_assert(sizeof(T) == sizeof(uint32_t)); constexpr std::size_t W = util::native_width64; - if (rd) { + if (rd && wf.cu().plugin_group().empty()) { rd->template simd_store_narrow(base, v, mask); return; } @@ -616,7 +640,7 @@ inline void write_simd64(const Op &op, Wavefront &wf, uint32_t lane_base, util:: static_assert(sizeof(T) == sizeof(uint64_t)); constexpr std::size_t W = util::native_width64; VgprStoragePair64 p = simd_dst_reg64(op, wf); - if (p.lo) { + if (p.lo && wf.cu().plugin_group().empty()) { p.lo->template simd_store64(*p.hi, lane_base, v, mask); return; } @@ -647,7 +671,7 @@ inline void write_simd64_at(VgprStoragePair64 p, const Op &op, Wavefront &wf, ui util::native v, uint64_t mask) { static_assert(sizeof(T) == sizeof(uint64_t)); constexpr std::size_t W = util::native_width64; - if (p.lo) { + if (p.lo && wf.cu().plugin_group().empty()) { p.lo->template simd_store64(*p.hi, base, v, mask); return; } @@ -3543,8 +3567,8 @@ template template requires(util::has_stdx_simd) [[nodiscard]] inline bool try_execute_vop3p_pk_binary_f32_simd(Inst &inst, Wavefront &wf, Op op) { - if (simd_force_scalar() || !inst.src0.simd_capable() || !inst.src1.simd_capable() || - !inst.vdst.simd_capable()) + if (simd_force_scalar() || !wf.cu().plugin_group().empty() || !inst.src0.simd_capable() || + !inst.src1.simd_capable() || !inst.vdst.simd_capable()) return false; if (inst.inst_.op_sel != 0u || inst.inst_.op_sel_hi != 3u) return false; @@ -3584,8 +3608,8 @@ template template requires(util::has_stdx_simd) [[nodiscard]] inline bool try_execute_vop3p_pk_ternary_f32_simd(Inst &inst, Wavefront &wf, Op op) { - if (simd_force_scalar() || !inst.src0.simd_capable() || !inst.src1.simd_capable() || - !inst.src2.simd_capable() || !inst.vdst.simd_capable()) + if (simd_force_scalar() || !wf.cu().plugin_group().empty() || !inst.src0.simd_capable() || + !inst.src1.simd_capable() || !inst.src2.simd_capable() || !inst.vdst.simd_capable()) return false; if (inst.inst_.op_sel != 0u || inst.inst_.op_sel_hi != 3u || inst.inst_.op_sel_hi_2 != 1u) return false; diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/instruction.h b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/instruction.h index d7be04f3fdcc..26ace49f9b5d 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/instruction.h +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/instruction.h @@ -51,6 +51,15 @@ enum InstFlags : uint64_t { class BasicBlock; +/// @brief Architectural VGPR write shape for instructions whose execution +/// implementation performs temporary destination writes. +struct VgprWriteHookFilter { + uint64_t lane_mask = 0; + /// Nonzero overrides the observed write mask; zero preserves it. + uint8_t byte_mask = 0xF; + bool suppress_destination_merge_reads = false; +}; + /// @brief Base class for opaque per-instruction dynamic state. /// /// @details ISA backends attach subclasses to carry mutable pipeline state (e.g., @@ -212,6 +221,18 @@ class Instruction : public util::IListNode branch_offset_bytes() const { return std::nullopt; } + /// @brief Describe the architectural VGPR write produced by this instruction. + /// + /// Most instructions report writes directly as they update the register file. + /// DPP and SDWA execution wrappers additionally write temporary results and + /// restore/merge the destination. Their encoding bases override this method so + /// the compute unit can collapse those implementation writes into one + /// architectural callback per destination lane. + [[nodiscard]] virtual std::optional + vgpr_write_hook_filter(uint32_t /*wf_size*/, uint64_t /*exec*/) const { + return std::nullopt; + } + /// @brief Add registers implicitly read by this instruction. /// /// @details Used for encoded fields that affect execution but are not part of diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/isa_operand_simd_inl.h b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/isa_operand_simd_inl.h index f21eb49f4992..11fc30cb31a7 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/isa_operand_simd_inl.h +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/isa_operand_simd_inl.h @@ -70,7 +70,7 @@ void AmdgpuIsaOperand::write_lane_chunk(amdgpu::Wavefront &wf, uint32_t lan uint32_t voff = wf.gpr_idx_en() ? amdgpu::apply_gpr_idx(wf, *off, true) : *off; uint32_t reg = wf.vgpr_alloc().base + voff; uint64_t full_mask = util::mask(static_cast(count)); - if ((mask & full_mask) == full_mask) { + if ((mask & full_mask) == full_mask && wf.cu().plugin_group().empty()) { uint8_t *dst = wf.cu().vgpr_data(reg); std::memcpy(dst + lane_base * sizeof(uint32_t), vals, count * sizeof(uint32_t)); return; diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/command_processor.cpp b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/command_processor.cpp index bd869353fe21..aeb8ee7f2153 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/command_processor.cpp +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/command_processor.cpp @@ -309,13 +309,13 @@ void CommandProcessor::init_wavefront_regs(ComputeUnitCore *cu, Wavefront *wf, for (uint32_t lane = 0; lane < cu->wf_size(); ++lane) { const WorkitemCoord id = workitem_local_coord(pkt, wf_index_in_wg, lane, cu->wf_size()); if (packed_tid_) { - cu->write_vgpr(vbase, lane, pack_workitem_id(id, pkt.enable_vgpr_workitem_id)); + cu->write_vgpr_raw(vbase, lane, pack_workitem_id(id, pkt.enable_vgpr_workitem_id)); } else { - cu->write_vgpr(vbase, lane, id.x); + cu->write_vgpr_raw(vbase, lane, id.x); if (pkt.enable_vgpr_workitem_id >= 1) - cu->write_vgpr(vbase + 1, lane, id.y); + cu->write_vgpr_raw(vbase + 1, lane, id.y); if (pkt.enable_vgpr_workitem_id >= 2) - cu->write_vgpr(vbase + 2, lane, id.z); + cu->write_vgpr_raw(vbase + 2, lane, id.z); } } diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/compute_unit.cpp b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/compute_unit.cpp index e8bc30ae8b97..eb176e90ecd1 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/compute_unit.cpp +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/compute_unit.cpp @@ -21,6 +21,7 @@ #include "util/log.h" #include +#include #include #include #include @@ -176,6 +177,119 @@ void ComputeUnitCore::reset_all_wf() { } } +void ComputeUnitCore::begin_vgpr_write_hook_filter(const Instruction &inst, const Wavefront &wf) { + cancel_vgpr_write_hook_filter(); + if (plugin_group_->empty()) + return; + + auto filter = inst.vgpr_write_hook_filter(wf.wf_size(), wf.exec()); + if (!filter) + return; + + auto &state = vgpr_write_hook_filter_; + for (int i = 0; i < inst.num_dst_operands(); ++i) { + const Operand *dst = inst.dst_operand(i); + if (!dst) + continue; + auto ref = dst->to_register_ref(); + if (!ref || ref->cls != RegClass::VGPR) + continue; + + uint32_t off = ref->index + (wf.vgpr_msb_for_role(dst->vgpr_msb_role()) << 8); + uint32_t voff = wf.gpr_idx_en() ? apply_gpr_idx(wf, off, true) : off; + for (uint32_t r = 0; r < ref->width; ++r) { + uint32_t physical_reg = wf.vgpr_alloc().base + voff + r; + auto duplicate = std::find_if(state.destinations.begin(), state.destinations.end(), + [physical_reg](const FilteredVgprHooks &entry) { + return entry.physical_reg == physical_reg; + }); + if (duplicate == state.destinations.end()) + state.destinations.push_back(FilteredVgprHooks{physical_reg}); + } + } + if (state.destinations.empty()) + return; + + state.wf = &wf; + state.lane_mask = filter->lane_mask; + state.byte_mask = filter->byte_mask; + state.suppress_destination_merge_reads = filter->suppress_destination_merge_reads; + state.initial_destination_merge_read_lanes = + filter->suppress_destination_merge_reads ? wf.exec() : 0; +} + +void ComputeUnitCore::flush_vgpr_write_hook_filter() { + auto &state = vgpr_write_hook_filter_; + if (!state.active()) + return; + + const Wavefront *wf = state.wf; + state.wf = nullptr; + for (const auto &dst : state.destinations) { + uint64_t pending = dst.pending_writes; + while (pending) { + uint32_t lane = static_cast(std::countr_zero(pending)); + uint8_t byte_mask = state.byte_mask ? state.byte_mask : dst.byte_masks[lane]; + plugin_group_->onAmdgpuWriteVgpr(wf, dst.physical_reg, lane, byte_mask); + pending &= pending - 1; + } + } + state.destinations.clear(); +} + +void ComputeUnitCore::cancel_vgpr_write_hook_filter() { vgpr_write_hook_filter_.clear(); } + +bool ComputeUnitCore::filter_vgpr_read_notification(const Wavefront *wf, uint32_t reg_idx, + uint32_t lane_begin, uint32_t lane_end, + uint8_t byte_mask) const { + auto &state = vgpr_write_hook_filter_; + if (!state.active() || state.wf != wf || !state.suppress_destination_merge_reads) + return false; + + auto dst = std::find_if( + state.destinations.begin(), state.destinations.end(), + [reg_idx](const FilteredVgprHooks &entry) { return entry.physical_reg == reg_idx; }); + if (dst == state.destinations.end()) + return false; + + for (uint32_t lane = lane_begin; lane < lane_end; ++lane) { + assert(lane < 64); + uint64_t bit = 1ULL << lane; + bool suppress = false; + if (dst->writes & bit) { + suppress = true; + } else if ((state.initial_destination_merge_read_lanes & bit) && !(dst->initial_reads & bit)) { + dst->initial_reads |= bit; + suppress = true; + } + if (!suppress) + plugin_group_->onAmdgpuReadVgprs(wf, reg_idx, lane, lane + 1, byte_mask); + } + return true; +} + +bool ComputeUnitCore::filter_vgpr_write_notification(const Wavefront *wf, uint32_t reg_idx, + uint32_t lane, uint8_t byte_mask) const { + auto &state = vgpr_write_hook_filter_; + if (!state.active() || state.wf != wf) + return false; + + auto dst = std::find_if( + state.destinations.begin(), state.destinations.end(), + [reg_idx](const FilteredVgprHooks &entry) { return entry.physical_reg == reg_idx; }); + if (dst == state.destinations.end()) + return false; + + assert(lane < 64); + uint64_t bit = 1ULL << lane; + dst->writes |= bit; + if (state.lane_mask & bit) { + dst->pending_writes |= bit; + dst->byte_masks[lane] |= byte_mask; + } + return true; +} + void ComputeUnitCore::retire_halted_wfs_no_lds_reset() { for (auto &w : wfs_) { if (w->is_halted() && w->sgpr_alloc().count > 0) { @@ -408,7 +522,14 @@ void ComputeUnitCore::issue_instruction(Wavefront *active) { } } - execute_instruction(inst, *active); + begin_vgpr_write_hook_filter(*inst, *active); + try { + execute_instruction(inst, *active); + } catch (...) { + cancel_vgpr_write_hook_filter(); + throw; + } + flush_vgpr_write_hook_filter(); plugin_group_->onAmdgpuAfterExecuteInstruction(active->pc, *inst, *active); if constexpr (util::Logger::group_enabled(util::Logger::GROUP_VM)) { diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/compute_unit.h b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/compute_unit.h index 64f3f7a663b8..13e65a84d83b 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/compute_unit.h +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/compute_unit.h @@ -384,22 +384,55 @@ class ComputeUnitCore : public simdojo::CompositeComponent { void notify_vgpr_read(const Wavefront *wf, uint32_t reg_idx, uint32_t lane_begin, uint32_t lane_end, uint8_t byte_mask = 0xF) const { - if (wf) + if (wf && !filter_vgpr_read_notification(wf, reg_idx, lane_begin, lane_end, byte_mask)) plugin_group_->onAmdgpuReadVgprs(wf, reg_idx, lane_begin, lane_end, byte_mask); } + void notify_vgpr_write(const Wavefront *wf, uint32_t reg_idx, uint32_t lane, + uint8_t byte_mask = 0xF) const { + if (wf && !filter_vgpr_write_notification(wf, reg_idx, lane, byte_mask)) + plugin_group_->onAmdgpuWriteVgpr(wf, reg_idx, lane, byte_mask); + } + /// @brief Read a vector register lane from the physical VGPR file. /// @param reg_idx Physical register index. /// @param lane Lane index within the wavefront. /// @returns Lane value. virtual uint32_t read_vgpr(uint32_t reg_idx, uint32_t lane) const = 0; + /// @brief Read a vector register lane while notifying plugins with a byte mask. + virtual uint32_t read_vgpr_masked(uint32_t reg_idx, uint32_t lane, uint8_t byte_mask) const = 0; + + /// @brief Read a vector register lane without plugin hooks. + /// + /// Used by architectural memory writeback merge paths that need the old + /// dword value to preserve untouched D16 bits. + virtual uint32_t read_vgpr_raw(uint32_t reg_idx, uint32_t lane) const = 0; + /// @brief Write a vector register lane in the physical VGPR file. /// @param reg_idx Physical register index. /// @param lane Lane index within the wavefront. /// @param val Value to write. virtual void write_vgpr(uint32_t reg_idx, uint32_t lane, uint32_t val) = 0; + /// @brief Write a vector register lane while notifying plugins with a byte mask. + virtual void write_vgpr_masked(uint32_t reg_idx, uint32_t lane, uint32_t val, + uint8_t byte_mask) = 0; + + /// @brief Write a vector register lane without plugin write hooks. + /// + /// Used by simulator setup and architectural writeback paths that must update + /// the physical register file without reporting an instruction destination write. + virtual void write_vgpr_raw(uint32_t reg_idx, uint32_t lane, uint32_t val) = 0; + + /// @brief Write a memory instruction result without plugin write hooks. + /// + /// The race detector observes memory results when the instruction is routed. + /// Completion writes must not look like a new instruction destination write. + void write_vgpr_from_memory(uint32_t reg_idx, uint32_t lane, uint32_t val) { + write_vgpr_raw(reg_idx, lane, val); + } + /// @brief Return a pointer to a wavefront's SGPR data in the physical file. /// @param base Base register index in the SGPR file. /// @returns Pointer to the contiguous SGPR data. @@ -492,6 +525,14 @@ class ComputeUnitCore : public simdojo::CompositeComponent { /// @brief Tick all memory pipelines (called at the start of step in clocked mode). void tick_pipelines(); + void begin_vgpr_write_hook_filter(const Instruction &inst, const Wavefront &wf); + void flush_vgpr_write_hook_filter(); + void cancel_vgpr_write_hook_filter(); + bool filter_vgpr_read_notification(const Wavefront *wf, uint32_t reg_idx, uint32_t lane_begin, + uint32_t lane_end, uint8_t byte_mask) const; + bool filter_vgpr_write_notification(const Wavefront *wf, uint32_t reg_idx, uint32_t lane, + uint8_t byte_mask) const; + /// @brief Route a memory instruction into the appropriate pipeline. /// @param inst The memory instruction (ownership transferred). /// @param wf The issuing wavefront. @@ -536,6 +577,30 @@ class ComputeUnitCore : public simdojo::CompositeComponent { std::shared_ptr plugin_group_ = ExecutionPluginGroup::empty_group(); + struct FilteredVgprHooks { + uint32_t physical_reg = 0; + uint64_t initial_reads = 0; + uint64_t writes = 0; + uint64_t pending_writes = 0; + std::array byte_masks{}; + }; + struct VgprWriteHookFilterState { + const Wavefront *wf = nullptr; + uint64_t lane_mask = 0; + uint64_t initial_destination_merge_read_lanes = 0; + uint8_t byte_mask = 0xF; + bool suppress_destination_merge_reads = false; + std::vector destinations; + + bool active() const { return wf != nullptr; } + void clear() { + wf = nullptr; + initial_destination_merge_read_lanes = 0; + destinations.clear(); + } + }; + mutable VgprWriteHookFilterState vgpr_write_hook_filter_; + /// Reverse lookup: physical SGPR index -> owning wavefront (for race detection). /// Populated at dispatch_wf time. Null entries mean "not allocated". std::vector sgpr_to_wave_; @@ -632,9 +697,15 @@ class IsaExecComputeUnit : public ExecComputeUnit { /// @returns Lane value from the VGPR file. uint32_t read_vgpr(uint32_t reg_idx, uint32_t lane) const override { - if (auto *wf = vgpr_to_wave_[reg_idx]) { - this->plugin_group_->onAmdgpuReadVgprs(wf, reg_idx, lane, lane + 1); - } + return read_vgpr_masked(reg_idx, lane, 0xF); + } + + uint32_t read_vgpr_masked(uint32_t reg_idx, uint32_t lane, uint8_t byte_mask) const override { + this->notify_vgpr_read(vgpr_to_wave_[reg_idx], reg_idx, lane, lane + 1, byte_mask); + return vgpr_file_[reg_idx][lane]; + } + + uint32_t read_vgpr_raw(uint32_t reg_idx, uint32_t lane) const override { return vgpr_file_[reg_idx][lane]; } @@ -644,6 +715,16 @@ class IsaExecComputeUnit : public ExecComputeUnit { /// @brief Write a value to the VGPR file. void write_vgpr(uint32_t reg_idx, uint32_t lane, uint32_t val) override { + write_vgpr_masked(reg_idx, lane, val, 0xF); + } + + void write_vgpr_masked(uint32_t reg_idx, uint32_t lane, uint32_t val, + uint8_t byte_mask) override { + this->notify_vgpr_write(vgpr_to_wave_[reg_idx], reg_idx, lane, byte_mask); + write_vgpr_raw(reg_idx, lane, val); + } + + void write_vgpr_raw(uint32_t reg_idx, uint32_t lane, uint32_t val) override { vgpr_file_[reg_idx][lane] = val; } 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 f085a2125846..f4ee7146ad4d 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 @@ -157,10 +157,10 @@ MemoryAccessCompletion vector_complete(VectorMemState &d, Wavefront &wf, 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 = cu.read_vgpr_raw(d.dst_reg_base + i, lane); val = d.d16_hi ? (old & 0x0000FFFFu) : (old & 0xFFFF0000u); } - cu.write_vgpr(d.dst_reg_base + i, lane, val); + cu.write_vgpr_from_memory(d.dst_reg_base + i, lane, val); } } } @@ -186,14 +186,14 @@ MemoryAccessCompletion vector_complete(VectorMemState &d, Wavefront &wf, else val = val & 0xFFFF; } else { - uint32_t old = cu.read_vgpr(d.dst_reg_base + i, lane); + uint32_t old = cu.read_vgpr_raw(d.dst_reg_base + i, lane); if (d.d16_hi) val = (old & 0xFFFF) | (val << 16); else val = (old & 0xFFFF0000) | (val & 0xFFFF); } } - cu.write_vgpr(d.dst_reg_base + i, lane, val); + cu.write_vgpr_from_memory(d.dst_reg_base + i, lane, val); } } return MemoryAccessCompletion::Complete; @@ -615,7 +615,7 @@ MemoryAccessCompletion LocalMemPipeline::complete_access(Instruction &inst, Wave uint32_t val = 0; uint32_t data_offset = lane * d.elem_size + i * 4; std::memcpy(&val, &d.ds2_response_data[data_offset], std::min(d.elem_size, 4u)); - cu.write_vgpr(d.ds2_dst_reg_base + i, lane, val); + cu.write_vgpr_from_memory(d.ds2_dst_reg_base + i, lane, val); } } // Per-lane DS read2 complete trace: show first 4 lanes with both accesses. diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/execution_plugin.h b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/execution_plugin.h index d6b087e73c34..8aa3624d6cd3 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/execution_plugin.h +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/execution_plugin.h @@ -114,6 +114,14 @@ class ExecutionPlugin { uint32_t /*lane_begin*/, uint32_t /*lane_end*/, uint8_t /*byte_mask*/ = kFullByteMask) {} + /// Called when a VGPR is written by instruction execution. + /// @param wf Owning wavefront, or nullptr if the register is unallocated. + /// @param physical_reg Physical register index in the VGPR file. + /// @param lane Lane index within the wavefront. + /// @param byte_mask Sub-dword byte mask (kFullByteMask = full dword). + virtual void onAmdgpuWriteVgpr(const amdgpu::Wavefront * /*wf*/, uint32_t /*physical_reg*/, + uint32_t /*lane*/, uint8_t /*byte_mask*/ = kFullByteMask) {} + /// Called when an SGPR is read during instruction execution. /// @param wf Owning wavefront, or nullptr if the register is unallocated. /// @param physical_reg Physical register index in the SGPR file. diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/execution_plugin_group.h b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/execution_plugin_group.h index 26ad0de77f74..c4a0ce0681c8 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/execution_plugin_group.h +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/execution_plugin_group.h @@ -132,6 +132,12 @@ class ExecutionPluginGroup { p->onAmdgpuReadVgprs(wf, physical_reg, lane_begin, lane_end, byte_mask); } + virtual void onAmdgpuWriteVgpr(const amdgpu::Wavefront *wf, uint32_t physical_reg, uint32_t lane, + uint8_t byte_mask = ExecutionPlugin::kFullByteMask) { + for (auto &p : plugins_) + p->onAmdgpuWriteVgpr(wf, physical_reg, lane, byte_mask); + } + virtual void onAmdgpuReadSgpr(const amdgpu::Wavefront *wf, uint32_t physical_reg) { for (auto &p : plugins_) p->onAmdgpuReadSgpr(wf, physical_reg); diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/profiled_execution_plugin_group.h b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/profiled_execution_plugin_group.h index 1fbfe73c6247..1796ed801e55 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/profiled_execution_plugin_group.h +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/profiled_execution_plugin_group.h @@ -49,7 +49,7 @@ class ProfiledExecutionPluginGroup : public ExecutionPluginGroup { } void onShutdown() { - if (prof_before_exec_.count > 0) + if (has_profile_data()) print_profile_summary(); ExecutionPluginGroup::onShutdown(); double total = std::chrono::duration(Clock::now() - start_time_).count(); @@ -76,7 +76,7 @@ class ProfiledExecutionPluginGroup : public ExecutionPluginGroup { } void onAmdgpuDispatchPacketProcessed(const KernelDispatchInfo &info) override { - if (prof_before_exec_.count > 0) + if (has_profile_data()) print_profile_summary(); current_kernel_name_ = info.kernel_name; dispatch_kernel_names_[info.dispatch_id] = info.kernel_name; @@ -85,7 +85,7 @@ class ProfiledExecutionPluginGroup : public ExecutionPluginGroup { } void onAmdgpuDispatchExecutionEnd(uint32_t dispatch_id) override { - if (prof_before_exec_.count > 0) { + if (has_profile_data()) { auto it = dispatch_kernel_names_.find(dispatch_id); if (it != dispatch_kernel_names_.end()) current_kernel_name_ = it->second; @@ -128,6 +128,13 @@ class ProfiledExecutionPluginGroup : public ExecutionPluginGroup { }); } + void onAmdgpuWriteVgpr(const amdgpu::Wavefront *wf, uint32_t physical_reg, uint32_t lane, + uint8_t byte_mask = ExecutionPlugin::kFullByteMask) override { + profiled_dispatch(prof_write_vgpr_, [&]() { + ExecutionPluginGroup::onAmdgpuWriteVgpr(wf, physical_reg, lane, byte_mask); + }); + } + void onAmdgpuReadSgpr(const amdgpu::Wavefront *wf, uint32_t physical_reg) override { profiled_dispatch(prof_read_sgpr_, [&]() { ExecutionPluginGroup::onAmdgpuReadSgpr(wf, physical_reg); }); @@ -154,10 +161,19 @@ class ProfiledExecutionPluginGroup : public ExecutionPluginGroup { } } + bool has_profile_data() const { + return prof_before_exec_.count > 0 || prof_after_exec_.count > 0 || prof_read_vgpr_.count > 0 || + prof_write_vgpr_.count > 0 || prof_read_sgpr_.count > 0 || prof_route_mem_.count > 0 || + prof_barrier_.count > 0 || prof_wg_dispatched_.count > 0 || + prof_wg_completed_.count > 0 || prof_wf_dispatched_.count > 0 || + prof_wf_halted_.count > 0; + } + void reset_profiles() { prof_before_exec_ = {}; prof_after_exec_ = {}; prof_read_vgpr_ = {}; + prof_write_vgpr_ = {}; prof_read_sgpr_ = {}; prof_route_mem_ = {}; prof_barrier_ = {}; @@ -179,6 +195,7 @@ class ProfiledExecutionPluginGroup : public ExecutionPluginGroup { print_hook("beforeExecuteInstruction", prof_before_exec_); print_hook("afterExecuteInstruction", prof_after_exec_); print_hook("readVgpr", prof_read_vgpr_); + print_hook("writeVgpr", prof_write_vgpr_); print_hook("readSgpr", prof_read_sgpr_); print_hook("routeMemoryInstruction", prof_route_mem_); print_hook("barrierResolved", prof_barrier_); @@ -199,6 +216,7 @@ class ProfiledExecutionPluginGroup : public ExecutionPluginGroup { HookProfile prof_before_exec_; HookProfile prof_after_exec_; HookProfile prof_read_vgpr_; + HookProfile prof_write_vgpr_; HookProfile prof_read_sgpr_; HookProfile prof_route_mem_; HookProfile prof_barrier_; diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/core/race_detector.h b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/core/race_detector.h index 7b7d13093d4c..10ce45cbeb3c 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/core/race_detector.h +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/core/race_detector.h @@ -53,8 +53,7 @@ class RaceDetector { void validateRead(int addr, WaveId, int lane, int nBytes) const; /// Check for WAR hazards: no outstanding LDS reads overlap the range. - /// TODO(newling): WAW detection (write vs outstanding writes) is not - /// implemented. + /// LDS WAW detection (write vs outstanding writes) is not implemented. void validateWrite(int addr, WaveId, int lane, int nBytes) const; const EventRegistry &events() const { return events_; } diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/core/wave_race_state.cpp b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/core/wave_race_state.cpp index 086a370face0..dd364d7c40c2 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/core/wave_race_state.cpp +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/core/wave_race_state.cpp @@ -186,6 +186,17 @@ void WaveRaceState::checkVgprRead(int reg, int lane, uint8_t byteMask) const { } } +void WaveRaceState::checkVgprWrite(int reg, int lane, uint8_t byteMask) const { + for (EventId eid : vgprMemoryEvents[reg]) { + if (isToVgpr(detector->events().type(eid)) && + (detector->events().byteMask(eid) & byteMask) != 0 && + detector->events().isActiveForLane(eid, lane)) { + detector->getRaceHandler()( + {RaceViolation::Space::VGPR, reg, waveId.value, lane, true, detector->getWorkgroupId()}); + } + } +} + // Like checkVgprRead but for instructions that read all lanes (e.g. cross-lane ops). // countr_zero picks the first active lane from the event's exec mask as the // representative lane for the violation report. diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/core/wave_race_state.h b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/core/wave_race_state.h index d01b7d18c486..66d95c12a07f 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/core/wave_race_state.h +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/core/wave_race_state.h @@ -69,6 +69,9 @@ class WaveRaceState { /// Check a full VGPR read for races. Calls the RaceHandler on violation. void checkVgprRead(int reg, int lane, uint8_t byteMask) const; + /// Check a VGPR write for write-after-load races. + void checkVgprWrite(int reg, int lane, uint8_t byteMask) const; + /// Check all lanes of a VGPR for races (used by bulk register reads). void checkVgprReadAllLanes(int reg) const; diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/plugin.cpp b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/plugin.cpp index 72cfca952349..077465e1afd4 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/plugin.cpp +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/plugin.cpp @@ -229,9 +229,10 @@ void RaceDetectorPlugin::onAmdgpuWorkgroupDispatched(uint32_t dispatch_id, uint3 const char *space = v.space == RaceViolation::Space::VGPR ? "VGPR" : v.space == RaceViolation::Space::SGPR ? "SGPR" : "LDS"; - sink().write(std::format( - "RACE type={} reg={} wave={} lane={} wg={},{},{} conflict=unknown\n{}END_RACE\n", space, - v.index, v.wave, v.lane, v.workgroupId.x, v.workgroupId.y, v.workgroupId.z, oss.str())); + sink().write(std::format("RACE type={} access={} reg={} wave={} lane={} wg={},{},{} " + "conflict=unknown\n{}END_RACE\n", + space, v.isWrite ? "write" : "read", v.index, v.wave, v.lane, + v.workgroupId.x, v.workgroupId.y, v.workgroupId.z, oss.str())); } } }; @@ -346,6 +347,15 @@ void RaceDetectorPlugin::onAmdgpuReadVgprs(const amdgpu::Wavefront *wf, uint32_t byte_mask); } +void RaceDetectorPlugin::onAmdgpuWriteVgpr(const amdgpu::Wavefront *wf, uint32_t physical_reg, + uint32_t lane, uint8_t byte_mask) { + auto *s = get_state(wf); + if (!s || !s->race_state) + return; + uint32_t logical_reg = physical_reg - wf->vgpr_alloc().base; + s->race_state->checkVgprWrite(static_cast(logical_reg), static_cast(lane), byte_mask); +} + void RaceDetectorPlugin::onAmdgpuReadSgpr(const amdgpu::Wavefront *wf, uint32_t physical_reg) { auto *s = get_state(wf); assert(s && s->race_state); diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/plugin.h b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/plugin.h index 6150965cf44f..a8026b188abf 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/plugin.h +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/plugin.h @@ -104,6 +104,9 @@ class RaceDetectorPlugin : public ExecutionPlugin { void onAmdgpuReadVgprs(const amdgpu::Wavefront *wf, uint32_t physical_reg, uint32_t lane_begin, uint32_t lane_end, uint8_t byte_mask = 0xF) override; + void onAmdgpuWriteVgpr(const amdgpu::Wavefront *wf, uint32_t physical_reg, uint32_t lane, + uint8_t byte_mask = 0xF) override; + void onAmdgpuReadSgpr(const amdgpu::Wavefront *wf, uint32_t physical_reg) override; void onAmdgpuBeforeExecuteInstruction(uint64_t pc, const Instruction &inst, diff --git a/emulation/rocjitsu/tests/execution_plugin_test.cpp b/emulation/rocjitsu/tests/execution_plugin_test.cpp index 96b8d6d89a4c..20e64153bd59 100644 --- a/emulation/rocjitsu/tests/execution_plugin_test.cpp +++ b/emulation/rocjitsu/tests/execution_plugin_test.cpp @@ -8,6 +8,9 @@ #include "embedded_schema.h" #include "rocjitsu/config/config_loader.h" +#include "rocjitsu/isa/arch/amdgpu/cdna4/machine_insts.h" +#include "rocjitsu/isa/arch/amdgpu/rdna4/machine_insts.h" +#include "rocjitsu/isa/arch/amdgpu/shared/dpp_sdwa_ops.h" #include "rocjitsu/isa/instruction.h" #include "rocjitsu/vm/soc.h" @@ -17,12 +20,16 @@ RJ_DIAGNOSTIC_IGNORE_PEDANTIC #include "hsa/AMDHSAKernelDescriptor.h" RJ_DIAGNOSTIC_POP +#include "rocjitsu/vm/plugins/profiled_execution_plugin_group.h" #include "rocjitsu/vm/plugins/race_detector/plugin.h" +#include "util/simd_test_hooks.h" #include #include +#include #include +#include #include #include #include @@ -42,6 +49,85 @@ constexpr uint32_t S_NOP = sopp(0); constexpr uint32_t S_ENDPGM = sopp(1); constexpr uint32_t S_BARRIER = sopp(10); +// VOP1 encoding: bits[31:25]=0x3F, bits[24:17]=vdst, +// bits[16:9]=op, bits[8:0]=src0. +constexpr uint32_t vop1(uint32_t op, uint32_t vdst, uint32_t src0) { + return (0x3Fu << 25) | (vdst << 17) | (op << 9) | src0; +} +constexpr uint32_t vgpr_src(uint32_t idx) { return 256 + idx; } +constexpr uint32_t v_mov_b32(uint32_t vdst, uint32_t src0) { return vop1(1, vdst, src0); } + +std::array v_mov_b32_dpp(uint32_t vdst, uint32_t vsrc0) { + cdna4::Vop1VopDppMachineInst raw{}; + raw.src0 = amdgpu::SRC_DPP; + raw.op = 1; + raw.vdst = vdst; + raw.encoding = 0x3F; + raw.vsrc0 = vsrc0; + raw.dpp_ctrl = amdgpu::dpp::ROW_BCAST15; + raw.bound_ctrl = 0; + raw.bank_mask = 0xF; + raw.row_mask = 0xF; + static_assert(sizeof(raw) == 2 * sizeof(uint32_t)); + std::array code{}; + std::memcpy(code.data(), &raw, sizeof(raw)); + code[2] = S_ENDPGM; + return code; +} + +std::array rdna4_v_add_f16_dpp_high(uint32_t vdst, uint32_t vsrc0, uint32_t vsrc1) { + rdna4::Vop3VopDpp16MachineInst raw{}; + raw.vdst = vdst; + raw.opsel = 0xB; + raw.op = 0x132; + raw.encoding = 0x35; + raw.src0 = amdgpu::SRC_DPP; + raw.src1 = vgpr_src(vsrc1); + raw.vsrc0 = vsrc0; + raw.dpp_ctrl = amdgpu::dpp::ROW_SHR1; + raw.fi = 1; + raw.bound_ctrl = 1; + raw.bank_mask = 0xF; + raw.row_mask = 0xF; + static_assert(sizeof(raw) == 3 * sizeof(uint32_t)); + std::array code{}; + std::memcpy(code.data(), &raw, sizeof(raw)); + code[3] = S_ENDPGM; + return code; +} + +std::array v_mov_b32_sdwa_byte1_preserve(uint32_t vdst, uint32_t vsrc0) { + cdna4::Vop1VopSdwaMachineInst raw{}; + raw.src0 = amdgpu::SRC_SDWA; + raw.op = 1; + raw.vdst = vdst; + raw.encoding = 0x3F; + raw.vsrc0 = vsrc0; + raw.dst_sel = amdgpu::sdwa::BYTE_1; + raw.dst_unused = amdgpu::sdwa::UNUSED_PRESERVE; + raw.src0_sel = amdgpu::sdwa::DWORD; + static_assert(sizeof(raw) == 2 * sizeof(uint32_t)); + std::array code{}; + std::memcpy(code.data(), &raw, sizeof(raw)); + code[2] = S_ENDPGM; + return code; +} + +constexpr std::array cdna_vop3(uint32_t op, uint32_t vdst, uint32_t src0, + uint32_t src1, uint32_t src2, uint32_t opsel = 0) { + return {vdst | ((opsel & 0xFu) << 11) | ((op & 0x3FFu) << 16) | (0x34u << 26), + (src0 & 0x1FFu) | ((src1 & 0x1FFu) << 9) | ((src2 & 0x1FFu) << 18)}; +} + +struct ForceScalarGuard { + explicit ForceScalarGuard(bool force_scalar) : old_force_scalar(util::force_scalar()) { + util::set_force_scalar_for_testing(force_scalar); + } + ~ForceScalarGuard() { util::set_force_scalar_for_testing(old_force_scalar); } + + bool old_force_scalar; +}; + struct HookEvent { enum Kind { DISPATCH_PACKET_PROCESSED, @@ -55,6 +141,7 @@ struct HookEvent { AFTER_INSTRUCTION, ROUTE_MEMORY, READ_VGPR, + WRITE_VGPR, READ_SGPR, BARRIER_RESOLVED, INIT, @@ -71,6 +158,9 @@ struct HookEvent { uint32_t physical_vgpr_count = 0; uint32_t sgpr_count = 0; uint64_t pc = 0; + uint32_t physical_reg = 0; + uint32_t lane = 0; + uint8_t byte_mask = 0; std::string mnemonic; }; @@ -121,6 +211,8 @@ class OrderingPlugin : public ExecutionPlugin { } void onAmdgpuWavefrontDispatched(amdgpu::Wavefront &wf) override { + if (dispatch_exec_mask) + wf.set_exec(*dispatch_exec_mask); HookEvent e{HookEvent::WAVEFRONT_DISPATCHED}; e.dispatch_id = wf.dispatch_id(); e.wg_id = wf.wg_id(); @@ -128,6 +220,8 @@ class OrderingPlugin : public ExecutionPlugin { events.push_back(e); } + std::optional dispatch_exec_mask; + void onAmdgpuWavefrontHalted(amdgpu::Wavefront &wf) override { HookEvent e{HookEvent::WAVEFRONT_HALTED}; e.dispatch_id = wf.dispatch_id(); @@ -167,14 +261,30 @@ class OrderingPlugin : public ExecutionPlugin { events.push_back(e); } - void onAmdgpuReadVgprs(const amdgpu::Wavefront *wf, uint32_t, uint32_t, uint32_t, - uint8_t) override { + void onAmdgpuReadVgprs(const amdgpu::Wavefront *wf, uint32_t physical_reg, uint32_t lane_begin, + uint32_t, uint8_t) override { HookEvent e{HookEvent::READ_VGPR}; if (wf) { e.dispatch_id = wf->dispatch_id(); e.wg_id = wf->wg_id(); e.wf_id = wf->wf_id(); } + e.physical_reg = physical_reg; + e.lane = lane_begin; + events.push_back(e); + } + + void onAmdgpuWriteVgpr(const amdgpu::Wavefront *wf, uint32_t physical_reg, uint32_t lane, + uint8_t byte_mask) override { + HookEvent e{HookEvent::WRITE_VGPR}; + if (wf) { + e.dispatch_id = wf->dispatch_id(); + e.wg_id = wf->wg_id(); + e.wf_id = wf->wf_id(); + } + e.physical_reg = physical_reg; + e.lane = lane; + e.byte_mask = byte_mask; events.push_back(e); } @@ -211,6 +321,7 @@ const char *kindName(HookEvent::Kind k) { "AFTER_INSTRUCTION", "ROUTE_MEMORY", "READ_VGPR", + "WRITE_VGPR", "READ_SGPR", "BARRIER_RESOLVED", "INIT", @@ -380,11 +491,15 @@ struct PluginFixture { std::unique_ptr engine; SoC *soc = nullptr; amdgpu::GpuMemory *mem = nullptr; + uint32_t wave_front_size = 64; - explicit PluginFixture(uint32_t num_wf_slots = 10) { + explicit PluginFixture(uint32_t num_wf_slots = 10, std::string_view arch = "cdna4", + uint32_t configured_wave_front_size = 64) + : wave_front_size(configured_wave_front_size) { + uint32_t sgprs_per_wf = arch == "rdna4" || arch == "gfx1250" ? 128 : 104; std::string json = std::format(R"({{ "max_ticks":10000,"num_threads":1,"exec_mode":"functional", - "vm":{{"arch":"cdna4","gpu":{{"device":{{"wave_front_size":64}}}}}}, + "vm":{{"arch":"{}","gpu":{{"device":{{"wave_front_size":{}}}}}}}, "topology":{{"root":{{"name":"soc","type":"soc","children":[ {{"name":"vram","type":"gpu_memory"}}, {{"name":"xcd0","type":"xcd","children":[ @@ -393,7 +508,7 @@ struct PluginFixture { {{"name":"se0","type":"shader_engine","children":[ {{"name":"cu[0:1]","type":"compute_unit","config":[ {{"key":"num_wf_slots","value":"{}"}}, - {{"key":"sgprs_per_wf","value":"104"}}, + {{"key":"sgprs_per_wf","value":"{}"}}, {{"key":"vgprs_per_wf","value":"256"}}, {{"key":"lds_size_kb","value":"64"}} ]}} @@ -404,7 +519,7 @@ struct PluginFixture { {{"src":"xcd0.se0.cu0.req","dst":"xcd0.l2.cpl_0","latency":1,"weight":10}} ]}}}} )", - num_wf_slots); + arch, wave_front_size, num_wf_slots, sgprs_per_wf); auto loaded = config::load_config_from_string(json, rocjitsu::kEmbeddedSchema); soc = loaded.soc(); mem = loaded.memory(); @@ -417,13 +532,18 @@ struct PluginFixture { amdgpu::ComputeUnitCore *cu() { return soc->xcd(0)->shader_engine(0)->compute_unit(0); } amdgpu::CommandProcessor *cp() { return soc->xcd(0)->command_processor(); } - uint64_t write_kernel(uint64_t addr, const uint32_t *code, size_t num_words) { + uint64_t write_kernel(uint64_t addr, const uint32_t *code, size_t num_words, + uint32_t enable_vgpr_workitem_id = 0) { using namespace rocr::llvm::amdhsa; kernel_descriptor_t kd{}; kd.kernel_code_entry_byte_offset = sizeof(kernel_descriptor_t); AMDHSA_BITS_SET(kd.compute_pgm_rsrc1, COMPUTE_PGM_RSRC1_GRANULATED_WORKITEM_VGPR_COUNT, 31); AMDHSA_BITS_SET(kd.compute_pgm_rsrc1, COMPUTE_PGM_RSRC1_GRANULATED_WAVEFRONT_SGPR_COUNT, 12); AMDHSA_BITS_SET(kd.compute_pgm_rsrc2, COMPUTE_PGM_RSRC2_USER_SGPR_COUNT, 2); + AMDHSA_BITS_SET(kd.compute_pgm_rsrc2, COMPUTE_PGM_RSRC2_ENABLE_VGPR_WORKITEM_ID, + enable_vgpr_workitem_id); + if (wave_front_size == 32) + AMDHSA_BITS_SET(kd.kernel_code_properties, KERNEL_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32, 1); mem->load_image(reinterpret_cast(&kd), sizeof(kd), addr); mem->load_image(reinterpret_cast(code), num_words * 4, addr + sizeof(kernel_descriptor_t)); @@ -454,8 +574,8 @@ struct PluginFixture { } void run_kernel(const uint32_t *code, size_t num_words, uint32_t grid = 64, - uint32_t workgroup = 64) { - uint64_t ko = write_kernel(0x1000, code, num_words); + uint32_t workgroup = 64, uint32_t enable_vgpr_workitem_id = 0) { + uint64_t ko = write_kernel(0x1000, code, num_words, enable_vgpr_workitem_id); test::AqlQueue queue(mem, cp()); queue.dispatch(ko, grid, workgroup); run_until_idle(); @@ -468,6 +588,19 @@ TEST(ExecutionPluginTest, NoPluginNoCrash) { f.run_kernel(code, 2); } +TEST(ExecutionPluginTest, ProfiledGroupReportsVgprWriteHook) { + ProfiledExecutionPluginGroup group; + StringSink sink; + group.add_sink(&sink); + + group.onInit(); + group.onAmdgpuWriteVgpr(nullptr, /*physical_reg=*/7, /*lane=*/3, /*byte_mask=*/0xC); + group.onShutdown(); + + EXPECT_NE(sink.str().find("writeVgpr"), std::string::npos) << sink.str(); + EXPECT_NE(sink.str().find("calls=1"), std::string::npos) << sink.str(); +} + // -- Ordering tests ---------------------------------------------------------- // // These tests use functional mode (the PluginFixture default). Tests that @@ -509,6 +642,144 @@ TEST(HookOrderingTest, WorkgroupDispatchedReportsPhysicalVgprBlockSize) { EXPECT_EQ(it->sgpr_count, f.cu()->config().sgprs_per_wf); } +TEST(HookOrderingTest, WorkitemIdSetupVgprWritesAreSilent) { + PluginFixture f; + auto *p = f.attach_ordering_plugin(); + const uint32_t code[] = {S_ENDPGM}; + f.run_kernel(code, 1, /*grid=*/64, /*workgroup=*/64, + /*enable_vgpr_workitem_id=*/2); + f.shutdown(); + + EventLog log(p->events); + EXPECT_EQ(log.count(HookEvent::WAVEFRONT_DISPATCHED), 1u); + EXPECT_EQ(log.count(HookEvent::WRITE_VGPR), 0u) + << "work-item ID setup writes must use the raw VGPR path and stay invisible to " + "instruction write hooks"; +} + +TEST(HookOrderingTest, SimdVgprWritesReachPlugins) { + PluginFixture f; + auto *p = f.attach_ordering_plugin(); + const uint32_t code[] = {v_mov_b32(/*vdst=*/2, vgpr_src(/*idx=*/0)), S_ENDPGM}; + f.run_kernel(code, 2); + f.shutdown(); + + EventLog log(p->events); + EXPECT_EQ(log.count(HookEvent::WRITE_VGPR), 64u); + for (const auto &event : p->events) { + if (event.kind == HookEvent::WRITE_VGPR) + EXPECT_EQ(event.byte_mask, 0xF); + } +} + +TEST(HookOrderingTest, DppReportsOnlyArchitecturalDestinationWrites) { + PluginFixture f; + auto *p = f.attach_ordering_plugin(); + auto code = v_mov_b32_dpp(/*vdst=*/2, /*vsrc0=*/0); + f.run_kernel(code.data(), code.size()); + f.shutdown(); + + EventLog log(p->events); + EXPECT_EQ(log.count(HookEvent::WRITE_VGPR), 48u); + uint32_t destination_reg = UINT32_MAX; + for (const auto &event : p->events) { + if (event.kind != HookEvent::WRITE_VGPR) + continue; + destination_reg = event.physical_reg; + EXPECT_GE(event.lane, 16u); + EXPECT_EQ(event.byte_mask, 0xF); + } + ASSERT_NE(destination_reg, UINT32_MAX); + for (const auto &event : p->events) + if (event.kind == HookEvent::READ_VGPR) + EXPECT_NE(event.physical_reg, destination_reg); +} + +TEST(HookOrderingTest, DppAliasedInactiveSourceReadRemainsVisible) { + PluginFixture f; + auto *p = f.attach_ordering_plugin(); + p->dispatch_exec_mask = ~(1ULL << 15); + auto code = v_mov_b32_dpp(/*vdst=*/0, /*vsrc0=*/0); + f.run_kernel(code.data(), code.size()); + f.shutdown(); + + uint32_t destination_reg = UINT32_MAX; + for (const auto &event : p->events) + if (event.kind == HookEvent::WRITE_VGPR) { + destination_reg = event.physical_reg; + break; + } + ASSERT_NE(destination_reg, UINT32_MAX); + EXPECT_TRUE(std::ranges::any_of(p->events, [destination_reg](const HookEvent &event) { + return event.kind == HookEvent::READ_VGPR && event.physical_reg == destination_reg && + event.lane == 15; + })); +} + +TEST(HookOrderingTest, DppTrue16PreservesDestinationByteMask) { + PluginFixture f(/*num_wf_slots=*/10, /*arch=*/"rdna4", /*wave_front_size=*/32); + auto *p = f.attach_ordering_plugin(); + auto code = rdna4_v_add_f16_dpp_high(/*vdst=*/2, /*vsrc0=*/0, /*vsrc1=*/1); + f.run_kernel(code.data(), code.size(), /*grid=*/32, /*workgroup=*/32); + f.shutdown(); + + EventLog log(p->events); + EXPECT_EQ(log.count(HookEvent::WRITE_VGPR), 32u); + for (const auto &event : p->events) + if (event.kind == HookEvent::WRITE_VGPR) + EXPECT_EQ(event.byte_mask, 0xC); +} + +TEST(HookOrderingTest, SdwaCoalescesMergeAndReportsSelectedBytes) { + PluginFixture f; + auto *p = f.attach_ordering_plugin(); + auto code = v_mov_b32_sdwa_byte1_preserve(/*vdst=*/2, /*vsrc0=*/0); + f.run_kernel(code.data(), code.size()); + f.shutdown(); + + EventLog log(p->events); + EXPECT_EQ(log.count(HookEvent::WRITE_VGPR), 64u); + uint32_t destination_reg = UINT32_MAX; + for (const auto &event : p->events) { + if (event.kind != HookEvent::WRITE_VGPR) + continue; + destination_reg = event.physical_reg; + EXPECT_EQ(event.byte_mask, 0x2); + } + ASSERT_NE(destination_reg, UINT32_MAX); + for (const auto &event : p->events) + if (event.kind == HookEvent::READ_VGPR) + EXPECT_NE(event.physical_reg, destination_reg); +} + +TEST(HookOrderingTest, True16VgprWritesReportSelectedByteMask) { + for (bool force_scalar : {false, true}) { + SCOPED_TRACE(force_scalar ? "scalar" : "simd"); + ForceScalarGuard guard(force_scalar); + PluginFixture f; + auto *p = f.attach_ordering_plugin(); + constexpr auto inst = cdna_vop3(/*op=*/0x12B, /*vdst=*/16, /*src0=*/256, + /*src1=*/257, /*src2=*/0, /*opsel=*/0xB); + const uint32_t code[] = {inst[0], inst[1], S_ENDPGM}; + f.run_kernel(code, 3); + f.shutdown(); + + EventLog log(p->events); + EXPECT_EQ(log.count(HookEvent::WRITE_VGPR), 64u); + uint32_t destination_reg = UINT32_MAX; + for (const auto &event : p->events) { + if (event.kind == HookEvent::WRITE_VGPR) { + destination_reg = event.physical_reg; + EXPECT_EQ(event.byte_mask, 0xC); + } + } + ASSERT_NE(destination_reg, UINT32_MAX); + for (const auto &event : p->events) + if (event.kind == HookEvent::READ_VGPR) + EXPECT_NE(event.physical_reg, destination_reg); + } +} + TEST(HookOrderingTest, FiveDispatchLifecycle) { PluginFixture f(/*num_wf_slots=*/1); auto *p = f.attach_ordering_plugin(); diff --git a/emulation/rocjitsu/tests/race-detector/CMakeLists.txt b/emulation/rocjitsu/tests/race-detector/CMakeLists.txt index 6947a9a95199..a95dc38332ab 100644 --- a/emulation/rocjitsu/tests/race-detector/CMakeLists.txt +++ b/emulation/rocjitsu/tests/race-detector/CMakeLists.txt @@ -86,6 +86,7 @@ rj_add_gfx950_race_test_case(dual_offset_lds) rj_add_gfx950_race_test_case(dual_offset_lds_race) rj_add_gfx950_race_test_case(scratch_vmcnt) rj_add_gfx950_race_test_case(scratch_vmcnt_race) +rj_add_gfx950_race_test_case(waw_global_load_then_alu) rj_add_gfx950_race_test_case(f64_vgpr_safe) rj_add_gfx950_race_test_case(f64_vgpr_race) diff --git a/emulation/rocjitsu/tests/race-detector/hip_race_gfx950_test.hip b/emulation/rocjitsu/tests/race-detector/hip_race_gfx950_test.hip index af03de4462e6..906c35c0a4e7 100644 --- a/emulation/rocjitsu/tests/race-detector/hip_race_gfx950_test.hip +++ b/emulation/rocjitsu/tests/race-detector/hip_race_gfx950_test.hip @@ -34,6 +34,7 @@ struct RaceRecord { std::string type; // "VGPR", "SGPR", "LDS" + std::string access; // "read" or "write" int reg = -1; // register index (VGPR/SGPR) or byte address (LDS) int wave = -1; int lane = -1; @@ -47,7 +48,7 @@ struct RaceRecord { /// /// Example log file content: /// -/// RACE type=VGPR reg=2 wave=0 lane=0 wg=0,0,0 conflict=global_load_dword +/// RACE type=VGPR access=read reg=2 wave=0 lane=0 wg=0,0,0 conflict=global_load_dword /// Race on VGPR v2 [workgroup (0, 0, 0), wave 0, lane 0] /// ==> 0x5400021940 global_load_dword v2, v[2:3], off /// 0x5400021948 s_nop 0 @@ -81,6 +82,7 @@ static std::vector parse_race_log() { std::string key = rest.substr(pos, eq - pos); std::string val = rest.substr(eq + 1, sp - eq - 1); if (key == "type") r.type = val; + else if (key == "access") r.access = val; else if (key == "reg") r.reg = std::stoi(val); else if (key == "wave") r.wave = val.empty() ? -1 : std::stoi(val); else if (key == "lane") r.lane = val.empty() ? -1 : std::stoi(val); @@ -1062,26 +1064,19 @@ TEST_F(Gfx950RaceTest, scratch_vmcnt_race) { EXPECT_NE(t.mark0.find("global_load_dword"), npos) << "mark0: " << t.mark0; } -// TODO(newling): D16 (partial-register) HIP tests. +// TODO(newling): Add D16 (partial-register) HIP tests. // -// Concrete gap: the rocjitsu CU's read_vgpr always reads the full 32-bit -// register, so checkVgprRead is called with byteMask=0xF. Sub-dword store -// instructions (ds_write_b16, global_store_short) only consume 16 bits of -// the source VGPR, but trigger a full-register race check — false positive -// when the other half has a pending D16 load. +// The core and generated CU paths have byte-mask support. The remaining gap is +// HIP coverage proving sub-dword consumers such as ds_write_b16 and +// global_store_short read/write with byteMask=0x3/0xC and do not falsely race +// with the unrelated half of a pending D16 load. // -// Fix: add read_vgpr_masked(reg, lane, byteMask) to the CU, and have the -// codegen (_gen_ds_write in codegen.py, esz <= 2) emit the masked variant. -// See compute_unit.h for details. The original race-emulator solved this -// with Wave::getHalfVgpr(reg, lane, hi) which passed byteMask=0x3/0xC. +// read_vgpr_masked(reg, lane, byteMask) follows the original race-emulator +// Wave::getHalfVgpr(reg, lane, hi) model, which passed byteMask=0x3/0xC. // See /home/jnewling/workspace/rocm-libraries/shared/race-emulator/tests/ // race_tests.cpp (D16_LdsLoad_NoFalsePositive, D16_HiLoadOutstanding_LoReadSafe) // for the original tests exercising this. // -// The write-side byte mask propagation IS done: registerEvent/registerLdsEvent -// pass 0x3/0xC from VectorMemState::d16_lo/hi. The core D16 tests using -// RaceTestBuilder (which bypasses read_vgpr) pass correctly. -// // Deeper observation: even if an instruction reads "garbage" from pending // bytes, it only matters if that garbage propagates to an observable side // effect (global store, host-visible memory). Garbage churned through @@ -1092,13 +1087,10 @@ TEST_F(Gfx950RaceTest, scratch_vmcnt_race) { // --- WAW: global load then ALU write to same register --- // -// TODO(newling): the detector does not currently flag this. A global load targets -// a VGPR, then an ALU instruction overwrites the same VGPR before -// s_waitcnt retires the load. The final register value is ambiguous — -// if the load completes after the ALU write, it silently clobbers the -// ALU result. This is a WAW (write-after-write) race that should be -// detected but isn't, because the detector only checks reads against -// pending loads (checkVgprRead), not ALU writes. +// A global load targets a VGPR, then an ALU instruction overwrites the same +// VGPR before s_waitcnt retires the load. The final register value is +// ambiguous: if the load completes after the ALU write, it silently clobbers +// the ALU result. __global__ void waw_global_load_then_alu_kernel(const float *src, float *dst) { int tid = threadIdx.x + blockIdx.x * blockDim.x; @@ -1106,7 +1098,7 @@ __global__ void waw_global_load_then_alu_kernel(const float *src, float *dst) { asm volatile( "global_load_dword %0, %1, off\n" // BUG: no s_waitcnt before ALU write to the same register. - "v_add_f32 %0, %0, %0\n" + "v_mov_b32 %0, 0\n" : "=v"(val) : "v"(&src[tid]) ); asm volatile("s_waitcnt vmcnt(0)\n" ::: "memory"); @@ -1117,12 +1109,13 @@ TEST_F(Gfx950RaceTest, waw_global_load_then_alu) { auto *src = allocWithData(64), *dst = alloc(64); waw_global_load_then_alu_kernel<<<1, 64>>>(src, dst); sync(); - // TODO(newling): this should detect a VGPR WAW race, but currently - // doesn't. When WAW detection is implemented, change to: - // auto r = records(); - // ASSERT_GE(r.size(), 1u); - // EXPECT_NE(r[0].type.find("VGPR"), npos); - ExpectNoRace(); + auto r = records(); + ASSERT_GE(r.size(), 1u); + EXPECT_EQ(r[0].access, "write"); + auto t = parseTrace(r[0]); + EXPECT_NE(t.header.find("VGPR"), npos); + EXPECT_NE(t.mark0.find("global_load_dword"), npos) << "mark0: " << t.mark0; + EXPECT_NE(t.mark1.find("v_mov_b32"), npos) << "mark1: " << t.mark1; } // --- f64 VGPR race (64-bit global load + use without waitcnt) --- diff --git a/emulation/rocjitsu/tests/race-detector/race_detector_tests.cpp b/emulation/rocjitsu/tests/race-detector/race_detector_tests.cpp index c4a9cef4230a..67143e8fffcb 100644 --- a/emulation/rocjitsu/tests/race-detector/race_detector_tests.cpp +++ b/emulation/rocjitsu/tests/race-detector/race_detector_tests.cpp @@ -218,6 +218,23 @@ TEST(RaceDetector, SameWave_WriteWaitcntOk) { EXPECT_FALSE(b.hasRace()); } +TEST(RaceDetector, SameWave_VgprWriteBeforeWaitcntRaces) { + RaceTestBuilder b(/*numWaves=*/1, /*vgprs=*/8, /*sgprs=*/8); + b.globalLoad(/*wave=*/0, /*vgprBase=*/2, /*numRegs=*/1); + b.checkVgprWrite(/*wave=*/0, /*reg=*/2, /*lane=*/0); + EXPECT_TRUE(b.hasVgprRace(2)); + ASSERT_EQ(b.raceCount(), 1); + EXPECT_TRUE(b.violations()[0].isWrite); +} + +TEST(RaceDetector, SameWave_VgprWriteAfterWaitcntOk) { + RaceTestBuilder b(/*numWaves=*/1, /*vgprs=*/8, /*sgprs=*/8); + b.globalLoad(/*wave=*/0, /*vgprBase=*/2, /*numRegs=*/1); + b.waitcnt(/*wave=*/0, /*vmcnt=*/0); + b.checkVgprWrite(/*wave=*/0, /*reg=*/2, /*lane=*/0); + EXPECT_FALSE(b.hasRace()); +} + TEST(RaceDetector, SameWave_DoubleWriteRace) { // Two LDS loads into same VGPR, waitcnt(1) drains oldest → newest RACE. RaceTestBuilder b(/*numWaves=*/1, /*vgprs=*/8, /*sgprs=*/8); @@ -354,6 +371,25 @@ TEST(RaceDetector, D16_LoOutstanding_FullRaces) { EXPECT_TRUE(b.hasVgprRace(2)); } +TEST(RaceDetector, D16_HiOutstanding_LoWriteSafe) { + // Hi half outstanding, low-half destination write does not overlap. + RaceTestBuilder b(/*numWaves=*/1, /*vgprs=*/8, /*sgprs=*/8); + b.ldsRead(/*wave=*/0, /*lane=*/0, /*addr=*/0, /*bytes=*/2, /*vgprDst=*/2, + /*byteMask=*/0xC); + b.checkVgprWrite(/*wave=*/0, /*reg=*/2, /*lane=*/0, + /*byteMask=*/0x3); + EXPECT_FALSE(b.hasRace()); +} + +TEST(RaceDetector, D16_HiOutstanding_HiWriteRaces) { + RaceTestBuilder b(/*numWaves=*/1, /*vgprs=*/8, /*sgprs=*/8); + b.ldsRead(/*wave=*/0, /*lane=*/0, /*addr=*/0, /*bytes=*/2, /*vgprDst=*/2, + /*byteMask=*/0xC); + b.checkVgprWrite(/*wave=*/0, /*reg=*/2, /*lane=*/0, + /*byteMask=*/0xC); + EXPECT_TRUE(b.hasVgprRace(2)); +} + // ---- Direct-to-LDS (GlobalToLds) ---- TEST(RaceDetector, Dtl_CrossWaveRace) { @@ -478,16 +514,15 @@ TEST(RaceDetector, LdsSameWave_ScatterBroadcastOk) { } TEST(RaceDetector, LdsSameWave_AllLanesWriteSameAddr) { - // All lanes write same address (WAW). Currently not detected as a race. - // TODO(newling): WAW detection is not implemented — this documents current - // behavior. + // All lanes write same LDS address (LDS WAW). Currently not detected as a + // race; this documents current LDS behavior. RaceTestBuilder b(/*numWaves=*/1, /*vgprs=*/8, /*sgprs=*/8); b.ldsWrite(/*wave=*/0, /*lane=*/0, /*addr=*/0, /*bytes=*/4); b.ldsWrite(/*wave=*/0, /*lane=*/1, /*addr=*/0, /*bytes=*/4); b.ldsWrite(/*wave=*/0, /*lane=*/2, /*addr=*/0, /*bytes=*/4); b.waitcnt(/*wave=*/0, /*vmcnt=*/-1, /*lgkmcnt=*/0); b.checkLdsRead(/*wave=*/0, /*lane=*/0, /*addr=*/0, /*bytes=*/4); - // Currently passes (WAW not flagged). If WAW detection is added, this + // Currently passes (LDS WAW not flagged). If LDS WAW detection is added, this // test should be updated to EXPECT_TRUE(b.hasRace()). EXPECT_FALSE(b.hasRace()); } diff --git a/emulation/rocjitsu/tests/race-detector/race_test_builder.h b/emulation/rocjitsu/tests/race-detector/race_test_builder.h index 1de0cb2203c9..e9069b5b4cac 100644 --- a/emulation/rocjitsu/tests/race-detector/race_test_builder.h +++ b/emulation/rocjitsu/tests/race-detector/race_test_builder.h @@ -130,6 +130,10 @@ class RaceTestBuilder { waves_[wave]->checkVgprRead(reg, lane, byteMask); } + void checkVgprWrite(int wave, int reg, int lane, uint8_t byteMask = 0xF) { + waves_[wave]->checkVgprWrite(reg, lane, byteMask); + } + void checkSgprRead(int wave, int reg) { waves_[wave]->checkSgprRead(reg); } void checkLdsRead(int wave, int lane, int addr, int bytes) { diff --git a/emulation/rocjitsu/tests/shared_infra_test.cpp b/emulation/rocjitsu/tests/shared_infra_test.cpp index 983e5aa501ba..7c2b9b8f4805 100644 --- a/emulation/rocjitsu/tests/shared_infra_test.cpp +++ b/emulation/rocjitsu/tests/shared_infra_test.cpp @@ -2753,6 +2753,17 @@ TEST(SdwaTest, DstMerge) { EXPECT_EQ(merged, 0x12345678u); } +TEST(SdwaTest, DstWriteMask) { + using namespace amdgpu::sdwa; + EXPECT_EQ(sdwa_dst_write_mask(BYTE_0, UNUSED_PRESERVE), 0x1); + EXPECT_EQ(sdwa_dst_write_mask(BYTE_3, UNUSED_PRESERVE), 0x8); + EXPECT_EQ(sdwa_dst_write_mask(WORD_0, UNUSED_PRESERVE), 0x3); + EXPECT_EQ(sdwa_dst_write_mask(WORD_1, UNUSED_PRESERVE), 0xC); + EXPECT_EQ(sdwa_dst_write_mask(DWORD, UNUSED_PRESERVE), 0xF); + EXPECT_EQ(sdwa_dst_write_mask(BYTE_1, UNUSED_PAD), 0xF); + EXPECT_EQ(sdwa_dst_write_mask(WORD_1, UNUSED_SEXT), 0xF); +} + // --------------------------------------------------------------------------- // Scratch address calculation tests // ---------------------------------------------------------------------------