From 0fcf829d2f07bfa360f14ef6b7be67efd5092c87 Mon Sep 17 00:00:00 2001 From: hanhanW Date: Sun, 23 Aug 2026 21:06:32 -0700 Subject: [PATCH 1/3] [rocjitsu] Add support for packed V_PK_LSHL_ADD_U64 execution. Teach the semantics derivation and packed code generator that `V_PK_LSHL_ADD_U64` operates on two U64 values and two U32 shift counts. Generate operand helpers that read two U64 elements from four consecutive VGPRs, broadcast the low U64 value for SGPR and non-register sources, read two U32 shift counts, and write two U64 results. Install the CDNA5 execution callback for this instruction. This fills the instruction's execution gap while following the public LLVM operand profile and `ThreeOpFrag` selection pattern. The instruction does not consume clamp or source modifiers. Public sources only define shift counts from 0 through 4; `lshl_masked` keeps unsupported host inputs well-defined but does not establish gfx1251 semantics outside that range. Add generator and instruction-level integration tests covering public source forms, SGPR broadcast behavior, wraparound, EXEC masking, operand order, register layout, and legal source/destination overlap. Keep LLVM's literal-101 vector as decode-only coverage because the available public sources do not define its arithmetic result. This enables only `V_PK_LSHL_ADD_U64`. The gfx1251 execution capability remains unavailable, normal simulator configuration still rejects gfx1251, and the tests do not validate behavior on physical gfx1251 hardware. - [x] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. Signed-off-by: hanhanW --- .../lib/python/amdisa/codegen/_generator.py | 50 +++++ .../python/amdisa/codegen/execute/__init__.py | 2 + .../python/amdisa/codegen/execute/packed.py | 32 +++ .../rocjitsu/lib/python/amdisa/semantics.py | 4 + .../tests/test_generator_profile_gates.py | 21 +- .../amdisa/tests/test_packed_codegen.py | 24 +++ .../python/amdisa/tests/test_sema_derive.py | 8 + .../generated/cdna5/execution_backend.h | 1 + .../cdna5/execution_backend_exec.cpp | 1 + .../isa/arch/amdgpu/generated/cdna5/vop3p.cpp | 3 +- .../isa/arch/amdgpu/generated/cdna5/vop3p.h | 1 + .../amdgpu/generated/cdna5/vop3p_exec.cpp | 66 +++++++ .../rocjitsu/tests/cdna5_execution_test.cpp | 187 +++++++++++++++++- .../isa/amdgpu_isa_cdna5_variants.json | 1 - 14 files changed, 393 insertions(+), 8 deletions(-) diff --git a/emulation/rocjitsu/lib/python/amdisa/codegen/_generator.py b/emulation/rocjitsu/lib/python/amdisa/codegen/_generator.py index 3516ed5e52e3..8917b0c9bfa6 100644 --- a/emulation/rocjitsu/lib/python/amdisa/codegen/_generator.py +++ b/emulation/rocjitsu/lib/python/amdisa/codegen/_generator.py @@ -5296,6 +5296,56 @@ def _emit_cdna5_matrix_fmt_helpers(self) -> _ImplOutputs: const auto pair = amdgpu::RegisterAccess(wf).read_lane_pair32(operand, lane); return {pair.lo, pair.hi}; } + + struct PkU64Pair { + uint64_t lo; + uint64_t hi; + }; + + struct PkU32Pair { + uint32_t lo; + uint32_t hi; + }; + + bool is_general_register(const std::optional ®) { + return reg && (reg->cls == RegClass::SGPR || reg->cls == RegClass::VGPR); + } + + Operand packed_register_dword_offset(const Operand &operand, uint32_t dword_offset) { + Operand shifted = operand; + shifted.encoding_value_ += static_cast(dword_offset); + return shifted; + } + + PkU64Pair read_pk_u64_pair(const Operand &operand, const amdgpu::Wavefront &wf, + uint32_t lane) { + const uint64_t lo = amdgpu::RegisterAccess(wf).read_lane64(operand, lane); + const auto reg = operand.to_register_ref(); + if (!reg || reg->cls != RegClass::VGPR) + return {lo, lo}; + + const Operand hi_operand = packed_register_dword_offset(operand, 2); + return {lo, amdgpu::RegisterAccess(wf).read_lane64(hi_operand, lane)}; + } + + PkU32Pair read_pk_u32_pair(const Operand &operand, const amdgpu::Wavefront &wf, + uint32_t lane) { + if (!is_general_register(operand.to_register_ref())) { + const uint32_t value = amdgpu::RegisterAccess(wf).read_lane(operand, lane); + return {value, value}; + } + + const uint64_t raw = amdgpu::RegisterAccess(wf).read_lane64(operand, lane); + return {static_cast(raw), static_cast(raw >> 32)}; + } + + void write_pk_u64_pair(const Operand &operand, amdgpu::Wavefront &wf, uint32_t lane, + PkU64Pair value) { + const Operand hi_operand = packed_register_dword_offset(operand, 2); + amdgpu::RegisterAccess access(wf); + access.write_lane64(operand, lane, value.lo); + access.write_lane64(hi_operand, lane, value.hi); + } ''') model = ( 'namespace {\n\n' diff --git a/emulation/rocjitsu/lib/python/amdisa/codegen/execute/__init__.py b/emulation/rocjitsu/lib/python/amdisa/codegen/execute/__init__.py index 12f0480b9246..82f84476971b 100644 --- a/emulation/rocjitsu/lib/python/amdisa/codegen/execute/__init__.py +++ b/emulation/rocjitsu/lib/python/amdisa/codegen/execute/__init__.py @@ -119,6 +119,7 @@ def _register_handlers() -> None: gen_pk_fmac_vop3, gen_pk_binop_f32, gen_pk_ternary_f32, + gen_pk_lshl_add_u64, gen_pk_mov_b32, gen_mad_mix_f32, gen_mad_mix_lo_hi, @@ -329,6 +330,7 @@ def _register_handlers() -> None: opsel_exprs=c.opsel_exprs, use_cdna5_helpers=c.arch_name == 'cdna5', ) + DISPATCH['pk_lshl_add_u64'] = lambda c: gen_pk_lshl_add_u64(c.dst_ops, c.src_ops) DISPATCH['pk_mov_b32'] = lambda c: gen_pk_mov_b32( c.dst_ops, c.src_ops, diff --git a/emulation/rocjitsu/lib/python/amdisa/codegen/execute/packed.py b/emulation/rocjitsu/lib/python/amdisa/codegen/execute/packed.py index 7e944b16902c..a131b13c6015 100644 --- a/emulation/rocjitsu/lib/python/amdisa/codegen/execute/packed.py +++ b/emulation/rocjitsu/lib/python/amdisa/codegen/execute/packed.py @@ -638,6 +638,38 @@ def gen_pk_ternary_f32( return '\n'.join(L) +def gen_pk_lshl_add_u64(dst: list[str], src: list[str]) -> str: + """Generate packed U64 shift-left-add over two independent elements. + + V_PK_LSHL_ADD_U64 has two 64-bit values in each 128-bit value/addend + operand and two 32-bit shift counts in its 64-bit shift operand. VGPR U64 + sources supply both elements, while SGPR and non-register U64 sources + broadcast the low 64 bits. Its public LLVM profile explicitly disables + clamp and source modifiers, so the body intentionally consumes neither + field. Public LLVM selection patterns and the public CDNA5 description + define shift counts only in the inclusive range 0..4. ``lshl_masked`` + keeps the host operation defined if an unsupported runtime value reaches + the callback; that fallback is not an architectural claim for counts above + 4. + """ + d, s0, s1, s2 = dst[0], src[0], src[1], src[2] + return '\n'.join( + [ + ' uint64_t exec = wf.exec();', + ' for (uint32_t lane = 0; lane < wf.wf_size(); ++lane) {', + ' if (!(exec & (1ULL << lane))) continue;', + f' const auto values = read_pk_u64_pair({s0}, wf, lane);', + f' const auto shifts = read_pk_u32_pair({s1}, wf, lane);', + f' const auto addends = read_pk_u64_pair({s2}, wf, lane);', + ' // Public semantics cover shift counts 0..4. The masked helper only avoids host undefined behavior for unsupported values.', + ' const uint64_t result_lo = amdgpu::lshl_masked(values.lo, static_cast(shifts.lo)) + addends.lo;', + ' const uint64_t result_hi = amdgpu::lshl_masked(values.hi, static_cast(shifts.hi)) + addends.hi;', + f' write_pk_u64_pair({d}, wf, lane, {{result_lo, result_hi}});', + ' }', + ] + ) + + def gen_pk_mov_b32( dst: list[str], src: list[str], diff --git a/emulation/rocjitsu/lib/python/amdisa/semantics.py b/emulation/rocjitsu/lib/python/amdisa/semantics.py index fe951179c1b1..fae9dcb3f263 100644 --- a/emulation/rocjitsu/lib/python/amdisa/semantics.py +++ b/emulation/rocjitsu/lib/python/amdisa/semantics.py @@ -1603,6 +1603,10 @@ def _derive_vop3p(name: str) -> InstructionSemantics | None: return InstructionSemantics( name, 'pk_binop_f32', operation='add', data_type='f32' ) + if name == 'V_PK_LSHL_ADD_U64': + return InstructionSemantics( + name, 'pk_lshl_add_u64', operation='lshl_add', data_type='u64' + ) if name == 'V_PK_MOV_B32': return InstructionSemantics(name, 'pk_mov_b32') 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 415d0c02b227..140e3cd8dd7d 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 @@ -5083,7 +5083,6 @@ def test_cdna5_model_only_variant_instructions_have_no_execution_callbacks( 'VPkSubNcU64Vop3p', 'VPkMaxNumF64Vop3p', 'VPkMinNumF64Vop3p', - 'VPkLshlAddU64Vop3p', 'VWmmaF6416x16x4F64Vop3p', ) header = (gfx1250_generated_root / 'vop3p.h').read_text() @@ -5103,6 +5102,26 @@ def test_cdna5_model_only_variant_instructions_have_no_execution_callbacks( assert class_name not in backend_source assert class_name not in execution_source + executable_class = 'VPkLshlAddU64Vop3p' + class_body = header.split(f'class {executable_class} ', 1)[1].split('\n};', 1)[0] + assert 'execute_impl' in class_body + constructor = model.split(f'{executable_class}::{executable_class}(', 1)[1].split( + '\n}', 1 + )[0] + assert 'selected_exec_fn(InstructionExecutionId::VPkLshlAddU64Vop3p)' in constructor + assert executable_class in backend_header + assert executable_class in backend_source + assert executable_class in execution_source + assert 'PkU64Pair read_pk_u64_pair' in execution_source + assert 'PkU32Pair read_pk_u32_pair' in execution_source + assert 'void write_pk_u64_pair' in execution_source + u64_read_helper = execution_source.split('PkU64Pair read_pk_u64_pair', 1)[1].split( + '\n}', 1 + )[0] + assert 'const auto reg = operand.to_register_ref();' in u64_read_helper + assert 'if (!reg || reg->cls != RegClass::VGPR)' in u64_read_helper + assert 'is_general_register' not in u64_read_helper + def test_generated_vop_execution_has_no_instruction_storage_bypass( amdgpu_generated_root: Path, diff --git a/emulation/rocjitsu/lib/python/amdisa/tests/test_packed_codegen.py b/emulation/rocjitsu/lib/python/amdisa/tests/test_packed_codegen.py index 6e3a16e548ab..2bc266862291 100644 --- a/emulation/rocjitsu/lib/python/amdisa/tests/test_packed_codegen.py +++ b/emulation/rocjitsu/lib/python/amdisa/tests/test_packed_codegen.py @@ -15,6 +15,7 @@ gen_pk_binop_f32, gen_pk_fmac_vop2, gen_pk_fmac_vop3, + gen_pk_lshl_add_u64, gen_pk_ternary, ) from amdisa.codegen.execute.simd_codegen import vop3p_local_simd_probe_line @@ -289,6 +290,29 @@ def test_cdna_pk_f32_reads_all_register_pairs(): assert 'encoding_value_ >= 256' not in cpp +def test_pk_lshl_add_u64_operates_on_two_independent_64_bit_elements(): + cpp = gen_pk_lshl_add_u64(['vdst'], ['src0', 'src1', 'src2']) + + assert 'const auto values = read_pk_u64_pair(src0, wf, lane);' in cpp + assert 'const auto shifts = read_pk_u32_pair(src1, wf, lane);' in cpp + assert 'const auto addends = read_pk_u64_pair(src2, wf, lane);' in cpp + assert 'Public semantics cover shift counts 0..4' in cpp + assert 'only avoids host undefined behavior for unsupported values' in cpp + assert ( + 'amdgpu::lshl_masked(values.lo, static_cast(shifts.lo)) + addends.lo' + in cpp + ) + assert ( + 'amdgpu::lshl_masked(values.hi, static_cast(shifts.hi)) + addends.hi' + in cpp + ) + assert 'lshl_masked(values.lo, shifts.lo)' not in cpp + assert 'lshl_masked(values.hi, shifts.hi)' not in cpp + assert 'write_pk_u64_pair(vdst, wf, lane, {result_lo, result_hi});' in cpp + assert 'inst_.neg' not in cpp + assert 'inst_.clamp' not in cpp + + def test_renamed_vop3p_packed_f32_probe_passes_profile_selectors(): probe = vop3p_local_simd_probe_line('v_pk_add_f32_vop3p', ('opsel', 'opsel_hi')) diff --git a/emulation/rocjitsu/lib/python/amdisa/tests/test_sema_derive.py b/emulation/rocjitsu/lib/python/amdisa/tests/test_sema_derive.py index 02b193ba8296..e963a575c0b5 100644 --- a/emulation/rocjitsu/lib/python/amdisa/tests/test_sema_derive.py +++ b/emulation/rocjitsu/lib/python/amdisa/tests/test_sema_derive.py @@ -2493,6 +2493,14 @@ def test_pk_binop_f32(self): block = derive_sema_block(sem) assert block is not None + def test_pk_lshl_add_u64_has_distinct_packed_width_semantics(self): + sem = derive_semantics('V_PK_LSHL_ADD_U64', 'ENC_VOP3P') + + assert sem is not None + assert sem.semantic_class == 'pk_lshl_add_u64' + assert sem.operation == 'lshl_add' + assert sem.data_type == 'u64' + def test_pk_mov_b32(self): sem = _FakeSem('V_PK_MOV_B32', 'pk_mov_b32') block = derive_sema_block(sem) diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/generated/cdna5/execution_backend.h b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/generated/cdna5/execution_backend.h index a344a36bb530..4d3552c164e6 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/generated/cdna5/execution_backend.h +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/generated/cdna5/execution_backend.h @@ -1269,6 +1269,7 @@ enum class InstructionExecutionId : size_t { VWmmaF1616x16x128Bf8Fp8Vop3p, VWmmaF1616x16x128Bf8Bf8Vop3p, VWmmaF3232x16x128F4Vop3p, + VPkLshlAddU64Vop3p, VWmmaScaleF32Vop3px2, DsAddU32Vds, DsSubU32Vds, diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/generated/cdna5/execution_backend_exec.cpp b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/generated/cdna5/execution_backend_exec.cpp index c4ef09e1530a..bd0cabd7e674 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/generated/cdna5/execution_backend_exec.cpp +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/generated/cdna5/execution_backend_exec.cpp @@ -1276,6 +1276,7 @@ constexpr InstructionCallbackTable kInstructionCallbacks{{ &execute_with_backend, &execute_with_backend, &execute_with_backend, + &execute_with_backend, &execute_with_backend, &execute_with_backend, &execute_with_backend, diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/generated/cdna5/vop3p.cpp b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/generated/cdna5/vop3p.cpp index 14a79b1a3ba3..49278cc58835 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/generated/cdna5/vop3p.cpp +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/generated/cdna5/vop3p.cpp @@ -4461,7 +4461,8 @@ DecodeResult decodeVPkMinNumF64Vop3p(const MachineInst *opcode, } // namespace detail VPkLshlAddU64Vop3p::VPkLshlAddU64Vop3p(const MachineInst *inst) - : Vop3p("v_pk_lshl_add_u64", reinterpret_cast(inst), nullptr), + : Vop3p("v_pk_lshl_add_u64", reinterpret_cast(inst), + selected_exec_fn(InstructionExecutionId::VPkLshlAddU64Vop3p)), vdst(128, OperandType::OPR_VGPR, reinterpret_cast(inst)->vdst), src0(128, OperandType::OPR_SRC, reinterpret_cast(inst)->src0), src1(64, OperandType::OPR_SRC, reinterpret_cast(inst)->src1), diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/generated/cdna5/vop3p.h b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/generated/cdna5/vop3p.h index 6c4ed8494f82..0752b60d9eb2 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/generated/cdna5/vop3p.h +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/generated/cdna5/vop3p.h @@ -992,6 +992,7 @@ class VPkMinNumF64Vop3p : public Vop3p { class VPkLshlAddU64Vop3p : public Vop3p { public: VPkLshlAddU64Vop3p(const MachineInst *inst); + void execute_impl(amdgpu::Wavefront &wf); Operand vdst; Operand src0; Operand src1; diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/generated/cdna5/vop3p_exec.cpp b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/generated/cdna5/vop3p_exec.cpp index 08be7bcd4409..2c8f8b96fde9 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/generated/cdna5/vop3p_exec.cpp +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/generated/cdna5/vop3p_exec.cpp @@ -36,6 +36,54 @@ PkF32Words read_pk_f32_words(const Operand &operand, const amdgpu::Wavefront &wf return {pair.lo, pair.hi}; } +struct PkU64Pair { + uint64_t lo; + uint64_t hi; +}; + +struct PkU32Pair { + uint32_t lo; + uint32_t hi; +}; + +bool is_general_register(const std::optional ®) { + return reg && (reg->cls == RegClass::SGPR || reg->cls == RegClass::VGPR); +} + +Operand packed_register_dword_offset(const Operand &operand, uint32_t dword_offset) { + Operand shifted = operand; + shifted.encoding_value_ += static_cast(dword_offset); + return shifted; +} + +PkU64Pair read_pk_u64_pair(const Operand &operand, const amdgpu::Wavefront &wf, uint32_t lane) { + const uint64_t lo = amdgpu::RegisterAccess(wf).read_lane64(operand, lane); + const auto reg = operand.to_register_ref(); + if (!reg || reg->cls != RegClass::VGPR) + return {lo, lo}; + + const Operand hi_operand = packed_register_dword_offset(operand, 2); + return {lo, amdgpu::RegisterAccess(wf).read_lane64(hi_operand, lane)}; +} + +PkU32Pair read_pk_u32_pair(const Operand &operand, const amdgpu::Wavefront &wf, uint32_t lane) { + if (!is_general_register(operand.to_register_ref())) { + const uint32_t value = amdgpu::RegisterAccess(wf).read_lane(operand, lane); + return {value, value}; + } + + const uint64_t raw = amdgpu::RegisterAccess(wf).read_lane64(operand, lane); + return {static_cast(raw), static_cast(raw >> 32)}; +} + +void write_pk_u64_pair(const Operand &operand, amdgpu::Wavefront &wf, uint32_t lane, + PkU64Pair value) { + const Operand hi_operand = packed_register_dword_offset(operand, 2); + amdgpu::RegisterAccess access(wf); + access.write_lane64(operand, lane, value.lo); + access.write_lane64(hi_operand, lane, value.hi); +} + uint16_t read_fma_mix_f16_bits(uint32_t raw, uint32_t src_selector, bool high_half) { switch (src_selector) { case OpSelSrc::OPR_SRC_FLOAT_HALF: @@ -2945,6 +2993,24 @@ void VWmmaF3232x16x128F4Vop3p::execute_impl(amdgpu::Wavefront &wf) { amdgpu::wmma_c_modifier(inst_.neg, inst_.neg_hi)); } +void VPkLshlAddU64Vop3p::execute_impl(amdgpu::Wavefront &wf) { + uint64_t exec = wf.exec(); + for (uint32_t lane = 0; lane < wf.wf_size(); ++lane) { + if (!(exec & (1ULL << lane))) + continue; + const auto values = read_pk_u64_pair(src0, wf, lane); + const auto shifts = read_pk_u32_pair(src1, wf, lane); + const auto addends = read_pk_u64_pair(src2, wf, lane); + // Public semantics cover shift counts 0..4. The masked helper only avoids host undefined + // behavior for unsupported values. + const uint64_t result_lo = + amdgpu::lshl_masked(values.lo, static_cast(shifts.lo)) + addends.lo; + const uint64_t result_hi = + amdgpu::lshl_masked(values.hi, static_cast(shifts.hi)) + addends.hi; + write_pk_u64_pair(vdst, wf, lane, {result_lo, result_hi}); + } +} + void VWmmaScaleF32Vop3px2::execute_impl(amdgpu::Wavefront &wf) { auto &cu = wf.cu(); uint32_t vb = wf.vgpr_alloc().base; diff --git a/emulation/rocjitsu/tests/cdna5_execution_test.cpp b/emulation/rocjitsu/tests/cdna5_execution_test.cpp index f6c74c4d5adf..05d65a15058d 100644 --- a/emulation/rocjitsu/tests/cdna5_execution_test.cpp +++ b/emulation/rocjitsu/tests/cdna5_execution_test.cpp @@ -1142,7 +1142,45 @@ TEST(Gfx1250LiteralOperandTest, PkF32LiteralReplicatesAndUsesAvailableSimdPath) run_case(false); } -TEST(Gfx1251ModelLiteralOperandTest, PackedU32LiteralReplicatesBothLanes) { +using PackedU64Pair = std::array; + +void write_vgpr_packed_u64(amdgpu::ComputeUnitCore &cu, const amdgpu::Wavefront &wf, + uint32_t first_vgpr, uint32_t lane, PackedU64Pair values) { + const uint32_t base = wf.vgpr_alloc().base + first_vgpr; + for (uint32_t element = 0; element < values.size(); ++element) { + cu.write_vgpr(base + element * 2, lane, static_cast(values[element])); + cu.write_vgpr(base + element * 2 + 1, lane, static_cast(values[element] >> 32)); + } +} + +PackedU64Pair read_vgpr_packed_u64(const amdgpu::ComputeUnitCore &cu, const amdgpu::Wavefront &wf, + uint32_t first_vgpr, uint32_t lane) { + const uint32_t base = wf.vgpr_alloc().base + first_vgpr; + PackedU64Pair values{}; + for (uint32_t element = 0; element < values.size(); ++element) { + values[element] = cu.read_vgpr(base + element * 2, lane); + values[element] |= static_cast(cu.read_vgpr(base + element * 2 + 1, lane)) << 32; + } + return values; +} + +void write_vgpr_packed_u32(amdgpu::ComputeUnitCore &cu, const amdgpu::Wavefront &wf, + uint32_t first_vgpr, uint32_t lane, std::array values) { + const uint32_t base = wf.vgpr_alloc().base + first_vgpr; + cu.write_vgpr(base, lane, values[0]); + cu.write_vgpr(base + 1, lane, values[1]); +} + +void write_sgpr_packed_u64(amdgpu::ComputeUnitCore &cu, amdgpu::Wavefront &wf, uint32_t first_sgpr, + PackedU64Pair values) { + for (uint32_t element = 0; element < values.size(); ++element) { + write_wave_sgpr(cu, wf, first_sgpr + element * 2, static_cast(values[element])); + write_wave_sgpr(cu, wf, first_sgpr + element * 2 + 1, + static_cast(values[element] >> 32)); + } +} + +TEST(Gfx1251PackedU64ExecutionTest, PublicLiteralEncodingReplicatesWithoutSemanticOracle) { constexpr uint32_t literal = 0x65u; constexpr uint64_t replicated = (static_cast(literal) << 32) | static_cast(literal); @@ -1150,15 +1188,15 @@ TEST(Gfx1251ModelLiteralOperandTest, PackedU32LiteralReplicatesBothLanes) { // v_pk_lshl_add_u64 v[4:7], v[8:11], 101, v[16:19]. constexpr std::array words{0xCC7E4004u, 0x1C41FF08u, literal}; - // The public gfx1251 provider deliberately has no execution backend. Link - // the existing operand backend directly for this isolated operand-read test; - // the gfx1251-only instruction still has no instruction execution callback. + // LLVM's public MC test proves that literal 101 is accepted by this source + // form, but public semantics only cover shift counts 0..4. Decode and inspect + // the operand without executing 101 as an architectural arithmetic oracle. auto decoder = make_isa_decoder(&cdna5::execution_backend(), cdna5::kGfx1251IsaFeatures); ASSERT_NE(decoder, nullptr); std::unique_ptr decoded(decode_valid(*decoder, words.data())); ASSERT_NE(decoded, nullptr); - EXPECT_EQ(decoded->execute, nullptr); + ASSERT_NE(decoded->execute, nullptr); ASSERT_EQ(decoded->num_src_operands(), 3); const Operand *packed_shift = decoded->src_operand(1); ASSERT_NE(packed_shift, nullptr); @@ -1170,6 +1208,145 @@ TEST(Gfx1251ModelLiteralOperandTest, PackedU32LiteralReplicatesBothLanes) { EXPECT_EQ(amdgpu::RegisterAccess(*wf).read_lane64(*packed_shift, 0), replicated); } +TEST(Gfx1251PackedU64ExecutionTest, PublicVgprVectorExecutesSupportedShiftsAndActiveLanes) { + // Public LLVM gfx1251_asm_vop3p.s encoding for + // v_pk_lshl_add_u64 v[4:7], v[8:11], v[12:13], v[16:19]. + constexpr std::array words{0xCC7E4004u, 0x1C421908u}; + auto decoder = + make_isa_decoder(&cdna5::execution_backend(), cdna5::kGfx1251IsaFeatures); + ASSERT_NE(decoder, nullptr); + std::unique_ptr decoded(decode_valid(*decoder, words.data())); + ASSERT_NE(decoded, nullptr); + ASSERT_NE(decoded->execute, nullptr); + + Gfx1250Sim sim; + auto *cu = sim.cu(); + auto *wf = cu->dispatch_wf(0, 0, kGfx1250ScalarSlots, 32); + ASSERT_NE(wf, nullptr); + wf->set_exec(0x7u); + + constexpr std::array kValues{{ + {2u, 3u}, + {1u, std::numeric_limits::max()}, + {std::numeric_limits::max(), 0x8000000000000000ULL}, + {0x1111111111111111ULL, 0x2222222222222222ULL}, + }}; + constexpr std::array, 4> kShifts{{ + {0u, 1u}, + {4u, 0u}, + {1u, 4u}, + {4u, 4u}, + }}; + constexpr std::array kAddends{{ + {5u, 7u}, + {std::numeric_limits::max(), 1u}, + {2u, std::numeric_limits::max()}, + {3u, 3u}, + }}; + constexpr PackedU64Pair kInactiveSeed{0xdeadbeefcafef00dULL, 0xbaadf00d12345678ULL}; + for (uint32_t lane = 0; lane < kValues.size(); ++lane) { + write_vgpr_packed_u64(*cu, *wf, 8, lane, kValues[lane]); + write_vgpr_packed_u32(*cu, *wf, 12, lane, kShifts[lane]); + EXPECT_LE(kShifts[lane][0], 4u); + EXPECT_LE(kShifts[lane][1], 4u); + write_vgpr_packed_u64(*cu, *wf, 16, lane, kAddends[lane]); + write_vgpr_packed_u64(*cu, *wf, 4, lane, kInactiveSeed); + } + + cu->execute_instruction(decoded.get(), *wf); + + EXPECT_EQ(read_vgpr_packed_u64(*cu, *wf, 4, 0), (PackedU64Pair{7u, 13u})); + EXPECT_EQ(read_vgpr_packed_u64(*cu, *wf, 4, 1), (PackedU64Pair{15u, 0u})); + EXPECT_EQ(read_vgpr_packed_u64(*cu, *wf, 4, 2), + (PackedU64Pair{0u, std::numeric_limits::max()})); + EXPECT_EQ(read_vgpr_packed_u64(*cu, *wf, 4, 3), kInactiveSeed); +} + +TEST(Gfx1251PackedU64ExecutionTest, ExecutesEveryPublicLlvmSourceForm) { + struct SourceFormCase { + std::string_view name; + std::array words; + PackedU64Pair expected; + }; + // Encodings come from public LLVM gfx1251_asm_vop3p.s. A zero third dword is + // ignored for eight-byte forms. LLVM's literal-shift vector uses unsupported + // value 101, so its literal payload is replaced with supported value 4 for + // arithmetic coverage; the exact vector remains decode-only above. + constexpr std::array kCases{ + SourceFormCase{"vgpr-vgpr-vgpr", {0xCC7E4004u, 0x1C421908u, 0u}, {39u, 59u}}, + SourceFormCase{"sgpr-sgpr-vgpr", {0xCC7E4004u, 0x1C401808u, 0u}, {17u, 31u}}, + SourceFormCase{"vgpr-sgpr-vgpr", {0xCC7E4004u, 0x1C401908u, 0u}, {11u, 23u}}, + SourceFormCase{"sgpr-vgpr-vgpr", {0xCC7E4004u, 0x1C421808u, 0u}, {87u, 91u}}, + SourceFormCase{"vgpr-null-vgpr", {0xCC7E4004u, 0x1C40F908u, 0u}, {9u, 14u}}, + SourceFormCase{"vgpr-inline-vgpr", {0xCC7E4004u, 0x1C410308u, 0u}, {11u, 17u}}, + SourceFormCase{"inline-vgpr-vgpr", {0xCC7E4004u, 0x1C421081u, 0u}, {15u, 27u}}, + SourceFormCase{"literal-vgpr-vgpr", {0xCC7E4004u, 0x1C4210FFu, 0x65u}, {815u, 1627u}}, + SourceFormCase{"vgpr-literal-vgpr", {0xCC7E4004u, 0x1C41FF08u, 4u}, {39u, 59u}}, + SourceFormCase{"vgpr-vgpr-sgpr", {0xCC7E4004u, 0x18421908u, 0u}, {45u, 61u}}, + SourceFormCase{"vgpr-vgpr-null", {0xCC7E4004u, 0x19F21908u, 0u}, {32u, 48u}}, + SourceFormCase{"vgpr-vgpr-inline", {0xCC7E4004u, 0x1A061908u, 0u}, {33u, 49u}}, + SourceFormCase{"vgpr-vgpr-literal", {0xCC7E4004u, 0x1BFE2108u, 0x65u}, {133u, 149u}}, + }; + + auto decoder = + make_isa_decoder(&cdna5::execution_backend(), cdna5::kGfx1251IsaFeatures); + ASSERT_NE(decoder, nullptr); + Gfx1250Sim sim; + auto *cu = sim.cu(); + auto *wf = cu->dispatch_wf(0, 0, kGfx1250ScalarSlots, 32); + ASSERT_NE(wf, nullptr); + wf->set_exec(1u); + constexpr PackedU64Pair kSgprSrc0WithPoison{5u, 0xdeadbeefcafef00dULL}; + constexpr PackedU64Pair kSgprSrc2WithPoison{13u, 0xbaadf00d12345678ULL}; + + for (const auto &test_case : kCases) { + SCOPED_TRACE(test_case.name); + write_vgpr_packed_u64(*cu, *wf, 8, 0, {2u, 3u}); + write_vgpr_packed_u32(*cu, *wf, 12, 0, {4u, 4u}); + write_vgpr_packed_u64(*cu, *wf, 16, 0, {7u, 11u}); + write_sgpr_packed_u64(*cu, *wf, 8, kSgprSrc0WithPoison); + write_wave_sgpr(*cu, *wf, 12, 1u); + write_wave_sgpr(*cu, *wf, 13, 2u); + write_sgpr_packed_u64(*cu, *wf, 16, kSgprSrc2WithPoison); + // The public inline/literal-src0 vectors use v[8:9] for their shifts. + // The final literal-src2 vector uses v[16:17] for its shifts. + if (test_case.name == "inline-vgpr-vgpr" || test_case.name == "literal-vgpr-vgpr") + write_vgpr_packed_u32(*cu, *wf, 8, 0, {3u, 4u}); + if (test_case.name == "vgpr-vgpr-literal") + write_vgpr_packed_u32(*cu, *wf, 16, 0, {4u, 4u}); + write_vgpr_packed_u64(*cu, *wf, 4, 0, {0u, 0u}); + + std::unique_ptr decoded(decode_valid(*decoder, test_case.words.data())); + ASSERT_NE(decoded, nullptr); + ASSERT_NE(decoded->execute, nullptr); + cu->execute_instruction(decoded.get(), *wf); + EXPECT_EQ(read_vgpr_packed_u64(*cu, *wf, 4, 0), test_case.expected); + } +} + +TEST(Gfx1251PackedU64ExecutionTest, ReadsOverlappingSourcesBeforeWritingDestination) { + const auto words = cdna5::build_vop3p(cdna5::kVPkLshlAddU64Vop3p, + {.vdst = 8, .src0 = 264, .src1 = 268, .src2 = 272}); + auto decoder = + make_isa_decoder(&cdna5::execution_backend(), cdna5::kGfx1251IsaFeatures); + ASSERT_NE(decoder, nullptr); + std::unique_ptr decoded(decode_valid(*decoder, words.data())); + ASSERT_NE(decoded, nullptr); + ASSERT_NE(decoded->execute, nullptr); + + Gfx1250Sim sim; + auto *cu = sim.cu(); + auto *wf = cu->dispatch_wf(0, 0, kGfx1250ScalarSlots, 32); + ASSERT_NE(wf, nullptr); + wf->set_exec(1u); + write_vgpr_packed_u64(*cu, *wf, 8, 0, {0x1000000000000001ULL, 3u}); + write_vgpr_packed_u32(*cu, *wf, 12, 0, {4u, 4u}); + write_vgpr_packed_u64(*cu, *wf, 16, 0, {7u, 11u}); + + cu->execute_instruction(decoded.get(), *wf); + EXPECT_EQ(read_vgpr_packed_u64(*cu, *wf, 8, 0), (PackedU64Pair{0x0000000000000017ULL, 59u})); +} + TEST(Gfx1250LiteralOperandTest, PkF32MixedLiteralVgprSourcesUseAvailableSimdPath) { if constexpr (!util::has_stdx_simd) GTEST_SKIP() << "requires stdx SIMD"; diff --git a/shared/machine-readable-isa/isa/amdgpu_isa_cdna5_variants.json b/shared/machine-readable-isa/isa/amdgpu_isa_cdna5_variants.json index 6d4707e56650..7352108e95d8 100644 --- a/shared/machine-readable-isa/isa/amdgpu_isa_cdna5_variants.json +++ b/shared/machine-readable-isa/isa/amdgpu_isa_cdna5_variants.json @@ -52,7 +52,6 @@ "V_PK_SUB_NC_U64", "V_PK_MAX_NUM_F64", "V_PK_MIN_NUM_F64", - "V_PK_LSHL_ADD_U64", "V_WMMA_F64_16X16X4_F64" ], "encodings": [ From bc62b32fe8c1c3c8344778f7c753b5266ecc9bde Mon Sep 17 00:00:00 2001 From: hanhanW Date: Wed, 26 Aug 2026 17:46:29 -0700 Subject: [PATCH 2/3] address comments Signed-off-by: hanhanW --- .../lib/python/amdisa/codegen/_generator.py | 14 ++- .../python/amdisa/codegen/execute/packed.py | 26 ++++-- .../tests/test_generator_profile_gates.py | 8 +- .../amdisa/tests/test_packed_codegen.py | 17 +++- .../amdgpu/generated/cdna5/vop3p_exec.cpp | 29 +++--- .../src/rocjitsu/vm/amdgpu/compute_unit.cpp | 9 ++ .../src/rocjitsu/vm/amdgpu/compute_unit.h | 1 + .../src/rocjitsu/vm/amdgpu/wavefront.h | 37 +++++++- .../rocjitsu/tests/cdna5_execution_test.cpp | 91 ++++++++++++++++++- 9 files changed, 192 insertions(+), 40 deletions(-) diff --git a/emulation/rocjitsu/lib/python/amdisa/codegen/_generator.py b/emulation/rocjitsu/lib/python/amdisa/codegen/_generator.py index 8917b0c9bfa6..232a2bbcf4d2 100644 --- a/emulation/rocjitsu/lib/python/amdisa/codegen/_generator.py +++ b/emulation/rocjitsu/lib/python/amdisa/codegen/_generator.py @@ -5307,10 +5307,6 @@ def _emit_cdna5_matrix_fmt_helpers(self) -> _ImplOutputs: uint32_t hi; }; - bool is_general_register(const std::optional ®) { - return reg && (reg->cls == RegClass::SGPR || reg->cls == RegClass::VGPR); - } - Operand packed_register_dword_offset(const Operand &operand, uint32_t dword_offset) { Operand shifted = operand; shifted.encoding_value_ += static_cast(dword_offset); @@ -5330,13 +5326,15 @@ def _emit_cdna5_matrix_fmt_helpers(self) -> _ImplOutputs: PkU32Pair read_pk_u32_pair(const Operand &operand, const amdgpu::Wavefront &wf, uint32_t lane) { - if (!is_general_register(operand.to_register_ref())) { + // GFX12+ single-SGPR-read operands read the first SGPR and replicate it. + // VGPRs and 64-bit special registers such as VCC and EXEC remain pairs. + const auto reg = operand.to_register_ref(); + if (reg && reg->cls == RegClass::SGPR) { const uint32_t value = amdgpu::RegisterAccess(wf).read_lane(operand, lane); return {value, value}; } - - const uint64_t raw = amdgpu::RegisterAccess(wf).read_lane64(operand, lane); - return {static_cast(raw), static_cast(raw >> 32)}; + const auto pair = amdgpu::RegisterAccess(wf).read_lane_pair32(operand, lane); + return {pair.lo, pair.hi}; } void write_pk_u64_pair(const Operand &operand, amdgpu::Wavefront &wf, uint32_t lane, diff --git a/emulation/rocjitsu/lib/python/amdisa/codegen/execute/packed.py b/emulation/rocjitsu/lib/python/amdisa/codegen/execute/packed.py index a131b13c6015..97bb39d5e866 100644 --- a/emulation/rocjitsu/lib/python/amdisa/codegen/execute/packed.py +++ b/emulation/rocjitsu/lib/python/amdisa/codegen/execute/packed.py @@ -643,28 +643,34 @@ def gen_pk_lshl_add_u64(dst: list[str], src: list[str]) -> str: V_PK_LSHL_ADD_U64 has two 64-bit values in each 128-bit value/addend operand and two 32-bit shift counts in its 64-bit shift operand. VGPR U64 - sources supply both elements, while SGPR and non-register U64 sources - broadcast the low 64 bits. Its public LLVM profile explicitly disables - clamp and source modifiers, so the body intentionally consumes neither - field. Public LLVM selection patterns and the public CDNA5 description - define shift counts only in the inclusive range 0..4. ``lshl_masked`` - keeps the host operation defined if an unsupported runtime value reaches - the callback; that fallback is not an architectural claim for counts above - 4. + sources supply both elements, while scalar-backed U64 sources broadcast the + low 64 bits. Its public LLVM profile explicitly disables clamp and source + modifiers, so the body intentionally consumes neither field. The currently + proven shift-count range is 0..4; execution fails closed before any + destination write if either count of any active lane is outside that range. """ d, s0, s1, s2 = dst[0], src[0], src[1], src[2] return '\n'.join( [ ' uint64_t exec = wf.exec();', + ' std::array results{};', ' for (uint32_t lane = 0; lane < wf.wf_size(); ++lane) {', ' if (!(exec & (1ULL << lane))) continue;', f' const auto values = read_pk_u64_pair({s0}, wf, lane);', f' const auto shifts = read_pk_u32_pair({s1}, wf, lane);', f' const auto addends = read_pk_u64_pair({s2}, wf, lane);', - ' // Public semantics cover shift counts 0..4. The masked helper only avoids host undefined behavior for unsupported values.', + ' if (shifts.lo > 4u || shifts.hi > 4u) {', + ' wf.report_instruction_execution_error(', + ' amdgpu::InstructionExecutionError::UnsupportedOperandValue);', + ' return;', + ' }', ' const uint64_t result_lo = amdgpu::lshl_masked(values.lo, static_cast(shifts.lo)) + addends.lo;', ' const uint64_t result_hi = amdgpu::lshl_masked(values.hi, static_cast(shifts.hi)) + addends.hi;', - f' write_pk_u64_pair({d}, wf, lane, {{result_lo, result_hi}});', + ' results[lane] = {result_lo, result_hi};', + ' }', + ' for (uint32_t lane = 0; lane < wf.wf_size(); ++lane) {', + ' if (!(exec & (1ULL << lane))) continue;', + f' write_pk_u64_pair({d}, wf, lane, results[lane]);', ' }', ] ) 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 140e3cd8dd7d..dd7e2a875bc9 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 @@ -5120,7 +5120,13 @@ def test_cdna5_model_only_variant_instructions_have_no_execution_callbacks( )[0] assert 'const auto reg = operand.to_register_ref();' in u64_read_helper assert 'if (!reg || reg->cls != RegClass::VGPR)' in u64_read_helper - assert 'is_general_register' not in u64_read_helper + u32_read_helper = execution_source.split('PkU32Pair read_pk_u32_pair', 1)[1].split( + '\n}', 1 + )[0] + assert 'const auto reg = operand.to_register_ref();' in u32_read_helper + assert 'if (reg && reg->cls == RegClass::SGPR)' in u32_read_helper + assert 'read_lane(operand, lane)' in u32_read_helper + assert 'read_lane_pair32(operand, lane)' in u32_read_helper def test_generated_vop_execution_has_no_instruction_storage_bypass( diff --git a/emulation/rocjitsu/lib/python/amdisa/tests/test_packed_codegen.py b/emulation/rocjitsu/lib/python/amdisa/tests/test_packed_codegen.py index 2bc266862291..97b094478b18 100644 --- a/emulation/rocjitsu/lib/python/amdisa/tests/test_packed_codegen.py +++ b/emulation/rocjitsu/lib/python/amdisa/tests/test_packed_codegen.py @@ -296,8 +296,6 @@ def test_pk_lshl_add_u64_operates_on_two_independent_64_bit_elements(): assert 'const auto values = read_pk_u64_pair(src0, wf, lane);' in cpp assert 'const auto shifts = read_pk_u32_pair(src1, wf, lane);' in cpp assert 'const auto addends = read_pk_u64_pair(src2, wf, lane);' in cpp - assert 'Public semantics cover shift counts 0..4' in cpp - assert 'only avoids host undefined behavior for unsupported values' in cpp assert ( 'amdgpu::lshl_masked(values.lo, static_cast(shifts.lo)) + addends.lo' in cpp @@ -308,11 +306,24 @@ def test_pk_lshl_add_u64_operates_on_two_independent_64_bit_elements(): ) assert 'lshl_masked(values.lo, shifts.lo)' not in cpp assert 'lshl_masked(values.hi, shifts.hi)' not in cpp - assert 'write_pk_u64_pair(vdst, wf, lane, {result_lo, result_hi});' in cpp + assert 'results[lane] = {result_lo, result_hi};' in cpp + assert 'write_pk_u64_pair(vdst, wf, lane, results[lane]);' in cpp assert 'inst_.neg' not in cpp assert 'inst_.clamp' not in cpp +def test_pk_lshl_add_u64_rejects_unproven_counts_before_any_write(): + cpp = gen_pk_lshl_add_u64(['vdst'], ['src0', 'src1', 'src2']) + + reject = 'if (shifts.lo > 4u || shifts.hi > 4u)' + report = 'wf.report_instruction_execution_error(' + write = 'write_pk_u64_pair(vdst, wf, lane, results[lane]);' + assert reject in cpp + assert report in cpp + assert write in cpp + assert cpp.index(reject) < cpp.index(report) < cpp.index(write) + + def test_renamed_vop3p_packed_f32_probe_passes_profile_selectors(): probe = vop3p_local_simd_probe_line('v_pk_add_f32_vop3p', ('opsel', 'opsel_hi')) diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/generated/cdna5/vop3p_exec.cpp b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/generated/cdna5/vop3p_exec.cpp index 2c8f8b96fde9..189e3ba345c4 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/generated/cdna5/vop3p_exec.cpp +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/generated/cdna5/vop3p_exec.cpp @@ -46,10 +46,6 @@ struct PkU32Pair { uint32_t hi; }; -bool is_general_register(const std::optional ®) { - return reg && (reg->cls == RegClass::SGPR || reg->cls == RegClass::VGPR); -} - Operand packed_register_dword_offset(const Operand &operand, uint32_t dword_offset) { Operand shifted = operand; shifted.encoding_value_ += static_cast(dword_offset); @@ -67,13 +63,15 @@ PkU64Pair read_pk_u64_pair(const Operand &operand, const amdgpu::Wavefront &wf, } PkU32Pair read_pk_u32_pair(const Operand &operand, const amdgpu::Wavefront &wf, uint32_t lane) { - if (!is_general_register(operand.to_register_ref())) { + // GFX12+ single-SGPR-read operands read the first SGPR and replicate it. + // VGPRs and 64-bit special registers such as VCC and EXEC remain pairs. + const auto reg = operand.to_register_ref(); + if (reg && reg->cls == RegClass::SGPR) { const uint32_t value = amdgpu::RegisterAccess(wf).read_lane(operand, lane); return {value, value}; } - - const uint64_t raw = amdgpu::RegisterAccess(wf).read_lane64(operand, lane); - return {static_cast(raw), static_cast(raw >> 32)}; + const auto pair = amdgpu::RegisterAccess(wf).read_lane_pair32(operand, lane); + return {pair.lo, pair.hi}; } void write_pk_u64_pair(const Operand &operand, amdgpu::Wavefront &wf, uint32_t lane, @@ -2995,19 +2993,28 @@ void VWmmaF3232x16x128F4Vop3p::execute_impl(amdgpu::Wavefront &wf) { void VPkLshlAddU64Vop3p::execute_impl(amdgpu::Wavefront &wf) { uint64_t exec = wf.exec(); + std::array results{}; for (uint32_t lane = 0; lane < wf.wf_size(); ++lane) { if (!(exec & (1ULL << lane))) continue; const auto values = read_pk_u64_pair(src0, wf, lane); const auto shifts = read_pk_u32_pair(src1, wf, lane); const auto addends = read_pk_u64_pair(src2, wf, lane); - // Public semantics cover shift counts 0..4. The masked helper only avoids host undefined - // behavior for unsupported values. + if (shifts.lo > 4u || shifts.hi > 4u) { + wf.report_instruction_execution_error( + amdgpu::InstructionExecutionError::UnsupportedOperandValue); + return; + } const uint64_t result_lo = amdgpu::lshl_masked(values.lo, static_cast(shifts.lo)) + addends.lo; const uint64_t result_hi = amdgpu::lshl_masked(values.hi, static_cast(shifts.hi)) + addends.hi; - write_pk_u64_pair(vdst, wf, lane, {result_lo, result_hi}); + results[lane] = {result_lo, result_hi}; + } + for (uint32_t lane = 0; lane < wf.wf_size(); ++lane) { + if (!(exec & (1ULL << lane))) + continue; + write_pk_u64_pair(vdst, wf, lane, results[lane]); } } 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 11d9ffe811d6..8e2d27372f96 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 @@ -884,6 +884,15 @@ void ComputeUnitCore::issue_instruction(Wavefront *active) { execute_instruction(inst, *active); + if (active->instruction_execution_failed()) { + util::Logger::warn("CU ", this->name(), ": wf", active->wf_id(), + " halted after unsupported operand value in ", inst->mnemonic(), " at pc=0x", + std::hex, active->pc); + active->halt(Wavefront::CpCompletionNotice::Suppress); + delete inst; + return; + } + // A terminating instruction (s_endpgm with no pending waits) halts the wave // inside execute_instruction, which frees and resets its slot. Its registers, // pc, and allocations are now zeroed, so the after-execute hook, result logging, 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 f0935f429e82..a5d40037967c 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 @@ -1328,6 +1328,7 @@ class IsaExecComputeUnit : public ExecComputeUnit { /// @brief Execute one instruction on the given wavefront via direct dispatch. void execute_instruction(Instruction *inst, Wavefront &wf) override { assert(inst->execute && "instruction execution backend is not linked"); + wf.clear_instruction_execution_error(); inst->execute(*inst, &wf); } diff --git a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/wavefront.h b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/wavefront.h index 0ad1f80b9f37..e0e7836d0e21 100644 --- a/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/wavefront.h +++ b/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/amdgpu/wavefront.h @@ -39,6 +39,12 @@ enum class WfState : uint8_t { ENDING, ///< s_endpgm executed but outstanding memory ops are draining. }; +/// @brief Simulator execution failures that are not architectural wave state. +enum class InstructionExecutionError : uint8_t { + None, + UnsupportedOperandValue, +}; + /// @brief Allocation slice within a register file. struct RegAllocation { uint32_t base = 0; ///< First register index in the physical file. @@ -609,6 +615,25 @@ class Wavefront : public ThreadContext { /// @retval false Slot is active (running, waiting, or at a barrier). bool is_halted() const { return state_ == WfState::HALTED; } + /// @brief Record a fail-closed instruction execution error. + /// @details Execution callbacks use this for inputs whose architectural + /// behavior is not implemented. Callers must not treat it as a hardware trap. + void report_instruction_execution_error(InstructionExecutionError error) { + instruction_execution_error_ = error; + } + + InstructionExecutionError instruction_execution_error() const { + return instruction_execution_error_; + } + + bool instruction_execution_failed() const { + return instruction_execution_error_ != InstructionExecutionError::None; + } + + void clear_instruction_execution_error() { + instruction_execution_error_ = InstructionExecutionError::None; + } + // -- KFD debugger (trap) state -- // // These model the wave-level state the AMD trap handler maintains for @@ -872,6 +897,7 @@ class Wavefront : public ThreadContext { t = 0; trapsts_ = 0; pending_alu_causes_ = 0; + instruction_execution_error_ = InstructionExecutionError::None; sleep_cycles_ = 0; in_trap_handler_ = false; trap_interrupt_sent_ = false; @@ -913,11 +939,12 @@ class Wavefront : public ThreadContext { uint32_t wave_in_group_ = 0; ///< Position of this wave within its workgroup (debugger). uint32_t process_id_ = 0; ///< Owning process ID (PASID analog, set per dispatch). uint32_t queue_id_ = 0; ///< KFD queue ID that launched this wave (debugger correlation). - uint32_t lds_base_ = 0; ///< Per-WG LDS base offset (set per dispatch). - uint32_t lds_size_ = 0; ///< Aligned per-WG LDS allocation size. - Lds *lds_ = nullptr; ///< Placement-selected LDS backing; nullptr means CU-local LDS. - uint32_t cluster_rank_ = 0; ///< Workgroup rank inside the dispatch cluster. - uint32_t cluster_size_ = 1; ///< Number of workgroups in the dispatch cluster. + InstructionExecutionError instruction_execution_error_ = InstructionExecutionError::None; + uint32_t lds_base_ = 0; ///< Per-WG LDS base offset (set per dispatch). + uint32_t lds_size_ = 0; ///< Aligned per-WG LDS allocation size. + Lds *lds_ = nullptr; ///< Placement-selected LDS backing; nullptr means CU-local LDS. + uint32_t cluster_rank_ = 0; ///< Workgroup rank inside the dispatch cluster. + uint32_t cluster_size_ = 1; ///< Number of workgroups in the dispatch cluster. uint32_t wf_size_ = 0; ///< Lanes in the current dispatched wavefront. uint32_t default_wf_size_ = 0; ///< ISA default wavefront width. diff --git a/emulation/rocjitsu/tests/cdna5_execution_test.cpp b/emulation/rocjitsu/tests/cdna5_execution_test.cpp index 05d65a15058d..0f0003d72abb 100644 --- a/emulation/rocjitsu/tests/cdna5_execution_test.cpp +++ b/emulation/rocjitsu/tests/cdna5_execution_test.cpp @@ -1262,6 +1262,51 @@ TEST(Gfx1251PackedU64ExecutionTest, PublicVgprVectorExecutesSupportedShiftsAndAc EXPECT_EQ(read_vgpr_packed_u64(*cu, *wf, 4, 3), kInactiveSeed); } +TEST(Gfx1251PackedU64ExecutionTest, RejectsUnprovenShiftCountsBeforeAnyDestinationWrite) { + // The currently proven execution range is 0..4. Public LLVM separately shows + // that larger values are encodable; until their arithmetic meaning is public, + // fail closed without partially committing an earlier active lane. + constexpr std::array words{0xCC7E4004u, 0x1C421908u}; + constexpr std::array kUnprovenCounts{5u, 63u, 64u, 101u}; + constexpr PackedU64Pair kLane0Seed{0x1111222233334444ULL, 0x5555666677778888ULL}; + constexpr PackedU64Pair kLane1Seed{0x9999aaaabbbbccccULL, 0xddddeeeeffff0000ULL}; + + auto decoder = + make_isa_decoder(&cdna5::execution_backend(), cdna5::kGfx1251IsaFeatures); + ASSERT_NE(decoder, nullptr); + std::unique_ptr decoded(decode_valid(*decoder, words.data())); + ASSERT_NE(decoded, nullptr); + ASSERT_NE(decoded->execute, nullptr); + + Gfx1250Sim sim; + auto *cu = sim.cu(); + auto *wf = cu->dispatch_wf(0, 0, kGfx1250ScalarSlots, 32); + ASSERT_NE(wf, nullptr); + wf->set_exec(0x3u); + + for (uint32_t case_index = 0; case_index < kUnprovenCounts.size(); ++case_index) { + SCOPED_TRACE(kUnprovenCounts[case_index]); + write_vgpr_packed_u64(*cu, *wf, 8, 0, {2u, 3u}); + write_vgpr_packed_u64(*cu, *wf, 8, 1, {5u, 7u}); + write_vgpr_packed_u32(*cu, *wf, 12, 0, {0u, 4u}); + const std::array invalid_shifts = + case_index % 2 == 0 ? std::array{kUnprovenCounts[case_index], 0u} + : std::array{0u, kUnprovenCounts[case_index]}; + write_vgpr_packed_u32(*cu, *wf, 12, 1, invalid_shifts); + write_vgpr_packed_u64(*cu, *wf, 16, 0, {11u, 13u}); + write_vgpr_packed_u64(*cu, *wf, 16, 1, {17u, 19u}); + write_vgpr_packed_u64(*cu, *wf, 4, 0, kLane0Seed); + write_vgpr_packed_u64(*cu, *wf, 4, 1, kLane1Seed); + + cu->execute_instruction(decoded.get(), *wf); + + EXPECT_EQ(wf->instruction_execution_error(), + amdgpu::InstructionExecutionError::UnsupportedOperandValue); + EXPECT_EQ(read_vgpr_packed_u64(*cu, *wf, 4, 0), kLane0Seed); + EXPECT_EQ(read_vgpr_packed_u64(*cu, *wf, 4, 1), kLane1Seed); + } +} + TEST(Gfx1251PackedU64ExecutionTest, ExecutesEveryPublicLlvmSourceForm) { struct SourceFormCase { std::string_view name; @@ -1274,8 +1319,8 @@ TEST(Gfx1251PackedU64ExecutionTest, ExecutesEveryPublicLlvmSourceForm) { // arithmetic coverage; the exact vector remains decode-only above. constexpr std::array kCases{ SourceFormCase{"vgpr-vgpr-vgpr", {0xCC7E4004u, 0x1C421908u, 0u}, {39u, 59u}}, - SourceFormCase{"sgpr-sgpr-vgpr", {0xCC7E4004u, 0x1C401808u, 0u}, {17u, 31u}}, - SourceFormCase{"vgpr-sgpr-vgpr", {0xCC7E4004u, 0x1C401908u, 0u}, {11u, 23u}}, + SourceFormCase{"sgpr-sgpr-vgpr", {0xCC7E4004u, 0x1C401808u, 0u}, {17u, 21u}}, + SourceFormCase{"vgpr-sgpr-vgpr", {0xCC7E4004u, 0x1C401908u, 0u}, {11u, 17u}}, SourceFormCase{"sgpr-vgpr-vgpr", {0xCC7E4004u, 0x1C421808u, 0u}, {87u, 91u}}, SourceFormCase{"vgpr-null-vgpr", {0xCC7E4004u, 0x1C40F908u, 0u}, {9u, 14u}}, SourceFormCase{"vgpr-inline-vgpr", {0xCC7E4004u, 0x1C410308u, 0u}, {11u, 17u}}, @@ -1324,6 +1369,48 @@ TEST(Gfx1251PackedU64ExecutionTest, ExecutesEveryPublicLlvmSourceForm) { } } +TEST(Gfx1251PackedU64ExecutionTest, SplitsVccAndExecShiftSourcesIntoLowAndHighDwords) { + struct SpecialSourceCase { + std::string_view name; + std::array words; + bool use_exec; + }; + // Public LLVM MC accepts both special-register forms for gfx1251: + // v_pk_lshl_add_u64 v[4:7], v[8:11], vcc, v[16:19] + // v_pk_lshl_add_u64 v[4:7], v[8:11], exec, v[16:19] + constexpr std::array kCases{ + SpecialSourceCase{"vcc", {0xCC7E4004u, 0x1C40D508u}, false}, + SpecialSourceCase{"exec", {0xCC7E4004u, 0x1C40FD08u}, true}, + }; + + auto decoder = + make_isa_decoder(&cdna5::execution_backend(), cdna5::kGfx1251IsaFeatures); + ASSERT_NE(decoder, nullptr); + Gfx1250Sim sim; + auto *cu = sim.cu(); + auto *wf = cu->dispatch_wf(0, 0, kGfx1250ScalarSlots, 32); + ASSERT_NE(wf, nullptr); + + for (const auto &test_case : kCases) { + SCOPED_TRACE(test_case.name); + wf->set_exec_raw(1u); + wf->set_vcc_raw(0u); + if (test_case.use_exec) + wf->set_exec_raw(0x0000000200000001ULL); + else + wf->set_vcc_raw(0x0000000200000001ULL); + write_vgpr_packed_u64(*cu, *wf, 8, 0, {2u, 3u}); + write_vgpr_packed_u64(*cu, *wf, 16, 0, {7u, 11u}); + write_vgpr_packed_u64(*cu, *wf, 4, 0, {0u, 0u}); + + std::unique_ptr decoded(decode_valid(*decoder, test_case.words.data())); + ASSERT_NE(decoded, nullptr); + ASSERT_NE(decoded->execute, nullptr); + cu->execute_instruction(decoded.get(), *wf); + EXPECT_EQ(read_vgpr_packed_u64(*cu, *wf, 4, 0), (PackedU64Pair{11u, 23u})); + } +} + TEST(Gfx1251PackedU64ExecutionTest, ReadsOverlappingSourcesBeforeWritingDestination) { const auto words = cdna5::build_vop3p(cdna5::kVPkLshlAddU64Vop3p, {.vdst = 8, .src0 = 264, .src1 = 268, .src2 = 272}); From 5d32be46e8733c7610dd1cae04624f29ba946350 Mon Sep 17 00:00:00 2001 From: hanhanW Date: Tue, 1 Sep 2026 14:36:23 -0700 Subject: [PATCH 3/3] [rocjitsu] Retire failed instruction dispatches. Record a nonzero SimulationEngine exit before resetting a wavefront that rejects an unsupported operand value. Retire the wave through the normal workgroup completion path so the command processor cannot retain the dispatch after its final wave is gone. Exercise the scheduler path with the public gfx1251 literal-101 encoding and verify both the simulator failure and AQL completion signal. Keep gfx1251 full-simulation capability disabled. Signed-off-by: hanhanW --- .../src/rocjitsu/vm/amdgpu/compute_unit.cpp | 22 ++++++++-- .../src/rocjitsu/vm/amdgpu/compute_unit.h | 13 ++++++ .../rocjitsu/tests/cdna5_execution_test.cpp | 43 +++++++++++++++++++ 3 files changed, 74 insertions(+), 4 deletions(-) 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 8e2d27372f96..70fca7854855 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 @@ -60,6 +60,16 @@ constexpr uint32_t kPrivilegedStatusBit = 1u << 5; bool is_privileged(const Wavefront &wf) { return (wf.status_raw() & kPrivilegedStatusBit) != 0; } +std::string_view instruction_execution_error_name(InstructionExecutionError error) { + switch (error) { + case InstructionExecutionError::None: + return "none"; + case InstructionExecutionError::UnsupportedOperandValue: + return "unsupported operand value"; + } + return "unknown instruction execution error"; +} + uint32_t pack_barrier_state(uint32_t member_count, uint32_t signal_count, uint32_t allocation_blocks = 0) { return 1u | ((member_count & 0x7fu) << 4) | ((signal_count & 0x7fu) << 16) | @@ -885,10 +895,14 @@ void ComputeUnitCore::issue_instruction(Wavefront *active) { execute_instruction(inst, *active); if (active->instruction_execution_failed()) { - util::Logger::warn("CU ", this->name(), ": wf", active->wf_id(), - " halted after unsupported operand value in ", inst->mnemonic(), " at pc=0x", - std::hex, active->pc); - active->halt(Wavefront::CpCompletionNotice::Suppress); + const InstructionExecutionError error = active->instruction_execution_error(); + const std::string failure = std::format("CU {}: wf{} could not execute {} at pc={:#x}: {}", + this->name(), active->wf_id(), inst->mnemonic(), + active->pc, instruction_execution_error_name(error)); + util::Logger::warn(failure); + if (auto *sim_engine = this->engine()) + sim_engine->request_exit(failure, /*code=*/1); + active->halt(); delete inst; return; } 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 a5d40037967c..371a5b001743 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 @@ -840,6 +840,19 @@ class ComputeUnitCore : public simdojo::CompositeComponent { /// @returns Const pointer to the ISA decoder. const Decoder *decoder() const { return decoder_.get(); } + /// @brief Replace the concrete decoder in scheduler-level tests. + /// @details Some concrete targets intentionally remain unavailable to full + /// simulator configuration while individual instructions acquire execution + /// support. This seam lets a production-topology test exercise fetch, decode, + /// issue, and completion for such an instruction without weakening target + /// capability validation. The CU must be idle when its decoder is replaced. + void replace_decoder_for_test(std::unique_ptr decoder) { + assert(decoder != nullptr); + assert(!has_active_wfs()); + decoder->enable_pool(); + decoder_ = std::move(decoder); + } + /// @brief Return the completer port (receives dispatch requests from CP). /// @returns Pointer to the completer port. simdojo::Port *cpl_port() { return cpl_; } diff --git a/emulation/rocjitsu/tests/cdna5_execution_test.cpp b/emulation/rocjitsu/tests/cdna5_execution_test.cpp index 0f0003d72abb..d409f6bd8b1e 100644 --- a/emulation/rocjitsu/tests/cdna5_execution_test.cpp +++ b/emulation/rocjitsu/tests/cdna5_execution_test.cpp @@ -1307,6 +1307,49 @@ TEST(Gfx1251PackedU64ExecutionTest, RejectsUnprovenShiftCountsBeforeAnyDestinati } } +TEST(Gfx1251PackedU64ExecutionTest, RejectedCountTerminatesDispatchWithSimulatorFailure) { + constexpr uint32_t kLiteral = 101u; + // Public LLVM gfx1251_asm_vop3p.s encoding for + // v_pk_lshl_add_u64 v[4:7], v[8:11], 101, v[16:19]. The s_endpgm must never + // be needed to retire the rejected instruction's workgroup. + constexpr std::array kCode{0xCC7E4004u, 0x1C41FF08u, kLiteral, S_ENDPGM_GFX12}; + constexpr uint64_t kSignalAddress = 0x20000; + constexpr uint32_t kSignalValueOffset = 8; + + Gfx1250Sim sim; + auto decoder = + make_isa_decoder(&cdna5::execution_backend(), cdna5::kGfx1251IsaFeatures); + ASSERT_NE(decoder, nullptr); + sim.cu()->replace_decoder_for_test(std::move(decoder)); + + sim.memory->write64(kSignalAddress + kSignalValueOffset, 1u); + const uint64_t kernel_object = sim.write_kernel(0x10000, kCode.data(), kCode.size()); + test::AqlQueue queue(sim.memory, sim.cp()); + hsa_kernel_dispatch_packet_t packet{}; + packet.header = HSA_PACKET_TYPE_KERNEL_DISPATCH; + packet.setup = 1; + packet.workgroup_size_x = 32; + packet.workgroup_size_y = 1; + packet.workgroup_size_z = 1; + packet.grid_size_x = 32; + packet.grid_size_y = 1; + packet.grid_size_z = 1; + packet.kernel_object = kernel_object; + packet.completion_signal.handle = kSignalAddress; + queue.submit(packet); + + while (sim.engine->step()) { + } + + const simdojo::ExitStatus &exit = sim.engine->last_exit(); + EXPECT_EQ(exit.reason, simdojo::ExitReason::EXIT_REQUEST); + EXPECT_EQ(exit.code, 1); + EXPECT_NE(exit.message.find("v_pk_lshl_add_u64"), std::string::npos); + EXPECT_FALSE(sim.cu()->has_active_wfs()); + EXPECT_EQ(sim.memory->read64(kSignalAddress + kSignalValueOffset), 0u) + << "the rejected dispatch must still publish its workgroup completion"; +} + TEST(Gfx1251PackedU64ExecutionTest, ExecutesEveryPublicLlvmSourceForm) { struct SourceFormCase { std::string_view name;