From 43aaf51eceff0eeda79c166b65c79dcd879cd171 Mon Sep 17 00:00:00 2001 From: Chris Millette Date: Sun, 1 Nov 2026 00:01:28 -0400 Subject: [PATCH 1/7] fix(rocke): refuse buffer_rsrc operands its lowerings cannot take buffer_rsrc lowers to the p1 overload of llvm.amdgcn.make.buffer.rsrc, which names its pointer ptr addrspace(1). A pointer parameter declared addr_space="constant" is ptr addrspace(4) in the kernel header, so the call did not match the header and llc rejected the module. The LLVM lowering of both engines now refuses it with the text the global-pointer ops already use: buffer_rsrc ptr must be a global pointer, got ptr addrspace(4) (its parameter's addr_space) At llvm20 the intrinsic takes an i32 num_records, and the lowering printed any num_bytes operand as i32, so an i64 operand did not match its own definition. Both engines now refuse a num_bytes that is not i32 at llvm20: tile.buffer_rsrc num_bytes must be i32 at llvm20, got i64; the llvm20 make.buffer.rsrc takes an i32 num_records, so pass the byte count as i32 llvm22 and later still take an i32 (zero-extended) or an i64. The HIP lowering calls __builtin_amdgcn_make_buffer_rsrc, whose num_records is an int, and cast num_bytes to (int): an i64 byte count was silently truncated to 32 bits, at every LLVM version. The HIP lowering of both engines now refuses a num_bytes that is not i32: tile.buffer_rsrc num_bytes must be i32 in the HIP lowering, got i64; the HIP builtin takes an int num_records, so the byte count would be truncated at every LLVM version; pass it as i32 The HIP lowering keeps a constant-space pointer: its kernel signature declares every pointer parameter the same way whatever its addr_space, and the builtin takes a void*, so the source compiles and the descriptor is built from the right address. No IR or HIP source changes for a kernel that lowered before. Tests: tests/core/test_buffer_rsrc_operands.py checks every refusal in both engines with the exact text, that the accepted forms lower to the same IR in both engines at every flavor, and that both HIP lowerings keep an i32 num_bytes and a constant-space pointer; the C++ test future_intrinsic_lowering gains the same cases. op_vocabulary.md states the rules. Co-Authored-By: Claude --- .../cpp/core/lower_hip/lower_hip_mem.cpp | 17 +++ .../platform/cpp/core/lower_llvm/mem.cpp | 16 +++ .../dsl_docs/reference/op_vocabulary.md | 8 ++ .../platform/python/rocke/core/lower_hip.py | 13 ++ .../platform/python/rocke/core/lower_llvm.py | 19 ++- .../tests/core/future_intrinsic_lowering.cpp | 123 ++++++++++++++++ .../tests/core/test_buffer_rsrc_operands.py | 133 ++++++++++++++++++ 7 files changed, 327 insertions(+), 2 deletions(-) create mode 100644 dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_operands.py diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp index d392495f9949..f420e21518f0 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp +++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp @@ -1074,6 +1074,23 @@ static rocke_status_t _op_tile_buffer_rsrc(rocke_h_lowerer_t* lw, const rocke_op } ptr = op->operands[0]; num_bytes = op->operands[1]; + /* num_records is an int, so the (int) cast would truncate an i64 byte + * count at every LLVM version. A constant-space ptr is kept: the HIP + * signature declares every pointer parameter the same way, and the builtin + * takes a void*. */ + { + const char* nb_ty + = (num_bytes->type && num_bytes->type->name) ? num_bytes->type->name : "(null)"; + if(strcmp(nb_ty, "i32") != 0) + { + return rocke_h_fail(lw, + ROCKE_ERR_VALUE, + "tile.buffer_rsrc num_bytes must be i32 in the HIP lowering, got " + "%s; the HIP builtin takes an int num_records, so the byte count " + "would be truncated at every LLVM version; pass it as i32", + nb_ty); + } + } /* The flags word is the rsrc DWORD3, whose format / out-of-bounds encoding * is ISA-specific. Take it from the target's ISA backend, the same source * the LLVM lowerer's tile.buffer_rsrc reads, so both routes agree. */ diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mem.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mem.cpp index b4a59605c3f1..4846a6057731 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mem.cpp +++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mem.cpp @@ -1016,6 +1016,10 @@ static void op_tile_buffer_rsrc(rocke_lower_t* L, const rocke_op_t* op) const char* intrinsic; const char* nb_text; int word3; + /* Python _op_tile_buffer_rsrc. The p1 overload names ptr addrspace(1), so + * a constant-space parameter (ptr addrspace(4) in the header) is refused. */ + if(!ll_require_global_header_ptr(L, "buffer_rsrc", ptr)) + return; rocke_ll_need(L, "make.buffer.rsrc.p1"); if(rocke_ll_flavor_is_modern(L->flavor)) { @@ -1040,7 +1044,19 @@ static void op_tile_buffer_rsrc(rocke_lower_t* L, const rocke_op_t* op) } else { + /* llvm20's intrinsic takes an i32 num_records; any other operand + * printed as i32 would not match its own definition. */ + const char* nb_ty = rocke_ll_llvm_type(L, num_bytes->type); intrinsic = "llvm.amdgcn.make.buffer.rsrc.p1"; + if(strcmp(nb_ty, "i32") != 0) + { + rocke_ll_fail( + L, + ROCKE_ERR_VALUE, + "tile.buffer_rsrc num_bytes must be i32 at llvm20, got %s; the llvm20 " + "make.buffer.rsrc takes an i32 num_records, so pass the byte count as i32", + nb_ty); + } nb_text = rocke_arena_printf(&L->arena, "i32 %s", rocke_ll_operand(L, num_bytes)); } word3 = L->backend->buffer_rsrc_word3; diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/op_vocabulary.md b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/op_vocabulary.md index bba2458e0706..25b8198140d9 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/op_vocabulary.md +++ b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/op_vocabulary.md @@ -154,6 +154,14 @@ access of 4 bytes or more per 4-byte component and a 1- or 2-byte access whole: `num_records` reads 0, a store past it writes nothing. The HIP lowering has no vector type for u8 or i64 elements, so those two lower through LLVM only. +`buffer_rsrc` takes a global pointer: the intrinsic's overload names `ptr addrspace(1)`, so a +parameter declared `addr_space="constant"` (`ptr addrspace(4)` in the kernel header) is refused by +the LLVM lowering of both engines, as the `global_ptr_*` ops refuse it. `num_bytes` is i32 or i64 at +llvm22 and later (an i32 is zero-extended); llvm20's intrinsic takes an i32 `num_records`, so an i64 +`num_bytes` is refused there. The HIP lowering's `__builtin_amdgcn_make_buffer_rsrc` takes an `int` +`num_records` at every LLVM version, so it refuses any `num_bytes` but i32; it keeps a constant-space +pointer, since the HIP signature declares every pointer parameter alike and the builtin takes a `void*`. + ## MFMA | Op name | Operands (per-lane) | Output (per-lane) | diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py index bf852c458734..fd3b64dd6616 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py +++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py @@ -2414,6 +2414,19 @@ def _op_tile_buffer_rsrc(self, op: Op) -> None: # the right word, out-of-range byte offsets return 0 on load and are # dropped on store; without it the compiler can lower buffer loads to # unbounded single-dword loads with no bounds check. + # + # ``num_records`` is an ``int``, so the ``(int)`` cast would truncate an + # i64 byte count at every LLVM version; anything but i32 is refused. + # The C++ engine refuses with the same text. A constant-space ``ptr`` + # is kept: the HIP signature declares every pointer parameter the same + # way whatever its ``addr_space``, and the builtin takes a ``void*``. + nb_ty = op.operands[1].type.name + if nb_ty != "i32": + raise ValueError( + f"tile.buffer_rsrc num_bytes must be i32 in the HIP lowering, got " + f"{nb_ty}; the HIP builtin takes an int num_records, so the byte " + "count would be truncated at every LLVM version; pass it as i32" + ) ptr, num_bytes = op.operands word3 = backend_for(self.arch).buffer_rsrc_word3 self._emit( diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py index 480ddd096c24..e0d26d26cf0d 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py +++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py @@ -5978,15 +5978,23 @@ def _op_tile_buffer_rsrc(self, op: Op) -> None: On the modern-flavor path we accept either an ``i32`` or an ``i64`` ``num_bytes`` operand and ``zext`` ``i32`` callers up; ``i64`` callers reach the full 64-bit range needed for >4 GiB KV - caches that would otherwise OOB-zero at the tail. + caches that would otherwise OOB-zero at the tail. LLVM 20 takes only + an ``i32``, so any other ``num_bytes`` is refused there: printed as + the ``i32`` operand, it would not match its own definition. + + The intrinsic is the ``p1`` overload, so ``ptr`` must be + ``ptr addrspace(1)`` in the kernel header. A parameter declared + ``addr_space="constant"`` is ``ptr addrspace(4)`` there, and is + refused as the global-pointer ops refuse it. Flags = the target ISA backend's ``buffer_rsrc_word3`` -- the rsrc DWORD3 that makes the descriptor bounds-checked. Its encoding is ISA-specific (``0x00027000`` on gfx9, ``0x31014000`` on gfx11 / gfx12); the HIP lowering takes the same word from the same backend. """ - self._need("make.buffer.rsrc.p1") ptr, num_bytes = op.operands + self._require_global_header_ptr("buffer_rsrc", ptr) + self._need("make.buffer.rsrc.p1") if _is_modern_flavor(self._flavor): intrinsic = "llvm.amdgcn.make.buffer.rsrc.p8.p1" nb_ty = _llvm_type(num_bytes.type) @@ -6004,6 +6012,13 @@ def _op_tile_buffer_rsrc(self, op: Op) -> None: nb_text = f"i64 {nb_arg}" else: intrinsic = "llvm.amdgcn.make.buffer.rsrc.p1" + nb_ty = _llvm_type(num_bytes.type) + if nb_ty != "i32": + raise ValueError( + f"tile.buffer_rsrc num_bytes must be i32 at llvm20, got {nb_ty}; " + "the llvm20 make.buffer.rsrc takes an i32 num_records, so pass " + "the byte count as i32" + ) nb_text = f"i32 {self._operand(num_bytes)}" word3 = self._backend.buffer_rsrc_word3 self._current().emit( diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/future_intrinsic_lowering.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/future_intrinsic_lowering.cpp index 9f0331c81f5a..2c4fa49916fd 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/future_intrinsic_lowering.cpp +++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/future_intrinsic_lowering.cpp @@ -1440,6 +1440,33 @@ rocke_status_t lower_expect_error(const char* name, return st; } +/* The HIP-source counterpart of lower_expect_error: the HIP lowerer leaves its + * message on the builder's error channel. */ +template +rocke_status_t + lower_hip_expect_error(const char* name, BuildFn build, const char* arch, std::string* err_out) +{ + rocke_ir_builder_t b; + if(rocke_ir_builder_init(&b, name) != ROCKE_OK) + { + fail("rocke_ir_builder_init", __LINE__); + return ROCKE_OK; + } + build(&b); + rocke_b_ret(&b); + rocke_strbuf_t out; + rocke_strbuf_init(&out, 0); + rocke_lower_hip_opts_t opts{}; + opts.arch = arch; + const rocke_status_t st + = rocke_lower_kernel_to_hip(&b, rocke_ir_builder_kernel(&b), &opts, &out); + const char* msg = rocke_ir_builder_error(&b); + err_out->assign(msg ? msg : ""); + rocke_strbuf_free(&out); + rocke_ir_builder_free(&b); + return st; +} + rocke_value_t* private_ptr_param(rocke_ir_builder_t* b, const char* name, const rocke_type_t* elem) { return rocke_b_param(b, name, rocke_ptr_type(b, elem, "private"), nullptr); @@ -3296,6 +3323,101 @@ void case_buffer_load_lds_async() EXPECT_IR(ir, "ptr addrspace(3) %lds_ptr"); } +/* ---- buffer_rsrc: pointer address space and num_bytes width ---- */ + +/* A buffer_rsrc over `X` with a `num_bytes` parameter of type `nb_type`; + * `X` is declared addr_space="constant" when `constant` is set. */ +static void build_buffer_rsrc(rocke_ir_builder_t* b, bool constant, const rocke_type_t* nb_type) +{ + rocke_param_opts_t opts{}; + if(constant) + opts.addr_space = "constant"; + rocke_value_t* X = rocke_b_param(b, "X", rocke_ptr_type(b, rocke_f32(), "global"), &opts); + rocke_value_t* n = rocke_b_param(b, "n", nb_type, nullptr); + rocke_value_t* rsrc = rocke_b_buffer_rsrc(b, X, n); + rocke_b_buffer_load_vN( + b, rsrc, rocke_b_const_i32(b, 0), rocke_b_const_i32(b, 0), rocke_f32(), 1); +} + +void case_buffer_rsrc_refusals() +{ + /* The make.buffer.rsrc overload names ptr addrspace(1); a constant-space + * parameter is ptr addrspace(4) in the header, at every flavor. */ + const rocke_llvm_flavor_t flavors[] = {ROCKE_LLVM_FLAVOR_LLVM20, + ROCKE_LLVM_FLAVOR_LLVM22, + ROCKE_LLVM_FLAVOR_LLVM23, + ROCKE_LLVM_FLAVOR_LLVM24}; + for(const rocke_llvm_flavor_t flavor : flavors) + { + std::string err; + const rocke_status_t st = lower_expect_error( + "rsrc_constant", + [](rocke_ir_builder_t* b) { build_buffer_rsrc(b, true, rocke_i32()); }, + "gfx942", + flavor, + &err); + if(st != ROCKE_ERR_VALUE + || err.find("buffer_rsrc ptr must be a global pointer, got ptr addrspace(4) " + "(its parameter's addr_space)") + == std::string::npos) + fail(("constant-space buffer_rsrc: " + err).c_str(), __LINE__); + } + + /* llvm20's make.buffer.rsrc.p1 takes an i32 num_records. */ + std::string err; + const rocke_status_t st = lower_expect_error( + "rsrc_i64_llvm20", + [](rocke_ir_builder_t* b) { build_buffer_rsrc(b, false, rocke_i64()); }, + "gfx942", + ROCKE_LLVM_FLAVOR_LLVM20, + &err); + if(st != ROCKE_ERR_VALUE + || err.find("tile.buffer_rsrc num_bytes must be i32 at llvm20, got i64; the llvm20 " + "make.buffer.rsrc takes an i32 num_records, so pass the byte count as i32") + == std::string::npos) + fail(("i64 num_bytes at llvm20: " + err).c_str(), __LINE__); + + /* The accepted forms still lower: i32 at llvm20, i64 (as is) at llvm22. */ + const std::string ir20 = lower_one( + "rsrc_i32_llvm20", + [](rocke_ir_builder_t* b) { build_buffer_rsrc(b, false, rocke_i32()); }, + "gfx942", + ROCKE_LLVM_FLAVOR_LLVM20); + EXPECT_IR(ir20, "@llvm.amdgcn.make.buffer.rsrc.p1(ptr addrspace(1) %X, i16 0, i32 %n, "); + const std::string ir22 = lower_one( + "rsrc_i64_llvm22", + [](rocke_ir_builder_t* b) { build_buffer_rsrc(b, false, rocke_i64()); }, + "gfx942", + ROCKE_LLVM_FLAVOR_LLVM22); + EXPECT_IR(ir22, "@llvm.amdgcn.make.buffer.rsrc.p8.p1(ptr addrspace(1) %X, i16 0, i64 %n, "); + + /* The HIP builtin takes an int num_records, so the HIP lowering refuses an + * i64 num_bytes whatever the LLVM version. */ + std::string hip_err; + const rocke_status_t hip_st = lower_hip_expect_error( + "rsrc_i64_hip", + [](rocke_ir_builder_t* b) { build_buffer_rsrc(b, false, rocke_i64()); }, + "gfx942", + &hip_err); + if(hip_st != ROCKE_ERR_VALUE + || hip_err.find("tile.buffer_rsrc num_bytes must be i32 in the HIP lowering, got i64; " + "the HIP builtin takes an int num_records, so the byte count would be " + "truncated at every LLVM version; pass it as i32") + == std::string::npos) + fail(("i64 num_bytes in HIP: " + hip_err).c_str(), __LINE__); + + /* The HIP signature declares a constant-space pointer like any other, and + * the builtin takes a void*, so the HIP lowering keeps it. */ + const std::string hip = lower_one_hip( + "rsrc_constant_hip", + [](rocke_ir_builder_t* b) { build_buffer_rsrc(b, true, rocke_i32()); }, + "gfx942"); + EXPECT_IR(hip, "void rsrc_constant_hip(float* __restrict__ X, int n)"); + EXPECT_IR(hip, + "__builtin_amdgcn_make_buffer_rsrc((void*)X, /*stride=*/(short)0, " + "/*num_records=*/(int)n, "); +} + /* ---- dtype-generic buffer_load_vN / buffer_store_vN (Python IRBuilder) ---- */ /* Run `build` on a fresh builder; return the refusal text, or "" when the @@ -4077,6 +4199,7 @@ const TestCase k_cases[] = { {"memory_caps", case_memory_caps}, {"vmcnt_bits", case_vmcnt_bits}, {"buffer_load_lds_async", case_buffer_load_lds_async}, + {"buffer_rsrc_refusals", case_buffer_rsrc_refusals}, {"buffer_load_vN_generic", case_buffer_load_vN_generic}, {"buffer_store_vN_generic", case_buffer_store_vN_generic}, {"buffer_vN_verify", case_buffer_vN_verify}, diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_operands.py b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_operands.py new file mode 100644 index 000000000000..0ad1407f05be --- /dev/null +++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_operands.py @@ -0,0 +1,133 @@ +# Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +# SPDX-License-Identifier: MIT + +"""The operands ``buffer_rsrc`` refuses in the LLVM and HIP lowerings of both engines. + +``buffer_rsrc(ptr, num_bytes)`` lowers to ``@llvm.amdgcn.make.buffer.rsrc``, +whose ``p1`` overload names ``ptr addrspace(1)``: + +- a pointer parameter declared ``addr_space="constant"`` is + ``ptr addrspace(4)`` in the kernel header, so the call would not match the + header. It is refused at every flavor, with the text the global-pointer ops + use. +- llvm20's intrinsic takes an i32 ``num_records``; an i64 ``num_bytes`` would be + printed as an i32 operand that does not match its own definition. It is + refused at llvm20, with the fix in the message. llvm22 and later take an i64 + and zero-extend an i32, so both widths still lower there. + +Before these refusals both engines emitted the call anyway, and an LLVM 20 +``llc`` rejects the module ("'%X' defined with type 'ptr addrspace(4)' but +expected 'ptr addrspace(1)'"; "'%n' defined with type 'i64' but expected +'i32'"). + +The HIP lowering calls ``__builtin_amdgcn_make_buffer_rsrc``, whose +``num_records`` is an ``int``: it cast an i64 ``num_bytes`` to ``(int)``, which +truncates the byte count at every LLVM version, so it refuses any width but i32. +It keeps a constant-space pointer: the HIP signature declares every pointer +parameter the same way, and the builtin takes a ``void*``. +""" + +from __future__ import annotations + +import pytest +from rocke.core.ir import F32, I32, I64, IRBuilder, PtrType +from rocke.tests.engines import ( + assert_both_engines_reject, + assert_both_engines_reject_hip, + lower_hip_with_both_engines, + lower_with_both_engines, +) + +FLAVORS = ("llvm20", "llvm22", "llvm23", "llvm24") +ARCH = "gfx942" + + +def _rsrc_kernel(name, *, constant, num_bytes_type): + """A buffer_rsrc over ``X`` with a ``num_bytes`` parameter of + ``num_bytes_type``, and one load through it; ``X`` is declared + ``addr_space="constant"`` when ``constant`` is set.""" + b = IRBuilder(name) + space = {"addr_space": "constant"} if constant else {} + x = b.param("X", PtrType(F32, "global"), **space) + y = b.param("Y", PtrType(F32, "global")) + n = b.param("n", num_bytes_type) + rsrc = b.buffer_rsrc(x, n) + value = b.buffer_load_vN(rsrc, b.const_i32(0), b.const_i32(0), F32, 1) + b.global_store(y, b.const_i32(0), b.vec_extract(value, 0)) + b.ret() + return b.kernel + + +@pytest.mark.parametrize("flavor", FLAVORS) +def test_both_llvm_lowerings_refuse_a_constant_space_pointer(flavor): + kernel = _rsrc_kernel("rsrc_constant", constant=True, num_bytes_type=I32) + assert_both_engines_reject( + kernel, + ARCH, + TypeError, + "buffer_rsrc ptr must be a global pointer, got ptr addrspace(4) " + "(its parameter's addr_space)", + llvm_flavor=flavor, + ) + + +def test_both_llvm_lowerings_refuse_an_i64_num_bytes_at_llvm20(): + kernel = _rsrc_kernel("rsrc_i64", constant=False, num_bytes_type=I64) + assert_both_engines_reject( + kernel, + ARCH, + ValueError, + "tile.buffer_rsrc num_bytes must be i32 at llvm20, got i64; the llvm20 " + "make.buffer.rsrc takes an i32 num_records, so pass the byte count as i32", + llvm_flavor="llvm20", + ) + + +@pytest.mark.parametrize("flavor", FLAVORS) +def test_an_i32_num_bytes_lowers_at_every_flavor(flavor): + kernel = _rsrc_kernel("rsrc_i32", constant=False, num_bytes_type=I32) + ir = lower_with_both_engines(kernel, ARCH, llvm_flavor=flavor) + if flavor == "llvm20": + assert ( + "@llvm.amdgcn.make.buffer.rsrc.p1(ptr addrspace(1) %X, i16 0, i32 %n," in ir + ) + else: + assert "= zext i32 %n to i64" in ir + assert ( + "@llvm.amdgcn.make.buffer.rsrc.p8.p1(ptr addrspace(1) %X, i16 0, i64 %" + in ir + ) + + +@pytest.mark.parametrize("flavor", FLAVORS[1:]) +def test_an_i64_num_bytes_lowers_after_llvm20(flavor): + kernel = _rsrc_kernel("rsrc_i64", constant=False, num_bytes_type=I64) + ir = lower_with_both_engines(kernel, ARCH, llvm_flavor=flavor) + assert ( + "@llvm.amdgcn.make.buffer.rsrc.p8.p1(ptr addrspace(1) %X, i16 0, i64 %n," in ir + ) + + +def test_both_hip_lowerings_refuse_an_i64_num_bytes(): + kernel = _rsrc_kernel("rsrc_i64", constant=False, num_bytes_type=I64) + assert_both_engines_reject_hip( + kernel, + ARCH, + ValueError, + "tile.buffer_rsrc num_bytes must be i32 in the HIP lowering, got i64; " + "the HIP builtin takes an int num_records, so the byte count would be " + "truncated at every LLVM version; pass it as i32", + ) + + +@pytest.mark.parametrize("constant", (False, True)) +def test_the_hip_lowering_keeps_an_i32_num_bytes_and_a_constant_space_pointer( + constant, +): + kernel = _rsrc_kernel("rsrc_hip", constant=constant, num_bytes_type=I32) + hip = lower_hip_with_both_engines(kernel, ARCH) + assert "void rsrc_hip(float* __restrict__ X, float* __restrict__ Y, int n)" in hip + assert ( + "__builtin_amdgcn_make_buffer_rsrc((void*)X, /*stride=*/(short)0, " + "/*num_records=*/(int)n, " in hip + ) From 61a805c511573c686c74b2adee5f7be6105af698 Mon Sep 17 00:00:00 2001 From: Chris Millette Date: Tue, 3 Nov 2026 01:29:21 -0500 Subject: [PATCH 2/7] fix(rocke): blame the (int) cast in the HIP buffer_rsrc refusal The HIP lowering refuses a buffer_rsrc num_bytes wider than i32. Its message and comments said the builtin takes an int num_records at every LLVM version, which is not true: newer clang declares num_records as a uint64_t. What truncates a wider byte count is the (int) cast the lowering emits, so the message, the comments and the op reference now say that, in both engines. The refusal itself is unchanged. Co-Authored-By: Claude --- .../platform/cpp/core/lower_hip/lower_hip_mem.cpp | 9 +++++---- .../platform/dsl_docs/reference/op_vocabulary.md | 7 ++++--- .../rocke/platform/python/rocke/core/lower_hip.py | 12 +++++++----- .../tests/core/future_intrinsic_lowering.cpp | 8 ++++---- .../platform/tests/core/test_buffer_rsrc_operands.py | 12 ++++++------ 5 files changed, 26 insertions(+), 22 deletions(-) diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp index f420e21518f0..611a250f491f 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp +++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp @@ -1074,8 +1074,9 @@ static rocke_status_t _op_tile_buffer_rsrc(rocke_h_lowerer_t* lw, const rocke_op } ptr = op->operands[0]; num_bytes = op->operands[1]; - /* num_records is an int, so the (int) cast would truncate an i64 byte - * count at every LLVM version. A constant-space ptr is kept: the HIP + /* num_records is cast to (int), which ROCm 7's builtin (int num_records) + * and later clang's (uint64_t num_records) both accept; the cast would + * truncate an i64 byte count. A constant-space ptr is kept: the HIP * signature declares every pointer parameter the same way, and the builtin * takes a void*. */ { @@ -1086,8 +1087,8 @@ static rocke_status_t _op_tile_buffer_rsrc(rocke_h_lowerer_t* lw, const rocke_op return rocke_h_fail(lw, ROCKE_ERR_VALUE, "tile.buffer_rsrc num_bytes must be i32 in the HIP lowering, got " - "%s; the HIP builtin takes an int num_records, so the byte count " - "would be truncated at every LLVM version; pass it as i32", + "%s; the HIP lowering casts num_records to (int), so a wider byte " + "count would be truncated; pass it as i32", nb_ty); } } diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/op_vocabulary.md b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/op_vocabulary.md index 25b8198140d9..2c1b6b1d30bf 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/op_vocabulary.md +++ b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/op_vocabulary.md @@ -158,9 +158,10 @@ i64 elements, so those two lower through LLVM only. parameter declared `addr_space="constant"` (`ptr addrspace(4)` in the kernel header) is refused by the LLVM lowering of both engines, as the `global_ptr_*` ops refuse it. `num_bytes` is i32 or i64 at llvm22 and later (an i32 is zero-extended); llvm20's intrinsic takes an i32 `num_records`, so an i64 -`num_bytes` is refused there. The HIP lowering's `__builtin_amdgcn_make_buffer_rsrc` takes an `int` -`num_records` at every LLVM version, so it refuses any `num_bytes` but i32; it keeps a constant-space -pointer, since the HIP signature declares every pointer parameter alike and the builtin takes a `void*`. +`num_bytes` is refused there. The HIP lowering passes `num_records` to +`__builtin_amdgcn_make_buffer_rsrc` cast to `(int)`, which would truncate a wider byte count, so it +refuses any `num_bytes` but i32. It keeps a constant-space pointer, since the HIP signature declares +every pointer parameter alike and the builtin takes a `void*`. ## MFMA diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py index fd3b64dd6616..9db4a7f7d073 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py +++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py @@ -2404,7 +2404,8 @@ def _op_tile_buffer_rsrc(self, op: Op) -> None: # # Signature on ROCm 7 amdclang: # ``__builtin_amdgcn_make_buffer_rsrc(void*, short stride, - # int num_records, int flags)`` + # int num_records, int flags)``; later clang takes a ``uint64_t`` + # ``num_records``. # # The flags word is the rsrc DWORD3. Its format / out-of-bounds # encoding is ISA-specific (the gfx9 value makes every raw buffer @@ -2415,8 +2416,9 @@ def _op_tile_buffer_rsrc(self, op: Op) -> None: # dropped on store; without it the compiler can lower buffer loads to # unbounded single-dword loads with no bounds check. # - # ``num_records`` is an ``int``, so the ``(int)`` cast would truncate an - # i64 byte count at every LLVM version; anything but i32 is refused. + # ``num_records`` is cast to ``(int)``, which either signature accepts; + # the cast would truncate an i64 byte count, so anything but i32 is + # refused. # The C++ engine refuses with the same text. A constant-space ``ptr`` # is kept: the HIP signature declares every pointer parameter the same # way whatever its ``addr_space``, and the builtin takes a ``void*``. @@ -2424,8 +2426,8 @@ def _op_tile_buffer_rsrc(self, op: Op) -> None: if nb_ty != "i32": raise ValueError( f"tile.buffer_rsrc num_bytes must be i32 in the HIP lowering, got " - f"{nb_ty}; the HIP builtin takes an int num_records, so the byte " - "count would be truncated at every LLVM version; pass it as i32" + f"{nb_ty}; the HIP lowering casts num_records to (int), so a wider " + "byte count would be truncated; pass it as i32" ) ptr, num_bytes = op.operands word3 = backend_for(self.arch).buffer_rsrc_word3 diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/future_intrinsic_lowering.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/future_intrinsic_lowering.cpp index 2c4fa49916fd..c8b1c6c23510 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/future_intrinsic_lowering.cpp +++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/future_intrinsic_lowering.cpp @@ -3391,8 +3391,8 @@ void case_buffer_rsrc_refusals() ROCKE_LLVM_FLAVOR_LLVM22); EXPECT_IR(ir22, "@llvm.amdgcn.make.buffer.rsrc.p8.p1(ptr addrspace(1) %X, i16 0, i64 %n, "); - /* The HIP builtin takes an int num_records, so the HIP lowering refuses an - * i64 num_bytes whatever the LLVM version. */ + /* The HIP lowering casts num_records to (int), so it refuses an i64 + * num_bytes whatever the LLVM version. */ std::string hip_err; const rocke_status_t hip_st = lower_hip_expect_error( "rsrc_i64_hip", @@ -3401,8 +3401,8 @@ void case_buffer_rsrc_refusals() &hip_err); if(hip_st != ROCKE_ERR_VALUE || hip_err.find("tile.buffer_rsrc num_bytes must be i32 in the HIP lowering, got i64; " - "the HIP builtin takes an int num_records, so the byte count would be " - "truncated at every LLVM version; pass it as i32") + "the HIP lowering casts num_records to (int), so a wider byte count " + "would be truncated; pass it as i32") == std::string::npos) fail(("i64 num_bytes in HIP: " + hip_err).c_str(), __LINE__); diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_operands.py b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_operands.py index 0ad1407f05be..7cda00da13e2 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_operands.py +++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_operands.py @@ -20,10 +20,10 @@ expected 'ptr addrspace(1)'"; "'%n' defined with type 'i64' but expected 'i32'"). -The HIP lowering calls ``__builtin_amdgcn_make_buffer_rsrc``, whose -``num_records`` is an ``int``: it cast an i64 ``num_bytes`` to ``(int)``, which -truncates the byte count at every LLVM version, so it refuses any width but i32. -It keeps a constant-space pointer: the HIP signature declares every pointer +The HIP lowering calls ``__builtin_amdgcn_make_buffer_rsrc`` with +``num_records`` cast to ``(int)`` (ROCm 7's builtin takes an ``int`` there, +later clang a ``uint64_t``). The cast would truncate an i64 ``num_bytes``, so +the HIP lowering refuses any width but i32. It keeps a constant-space pointer: the HIP signature declares every pointer parameter the same way, and the builtin takes a ``void*``. """ @@ -115,8 +115,8 @@ def test_both_hip_lowerings_refuse_an_i64_num_bytes(): ARCH, ValueError, "tile.buffer_rsrc num_bytes must be i32 in the HIP lowering, got i64; " - "the HIP builtin takes an int num_records, so the byte count would be " - "truncated at every LLVM version; pass it as i32", + "the HIP lowering casts num_records to (int), so a wider byte count " + "would be truncated; pass it as i32", ) From 738c9095fcc888f52ba63f8d696c2fb40d7430cf Mon Sep 17 00:00:00 2001 From: Chris Millette Date: Tue, 3 Nov 2026 01:29:34 -0500 Subject: [PATCH 3/7] fix(rocke): match global memory ops to a constant-space parameter's type A pointer parameter declared addr_space="constant" keeps its global IR type, but the kernel header types it ptr addrspace(4). The global loads, stores, atomics, flag ops, cooperative store and global/LDS copies named it ptr addrspace(1) anyway, so LLVM rejected the module. The global loads now name the pointer's own type, so a kernel can read through a constant-space parameter; a flat or LDS pointer, which they also named ptr addrspace(1), is refused. Every other op that addresses global memory through the pointer refuses a constant-space parameter in the LLVM lowering of both engines, as buffer_rsrc and the global_ptr ops already do. The C++ global_store_async_from_lds and global_load_tr16_b128 lowerings now check the header type as the Python engine does. The shared pointer-space resolver's message says "op" rather than "intrinsic", since plain loads use it too. Co-Authored-By: Claude --- .../platform/cpp/core/lower_llvm/core.cpp | 8 +- .../platform/cpp/core/lower_llvm/data.cpp | 10 + .../platform/cpp/core/lower_llvm/mem.cpp | 104 ++++--- .../cpp/include/rocke/lower_llvm_internal.h | 21 +- .../dsl_docs/reference/op_vocabulary.md | 6 + .../platform/python/rocke/core/lower_llvm.py | 52 +++- .../core/test_constant_space_global_ops.py | 278 ++++++++++++++++++ .../rocke/platform/tests/test_rocke.py | 4 +- 8 files changed, 411 insertions(+), 72 deletions(-) create mode 100644 dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_constant_space_global_ops.py diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp index 7a4a3b800852..7fa183c7c800 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp +++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp @@ -700,12 +700,8 @@ int rocke_ll_anyptr_space(rocke_lower_t* L, break; used += (size_t)n; } - rocke_ll_fail(L, - ROCKE_ERR_VALUE, - "%s: pointer operand is %s, but the intrinsic accepts only %s", - op, - ty, - list); + rocke_ll_fail( + L, ROCKE_ERR_VALUE, "%s: pointer operand is %s, but the op accepts only %s", op, ty, list); } const char* rocke_ll_smem_storage_type(rocke_lower_t* L, const rocke_type_t* smem) diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/data.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/data.cpp index f05efcd07436..0063b077b2c5 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/data.cpp +++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/data.cpp @@ -469,6 +469,16 @@ const rocke_ll_anyptr_space_t ROCKE_LL_AV_B128_PTR_TYPES[] = { const int ROCKE_LL_AV_B128_PTR_TYPES_COUNT = (int)(sizeof(ROCKE_LL_AV_B128_PTR_TYPES) / sizeof(ROCKE_LL_AV_B128_PTR_TYPES[0])); +/* A memref.global_load* is a plain load, which reads through a global pointer + * or through a parameter declared addr_space="constant". It names the + * pointer's own type, so no entry needs a declare. */ +const rocke_ll_anyptr_space_t ROCKE_LL_GLOBAL_LOAD_PTR_TYPES[] = { + {1, "ptr addrspace(1)"}, + {4, "ptr addrspace(4)"}, +}; +const int ROCKE_LL_GLOBAL_LOAD_PTR_TYPES_COUNT + = (int)(sizeof(ROCKE_LL_GLOBAL_LOAD_PTR_TYPES) / sizeof(ROCKE_LL_GLOBAL_LOAD_PTR_TYPES[0])); + /* ---------------------------------------------------------------------- */ /* Per-generation overrides. Each array holds only the entries whose text */ /* changed AT that declaration ABI generation; a generation inherits every */ diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mem.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mem.cpp index 4846a6057731..79bfccce0a17 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mem.cpp +++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mem.cpp @@ -108,40 +108,71 @@ static void op_tile_smem_alloc(rocke_lower_t* L, const rocke_op_t* op) (void)op; } +/* Python _require_global_header_ptr; the text is the Python engine's. The + * global-pointer ops emit against ptr addrspace(1); the IR rules check the IR + * type, but a parameter's addr_space override (such as "constant") moves its + * header type away from it. */ +static void + ll_require_global_header_ptr(rocke_lower_t* L, const char* op_name, const rocke_value_t* ptr) +{ + const char* header = rocke_ll_value_ptr_type(L, ptr); + if(strcmp(header, "ptr addrspace(1)") != 0) + rocke_ll_fail(L, + ROCKE_ERR_VALUE, + "%s ptr must be a global pointer, got %s (its parameter's addr_space)", + op_name, + header); +} + +/* Python _anyptr_space over _GLOBAL_LOAD_PTR_TYPES: the pointer type a global + * load names, a global pointer or a parameter declared addr_space="constant". */ +static const char* + ll_global_load_ptr_type(rocke_lower_t* L, const char* op_name, const rocke_value_t* ptr) +{ + const char* ptr_ty = NULL; + rocke_ll_anyptr_space(L, + op_name, + ptr, + ROCKE_LL_GLOBAL_LOAD_PTR_TYPES, + ROCKE_LL_GLOBAL_LOAD_PTR_TYPES_COUNT, + &ptr_ty); + return ptr_ty; +} + static void op_memref_global_load(rocke_lower_t* L, const rocke_op_t* op) { const rocke_value_t* ptr = op->operands[0]; const rocke_value_t* idx = op->operands[1]; + const char* ptr_ty = ll_global_load_ptr_type(L, "global_load", ptr); const char* gep = rocke_ll_fresh(L, "gep"); int64_t align = ll_attr_int(op, "align", 2); rocke_ll_emitf(L, - " %s = getelementptr inbounds half, ptr addrspace(1) %s, i32 %s", + " %s = getelementptr inbounds half, %s %s, i32 %s", gep, + ptr_ty, rocke_ll_operand(L, ptr), rocke_ll_operand(L, idx)); rocke_ll_emitf( - L, " %s = load half, ptr addrspace(1) %s, align %lld", ll_res(op), gep, (long long)align); + L, " %s = load half, %s %s, align %lld", ll_res(op), ptr_ty, gep, (long long)align); } static void op_memref_global_load_typed(rocke_lower_t* L, const rocke_op_t* op) { const rocke_value_t* ptr = op->operands[0]; const rocke_value_t* idx = op->operands[1]; + const char* ptr_ty = ll_global_load_ptr_type(L, "global_load_typed", ptr); const char* elem_ty = rocke_ll_llvm_type(L, op->results[0]->type); const char* gep = rocke_ll_fresh(L, "gep"); int64_t align = ll_attr_int(op, "align", 1); rocke_ll_emitf(L, - " %s = getelementptr inbounds %s, ptr addrspace(1) %s, i32 %s", + " %s = getelementptr inbounds %s, %s %s, i32 %s", gep, elem_ty, + ptr_ty, rocke_ll_operand(L, ptr), rocke_ll_operand(L, idx)); - rocke_ll_emitf(L, - " %s = load %s, ptr addrspace(1) %s, align %lld", - ll_res(op), - elem_ty, - gep, - (long long)align); + rocke_ll_emitf( + L, " %s = load %s, %s %s, align %lld", ll_res(op), elem_ty, ptr_ty, gep, (long long)align); } static void op_memref_global_store(rocke_lower_t* L, const rocke_op_t* op) @@ -149,6 +180,7 @@ static void op_memref_global_store(rocke_lower_t* L, const rocke_op_t* op) const rocke_value_t* ptr = op->operands[0]; const rocke_value_t* idx = op->operands[1]; const rocke_value_t* val = op->operands[2]; + ll_require_global_header_ptr(L, "global_store", ptr); const char* gep = rocke_ll_fresh(L, "gep"); int64_t align = ll_attr_int(op, "align", 2); rocke_ll_emitf(L, @@ -168,6 +200,7 @@ static void op_memref_global_store_typed(rocke_lower_t* L, const rocke_op_t* op) const rocke_value_t* ptr = op->operands[0]; const rocke_value_t* idx = op->operands[1]; const rocke_value_t* val = op->operands[2]; + ll_require_global_header_ptr(L, "global_store_typed", ptr); const char* elem_ty = rocke_ll_llvm_type(L, val->type); const char* gep = rocke_ll_fresh(L, "gep"); int64_t align = ll_attr_int(op, "align", 1); @@ -233,6 +266,7 @@ static void op_memref_global_flag_store(rocke_lower_t* L, const rocke_op_t* op) const rocke_value_t* ptr = op->operands[0]; const rocke_value_t* idx = op->operands[1]; const rocke_value_t* val = op->operands[2]; + ll_require_global_header_ptr(L, "global_flag_store", ptr); const char* scope = ll_sync_attr(L, op, "scope", k_ll_sync_scopes); const char* ordering = ll_sync_attr(L, op, "ordering", k_ll_flag_store_orderings); const char* gep = rocke_ll_fresh(L, "gep"); @@ -258,6 +292,7 @@ static void op_memref_global_flag_wait_eq(rocke_lower_t* L, const rocke_op_t* op const rocke_value_t* ptr = op->operands[0]; const rocke_value_t* idx = op->operands[1]; const rocke_value_t* expect = op->operands[2]; + ll_require_global_header_ptr(L, "global_flag_wait_eq", ptr); const char* scope = ll_sync_attr(L, op, "scope", k_ll_sync_scopes); const char* gep = rocke_ll_fresh(L, "gep"); rocke_ll_block_t* cur = rocke_ll_current(L); @@ -305,6 +340,7 @@ static void op_memref_global_atomic_add(rocke_lower_t* L, const rocke_op_t* op) bool is_f32; const char* rmw_op; const char* md = ""; + ll_require_global_header_ptr(L, "global_atomic_add", ptr); rocke_ll_emitf(L, " %s = getelementptr inbounds %s, ptr addrspace(1) %s, i32 %s", gep, @@ -337,6 +373,7 @@ static void op_memref_global_atomic_add_f32(rocke_lower_t* L, const rocke_op_t* const rocke_value_t* ptr = op->operands[0]; const rocke_value_t* idx = op->operands[1]; const rocke_value_t* val = op->operands[2]; + ll_require_global_header_ptr(L, "global_atomic_add_f32", ptr); const char* gep = rocke_ll_fresh(L, "gep"); rocke_ll_emitf(L, " %s = getelementptr inbounds float, ptr addrspace(1) %s, i32 %s", @@ -373,6 +410,7 @@ static void op_memref_global_atomic_add_pk_bf16(rocke_lower_t* L, const rocke_op const char* idx_ty = rocke_ll_llvm_type(L, idx->type); const char* gep = rocke_ll_fresh(L, "gep"); const char* ordering = ll_attr_str(op, "ordering", "monotonic"); + ll_require_global_header_ptr(L, "global_atomic_add_pk_bf16", ptr); rocke_ll_emitf(L, " %s = getelementptr inbounds bfloat, ptr addrspace(1) %s, %s %s", gep, @@ -397,6 +435,7 @@ static void op_memref_global_atomic_add_pk_f16(rocke_lower_t* L, const rocke_op_ const rocke_value_t* idx = op->operands[1]; const rocke_value_t* val = op->operands[2]; const char* gep; + ll_require_global_header_ptr(L, "global_atomic_add_pk_f16", ptr); rocke_ll_need(L, "global.atomic.fadd.v2f16"); gep = rocke_ll_fresh(L, "gep"); rocke_ll_emitf(L, @@ -451,24 +490,27 @@ static void op_memref_global_load_vN(rocke_lower_t* L, const rocke_op_t* op) { const rocke_value_t* ptr = op->operands[0]; const rocke_value_t* idx = op->operands[1]; + const char* ptr_ty = ll_global_load_ptr_type(L, "global_load_vN", ptr); int64_t vec = ll_attr_int(op, "vec", 0); const char* elem_ty = rocke_ll_llvm_type(L, op->results[0]->type->elem); const char* idx_ty = rocke_ll_llvm_type(L, idx->type); const char* gep = rocke_ll_fresh(L, "gep"); int64_t align; rocke_ll_emitf(L, - " %s = getelementptr inbounds %s, ptr addrspace(1) %s, %s %s", + " %s = getelementptr inbounds %s, %s %s, %s %s", gep, elem_ty, + ptr_ty, rocke_ll_operand(L, ptr), idx_ty, rocke_ll_operand(L, idx)); align = ll_attr_int(op, "align", vec * 2); rocke_ll_emitf(L, - " %s = load <%lld x %s>, ptr addrspace(1) %s, align %lld%s", + " %s = load <%lld x %s>, %s %s, align %lld%s", ll_res(op), (long long)vec, elem_ty, + ptr_ty, gep, (long long)align, ll_nontemporal_md(L, op)); @@ -479,6 +521,7 @@ static void op_memref_global_store_vN(rocke_lower_t* L, const rocke_op_t* op) const rocke_value_t* ptr = op->operands[0]; const rocke_value_t* idx = op->operands[1]; const rocke_value_t* val = op->operands[2]; + ll_require_global_header_ptr(L, "global_store_vN", ptr); int64_t vec = ll_attr_int(op, "vec", 0); const char* gep = rocke_ll_fresh(L, "gep"); const char* elem_ty = ll_is_vec(val->type) ? rocke_ll_llvm_type(L, val->type->elem) @@ -516,6 +559,7 @@ static void op_memref_cooperative_global_store(rocke_lower_t* L, const rocke_op_ const char* addrs_ty; const char* values_ty; int64_t i; + ll_require_global_header_ptr(L, "cooperative_global_store", ptr); if(!ll_is_vec(values->type)) { rocke_ll_fail(L, ROCKE_ERR_NOTIMPL, "cooperative_global_store requires vector values"); @@ -919,30 +963,11 @@ static void op_tile_smem_ptr_add(rocke_lower_t* L, const rocke_op_t* op) rocke_ll_operand(L, off)); } -/* Python _require_global_header_ptr; the text is the Python engine's. The - * global-pointer ops emit against ptr addrspace(1); the IR rules check the IR - * type, but a parameter's addr_space override (such as "constant") moves its - * header type away from it. Returns false after failing the lowering. */ -static bool - ll_require_global_header_ptr(rocke_lower_t* L, const char* op_name, const rocke_value_t* ptr) -{ - const char* header = rocke_ll_value_ptr_type(L, ptr); - if(strcmp(header, "ptr addrspace(1)") == 0) - return true; - rocke_ll_fail(L, - ROCKE_ERR_VALUE, - "%s ptr must be a global pointer, got %s (its parameter's addr_space)", - op_name, - header); - return false; -} - static void op_tile_global_ptr_add(rocke_lower_t* L, const rocke_op_t* op) { const rocke_value_t* ptr = op->operands[0]; const rocke_value_t* off = op->operands[1]; - if(!ll_require_global_header_ptr(L, "global_ptr_add", ptr)) - return; + ll_require_global_header_ptr(L, "global_ptr_add", ptr); const char* off_ty = rocke_ll_llvm_type(L, off->type); const char* off64; if(strcmp(off_ty, "i32") == 0) @@ -978,8 +1003,7 @@ static void op_tile_global_ptr_offset(rocke_lower_t* L, const rocke_op_t* op) rocke_ll_fail(L, ROCKE_ERR_VALUE, "%s", error); return; } - if(!ll_require_global_header_ptr(L, "global_ptr_offset", op->operands[0])) - return; + ll_require_global_header_ptr(L, "global_ptr_offset", op->operands[0]); rocke_ll_emitf(L, " %s = getelementptr i8, ptr addrspace(1) %s, i64 %s", ll_res(op), @@ -999,8 +1023,7 @@ static void op_tile_global_ptr_to_i64(rocke_lower_t* L, const rocke_op_t* op) return; } ptr = op->operands[0]; - if(!ll_require_global_header_ptr(L, "global_ptr_to_i64", ptr)) - return; + ll_require_global_header_ptr(L, "global_ptr_to_i64", ptr); rocke_ll_emitf( L, " %s = ptrtoint ptr addrspace(1) %s to i64", ll_res(op), rocke_ll_operand(L, ptr)); } @@ -1018,8 +1041,7 @@ static void op_tile_buffer_rsrc(rocke_lower_t* L, const rocke_op_t* op) int word3; /* Python _op_tile_buffer_rsrc. The p1 overload names ptr addrspace(1), so * a constant-space parameter (ptr addrspace(4) in the header) is refused. */ - if(!ll_require_global_header_ptr(L, "buffer_rsrc", ptr)) - return; + ll_require_global_header_ptr(L, "buffer_rsrc", ptr); rocke_ll_need(L, "make.buffer.rsrc.p1"); if(rocke_ll_flavor_is_modern(L->flavor)) { @@ -1684,6 +1706,7 @@ static void op_tile_global_load_async_to_lds(rocke_lower_t* L, const rocke_op_t* const char* idx_ty; rocke_strbuf_t gep; int i; + ll_require_global_header_ptr(L, "global_load_async_to_lds", src_ptr); if(width == 1) suffix = "b8"; @@ -1768,8 +1791,7 @@ static void op_tile_global_store_async_from_lds(rocke_lower_t* L, const rocke_op rocke_ll_fail(L, ROCKE_ERR_VALUE, "global_store_async_from_lds expects two operands"); dst = op->operands[0]; lds = op->operands[1]; - if(!dst->type || dst->type->kind != ROCKE_TYPE_PTR || !dst->type->space - || strcmp(dst->type->space, "global") != 0) + if(strcmp(rocke_ll_value_ptr_type(L, dst), "ptr addrspace(1)") != 0) rocke_ll_fail( L, ROCKE_ERR_VALUE, "global_store_async_from_lds dst_ptr must be a global pointer"); width = ll_attr_int(op, "width_bytes", 0); @@ -1817,8 +1839,7 @@ static void op_tile_global_load_tr16_b128(rocke_lower_t* L, const rocke_op_t* op if(op->num_operands != 1 || op->num_results != 1) rocke_ll_fail(L, ROCKE_ERR_VALUE, "global_load_tr16_b128 expects one operand and result"); src = op->operands[0]; - if(!src->type || src->type->kind != ROCKE_TYPE_PTR || !src->type->space - || strcmp(src->type->space, "global") != 0) + if(strcmp(rocke_ll_value_ptr_type(L, src), "ptr addrspace(1)") != 0) rocke_ll_fail(L, ROCKE_ERR_VALUE, "global_load_tr16_b128 src_ptr must be a global pointer"); dtype = ll_attr_str(op, "dtype", ""); if(strcmp(dtype, "f16") == 0) @@ -2052,6 +2073,7 @@ static void op_tile_global_load_lds(rocke_lower_t* L, const rocke_op_t* op) int64_t aux = ll_attr_int(op, "aux", 0); const char* gep; const char* ptr_name; + ll_require_global_header_ptr(L, "global_load_lds", src_ptr); rocke_ll_need(L, "global.load.lds"); gep = rocke_ll_fresh(L, "gld_src"); rocke_ll_emitf(L, diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_llvm_internal.h b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_llvm_internal.h index ff0c5e1761fd..9e621a2bf7ec 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_llvm_internal.h +++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_llvm_internal.h @@ -104,10 +104,11 @@ const char* rocke_ll_matrix_decl_key_at(int index); /* ---------------------------------------------------------- anyptr overloads */ -/* One accepted address space of an llvm_anyptr_ty intrinsic, with the LLVM - * pointer text that names it. Mirrors one entry of the Python - * _S_PREFETCH_INST_PTR_TYPES / _AV_B128_PTR_TYPES dicts; every entry needs a - * matching ".p" row in ROCKE_LL_INTRINSIC_DECLS. */ +/* One accepted address space of a pointer operand, with the LLVM pointer text + * that names it. Mirrors one entry of the Python _S_PREFETCH_INST_PTR_TYPES / + * _AV_B128_PTR_TYPES / _GLOBAL_LOAD_PTR_TYPES dicts; an entry of an + * llvm_anyptr_ty intrinsic's table needs a matching ".p" row in + * ROCKE_LL_INTRINSIC_DECLS. */ typedef struct rocke_ll_anyptr_space { int space; @@ -120,6 +121,8 @@ extern const rocke_ll_anyptr_space_t ROCKE_LL_S_PREFETCH_DATA_PTR_TYPES[]; extern const int ROCKE_LL_S_PREFETCH_DATA_PTR_TYPES_COUNT; extern const rocke_ll_anyptr_space_t ROCKE_LL_AV_B128_PTR_TYPES[]; extern const int ROCKE_LL_AV_B128_PTR_TYPES_COUNT; +extern const rocke_ll_anyptr_space_t ROCKE_LL_GLOBAL_LOAD_PTR_TYPES[]; +extern const int ROCKE_LL_GLOBAL_LOAD_PTR_TYPES_COUNT; /* The resolver that consumes these tables is rocke_ll_anyptr_space, declared * with the other rocke_lower_t helpers below. */ @@ -585,11 +588,11 @@ const char* rocke_ll_param_llvm_type(rocke_lower_t* L, const rocke_param_t* p); * _Lowerer._ptr_llvm_type). */ const char* rocke_ll_value_ptr_type(rocke_lower_t* L, const rocke_value_t* v); -/* Address space of an llvm_anyptr_ty operand, validated against `allowed` - * (Python _Lowerer._anyptr_space). Also writes the matching pointer text to - * *out_ptr_ty when non-NULL. Fails (does not return) for a space the intrinsic - * does not accept: the space is part of the overload, so the mangled name, the - * declare, and the call site all have to agree with the pointer's real type. */ +/* Address space of a pointer operand, validated against `allowed` (Python + * _Lowerer._anyptr_space). Also writes the matching pointer text to + * *out_ptr_ty when non-NULL. Fails (does not return) for a space the op does + * not accept: the space is part of an llvm_anyptr_ty intrinsic's overload and + * of a plain load, so the call site has to name the pointer's real type. */ int rocke_ll_anyptr_space(rocke_lower_t* L, const char* op, const rocke_value_t* ptr, diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/op_vocabulary.md b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/op_vocabulary.md index 2c1b6b1d30bf..304b7d5bf3ff 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/op_vocabulary.md +++ b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/op_vocabulary.md @@ -106,6 +106,12 @@ Complete reference of operations recognized by `core/ir.py` and lowered to AMDGP `masked_global_load(ptr, idx, mask, other, dtype)` clamps the false-lane index to 0 before loading, then selects `other` afterward. +A pointer parameter declared `addr_space="constant"` is `ptr addrspace(4)` in the kernel header. The +three global loads above read through it as `load addrspace(4)`; they refuse a flat or LDS pointer. +Every other op that addresses global memory through it (the stores, atomics and flag ops above, +`cooperative_global_store`, `global_load_tr16_b128`, and the copies between global memory and LDS) is +refused by the LLVM lowering of both engines. + ### Global pointers | Op name | Builder | LLVM emission | HIP emission | diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py index e0d26d26cf0d..1cbfd67dd772 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py +++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py @@ -964,6 +964,13 @@ def _resolve_llvm_flavor() -> str: 0: "ptr", 1: "ptr addrspace(1)", } +# A ``memref.global_load*`` is a plain ``load``, which reads through a global +# pointer or through a parameter declared ``addr_space="constant"``. It names +# the pointer's own type, so no table entry needs a declare. +_GLOBAL_LOAD_PTR_TYPES: dict[int, str] = { + 1: "ptr addrspace(1)", + 4: "ptr addrspace(4)", +} # Declaration overrides, one table per declaration ABI generation @@ -1867,20 +1874,22 @@ def _lower_composed_op( def _anyptr_space( self, op: str, ptr: Value, allowed: Dict[int, str] ) -> Tuple[int, str]: - """Resolve an ``llvm_anyptr_ty`` operand to its (address space, type). - - For these intrinsics the address space is part of the overload, so the - mangled name, the declare, and the call site all have to agree with the - pointer's real type. Naming a bare ``ptr`` for an ``addrspace(N)`` value - is not a lax spelling -- LLVM rejects the module outright with - "defined with type 'ptr addrspace(N)' but expected 'ptr'". + """Resolve a pointer operand that may be in any of the ``allowed`` + address spaces to its (address space, type). + + For an ``llvm_anyptr_ty`` intrinsic the address space is part of the + overload, and for a plain ``load`` it is part of the instruction, so + the call site has to name the pointer's real type. Naming a bare + ``ptr`` for an ``addrspace(N)`` value is not a lax spelling -- LLVM + rejects the module outright with "defined with type + 'ptr addrspace(N)' but expected 'ptr'". """ ty = self._ptr_llvm_type(ptr) for space, text in allowed.items(): if text == ty: return space, text raise ValueError( - f"{op}: pointer operand is {ty}, but the intrinsic accepts only " + f"{op}: pointer operand is {ty}, but the op accepts only " f"{', '.join(allowed[s] for s in sorted(allowed))}" ) @@ -3429,31 +3438,34 @@ def _smem_global_name(self, smem_value: Value) -> Tuple[str, SmemType]: def _op_memref_global_load(self, op: Op) -> None: ptr, idx = op.operands + _, ptr_ty = self._anyptr_space("global_load", ptr, _GLOBAL_LOAD_PTR_TYPES) gep = self._fresh("gep") align = int(op.attrs.get("align", 2)) self._current().emit( - f" {gep} = getelementptr inbounds half, ptr addrspace(1) " + f" {gep} = getelementptr inbounds half, {ptr_ty} " f"{self._operand(ptr)}, i32 {self._operand(idx)}" ) self._current().emit( - f" {op.result.name} = load half, ptr addrspace(1) {gep}, align {align}" + f" {op.result.name} = load half, {ptr_ty} {gep}, align {align}" ) def _op_memref_global_load_typed(self, op: Op) -> None: ptr, idx = op.operands + _, ptr_ty = self._anyptr_space("global_load_typed", ptr, _GLOBAL_LOAD_PTR_TYPES) elem_ty = _llvm_type(op.result.type) gep = self._fresh("gep") align = int(op.attrs.get("align", 1)) self._current().emit( - f" {gep} = getelementptr inbounds {elem_ty}, ptr addrspace(1) " + f" {gep} = getelementptr inbounds {elem_ty}, {ptr_ty} " f"{self._operand(ptr)}, i32 {self._operand(idx)}" ) self._current().emit( - f" {op.result.name} = load {elem_ty}, ptr addrspace(1) {gep}, align {align}" + f" {op.result.name} = load {elem_ty}, {ptr_ty} {gep}, align {align}" ) def _op_memref_global_store(self, op: Op) -> None: ptr, idx, val = op.operands + self._require_global_header_ptr("global_store", ptr) gep = self._fresh("gep") align = int(op.attrs.get("align", 2)) self._current().emit( @@ -3466,6 +3478,7 @@ def _op_memref_global_store(self, op: Op) -> None: def _op_memref_global_store_typed(self, op: Op) -> None: ptr, idx, val = op.operands + self._require_global_header_ptr("global_store_typed", ptr) elem_ty = _llvm_type(val.type) gep = self._fresh("gep") align = int(op.attrs.get("align", 1)) @@ -3485,6 +3498,7 @@ def _op_memref_global_atomic_add(self, op: Op) -> None: ``atomicrmw`` semantics, which match HIP's ``atomicAdd``). """ ptr, idx, val = op.operands + self._require_global_header_ptr("global_atomic_add", ptr) elem_ty = _llvm_type(val.type) gep = self._fresh("gep") self._current().emit( @@ -3529,6 +3543,7 @@ def _op_memref_fence(self, op: Op) -> None: def _op_memref_global_flag_store(self, op: Op) -> None: """Lower ``memref.global_flag_store`` to a scoped ``store atomic i32``.""" ptr, idx, val = op.operands + self._require_global_header_ptr("global_flag_store", ptr) scope = _sync_attr(op, "scope", _FENCE_SCOPES) ordering = _sync_attr(op, "ordering", _FLAG_STORE_ORDERINGS) gep = self._fresh("gep") @@ -3549,6 +3564,7 @@ def _op_memref_global_flag_wait_eq(self, op: Op) -> None: orders every later global load after the observed flag value. """ ptr, idx, expect = op.operands + self._require_global_header_ptr("global_flag_wait_eq", ptr) scope = _sync_attr(op, "scope", _FENCE_SCOPES) gep = self._fresh("gep") cur = self._current() @@ -3584,6 +3600,7 @@ def _op_memref_global_atomic_add_pk_bf16(self, op: Op) -> None: C-scatter address fix, so match its width). """ ptr, idx, val = op.operands + self._require_global_header_ptr("global_atomic_add_pk_bf16", ptr) idx_ty = _llvm_type(idx.type) gep = self._fresh("gep") self._current().emit( @@ -3610,6 +3627,7 @@ def _op_memref_global_atomic_add_pk_f16(self, op: Op) -> None: ``llvm.amdgcn.global.atomic.fadd.v2f16.p1`` (gfx940+). """ ptr, idx, val = op.operands + self._require_global_header_ptr("global_atomic_add_pk_f16", ptr) self._need("global.atomic.fadd.v2f16") gep = self._fresh("gep") self._current().emit( @@ -3626,6 +3644,7 @@ def _op_memref_global_load_vN(self, op: Op) -> None: """Vectorised load: a single naturally-aligned global_load_dwordx{1,2,4} on AMDGPU when the address is aligned.""" ptr, idx = op.operands + _, ptr_ty = self._anyptr_space("global_load_vN", ptr, _GLOBAL_LOAD_PTR_TYPES) vec = int(op.attrs["vec"]) elem_ty = _llvm_type(op.result.type.elem) # type: ignore[attr-defined] # GEP index width must match the operand type -- i64 indices are @@ -3636,12 +3655,12 @@ def _op_memref_global_load_vN(self, op: Op) -> None: # elem, then bitcast to ptr to . This is exactly the # pattern Clang emits for `*(__fp16x4_t*)(ptr + idx)`. self._current().emit( - f" {gep} = getelementptr inbounds {elem_ty}, ptr addrspace(1) " + f" {gep} = getelementptr inbounds {elem_ty}, {ptr_ty} " f"{self._operand(ptr)}, {idx_ty} {self._operand(idx)}" ) align = int(op.attrs.get("align", vec * 2)) self._current().emit( - f" {op.result.name} = load <{vec} x {elem_ty}>, ptr addrspace(1) {gep}, " + f" {op.result.name} = load <{vec} x {elem_ty}>, {ptr_ty} {gep}, " f"align {align}{self._nontemporal_md(op)}" ) @@ -4512,6 +4531,7 @@ def _op_memref_cooperative_global_store(self, op: Op) -> None: addresses form a stride-1 pattern across active lanes. """ ptr, addrs, values = op.operands + self._require_global_header_ptr("cooperative_global_store", ptr) n = int(op.attrs["vec"]) if not isinstance(values.type, VectorType): raise NotImplementedError("cooperative_global_store requires vector values") @@ -5555,6 +5575,7 @@ def _op_tile_global_load_async_to_lds(self, op: Op) -> None: src_index = op.operands[1] lds_smem = op.operands[2] lds_indices = op.operands[3:] + self._require_global_header_ptr("global_load_async_to_lds", src_ptr) width = int(op.attrs["width_bytes"]) cpol = int(op.attrs.get("cpol", 0)) ioff = int(op.attrs.get("offset_bytes", 0)) @@ -6432,6 +6453,7 @@ def _op_tile_buffer_load_lds_async(self, op: Op) -> None: def _op_tile_global_load_lds(self, op: Op) -> None: src_ptr, byte_off, lds_addr = op.operands + self._require_global_header_ptr("global_load_lds", src_ptr) size_bytes = int(op.attrs["size_bytes"]) aux = int(op.attrs.get("aux", 0)) self._need("global.load.lds") @@ -6539,6 +6561,7 @@ def _op_vector_trunc_f32_to(self, op: Op) -> None: def _op_memref_global_store_vN(self, op: Op) -> None: ptr, idx, val = op.operands + self._require_global_header_ptr("global_store_vN", ptr) vec = int(op.attrs["vec"]) gep = self._fresh("gep") elem_ty = ( @@ -6580,6 +6603,7 @@ def _op_memref_global_store_vN(self, op: Op) -> None: def _op_memref_global_atomic_add_f32(self, op: Op) -> None: ptr, idx, val = op.operands + self._require_global_header_ptr("global_atomic_add_f32", ptr) gep = self._fresh("gep") self._current().emit( f" {gep} = getelementptr inbounds float, ptr addrspace(1) " diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_constant_space_global_ops.py b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_constant_space_global_ops.py new file mode 100644 index 000000000000..8b899fb27435 --- /dev/null +++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_constant_space_global_ops.py @@ -0,0 +1,278 @@ +# Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +# SPDX-License-Identifier: MIT + +"""Global-memory ops on a pointer parameter declared ``addr_space="constant"``. + +Such a parameter keeps its IR type (a global pointer), but the kernel header +types it ``ptr addrspace(4)``. An instruction that names it +``ptr addrspace(1)`` does not match the header, and LLVM rejects the module. + +- The global loads (``global_load_f16``, ``global_load``, ``global_load_vN``) + name the pointer's own type, so they read through a constant-space parameter + and the module compiles. They refuse a pointer in any other space (flat or + LDS), which they used to name ``ptr addrspace(1)`` as well. +- Every other op that addresses global memory through the pointer -- stores, + atomics, flag ops, the cooperative store, the transposed load and the copies + between global memory and LDS -- is refused by the LLVM lowering of both + engines. Most use the text the ``global_ptr_*`` ops use; the two ops that + already refused a non-global pointer keep their own text. +""" + +from __future__ import annotations + +from collections.abc import Callable + +import pytest +from rocke.core.ir import ( + BF16, + F16, + F32, + I16, + I32, + IRBuilder, + KernelDef, + PtrType, + Type, + Value, +) +from rocke.tests.engines import assert_both_engines_reject, lower_with_both_engines +from rocke.tests.ir_validity import assert_target_accepts + +FLAVORS = ("llvm20", "llvm22", "llvm23", "llvm24") +ARCH = "gfx942" + + +def _constant_param(b: IRBuilder, dtype: Type) -> Value: + return b.param("X", PtrType(dtype, "global"), addr_space="constant") + + +# -------------------------------------------------------------------------- +# Loads read through a constant-space parameter +# -------------------------------------------------------------------------- + + +def _load_kernel( + name: str, + load: Callable[[IRBuilder, Value], Value], + dtype: Type, + x_space: str = "constant", +) -> KernelDef: + """Load from ``X`` with ``load`` and store the value to the global ``Y``. + + ``X`` is a global pointer parameter declared ``addr_space="constant"``, or, + for any other ``x_space``, a pointer parameter in that space.""" + b = IRBuilder(name) + if x_space == "constant": + x = _constant_param(b, dtype) + else: + x = b.param("X", PtrType(dtype, x_space)) + y = b.param("Y", PtrType(dtype, "global")) + value = load(b, x) + if value.type == dtype: + b.global_store(y, b.const_i32(0), value) + else: + b.global_store_vN(y, b.const_i32(0), value, 4) + b.ret() + return b.kernel + + +LOADS = { + "global_load": ( + lambda b, x: b.global_load_f16(x, b.const_i32(1)), + F16, + "half", + "half", + ), + "global_load_typed": ( + lambda b, x: b.global_load(x, b.const_i32(1), F32, align=4), + F32, + "float", + "float", + ), + "global_load_vN": ( + lambda b, x: b.global_load_vN(x, b.const_i32(4), F16, 4), + F16, + "half", + "<4 x half>", + ), +} + + +@pytest.mark.parametrize("flavor", FLAVORS) +@pytest.mark.parametrize("op_name", sorted(LOADS)) +def test_both_llvm_lowerings_load_through_a_constant_space_parameter(op_name, flavor): + load, dtype, elem_ty, loaded_ty = LOADS[op_name] + ir = lower_with_both_engines( + _load_kernel(f"{op_name}_constant", load, dtype), ARCH, llvm_flavor=flavor + ) + assert f"@{op_name}_constant(ptr addrspace(4) nocapture %X, " in ir + assert f"getelementptr inbounds {elem_ty}, ptr addrspace(4) %X, i32 " in ir + assert f"= load {loaded_ty}, ptr addrspace(4) %gep" in ir + assert "ptr addrspace(1) %X" not in ir + + +@pytest.mark.parametrize("op_name", sorted(LOADS)) +def test_a_load_through_a_constant_space_parameter_compiles(op_name): + load, dtype, _, _ = LOADS[op_name] + ir = lower_with_both_engines(_load_kernel(f"{op_name}_compiles", load, dtype), ARCH) + assert_target_accepts(ir, ARCH) + + +@pytest.mark.parametrize( + "x_space, header_type", [("private", "ptr"), ("lds", "ptr addrspace(3)")] +) +@pytest.mark.parametrize("op_name", sorted(LOADS)) +def test_both_llvm_lowerings_refuse_a_load_through_another_space( + op_name, x_space, header_type +): + load, dtype, _, _ = LOADS[op_name] + kernel = _load_kernel(f"{op_name}_{x_space}", load, dtype, x_space) + assert_both_engines_reject( + kernel, + ARCH, + ValueError, + f"{op_name}: pointer operand is {header_type}, but the op accepts only " + "ptr addrspace(1), ptr addrspace(4)", + ) + + +# -------------------------------------------------------------------------- +# Every other global-memory op refuses a constant-space parameter +# -------------------------------------------------------------------------- + + +def _header_refusal(op_name: str) -> str: + return ( + f"{op_name} ptr must be a global pointer, got ptr addrspace(4) " + "(its parameter's addr_space)" + ) + + +def _lds_address(b: IRBuilder, dtype: Type) -> Value: + return b.smem_addr_of(b.smem_alloc(dtype, [64], name_hint="stage")) + + +# op name -> (element type of X, the op applied to X, arch, flavor, message) +REFUSALS = { + "global_store": ( + F16, + lambda b, x: b.store_f16( + x, b.const_i32(0), b.global_load_f16(x, b.const_i32(1)) + ), + ARCH, + None, + _header_refusal("global_store"), + ), + "global_store_typed": ( + F32, + lambda b, x: b.global_store(x, b.const_i32(0), b.const_f32(1.0)), + ARCH, + None, + _header_refusal("global_store_typed"), + ), + "global_store_vN": ( + F16, + lambda b, x: b.global_store_vN(x, b.const_i32(0), b.zero_vec(F16, 4), 4), + ARCH, + None, + _header_refusal("global_store_vN"), + ), + "global_atomic_add": ( + I32, + lambda b, x: b.global_atomic_add(x, b.const_i32(0), b.const_i32(1)), + ARCH, + None, + _header_refusal("global_atomic_add"), + ), + "global_atomic_add_f32": ( + F32, + lambda b, x: b.global_atomic_add_f32(x, b.const_i32(0), b.const_f32(1.0)), + ARCH, + None, + _header_refusal("global_atomic_add_f32"), + ), + "global_atomic_add_pk_bf16": ( + BF16, + lambda b, x: b.global_atomic_add_pk_bf16( + x, b.const_i32(0), b.zero_vec(BF16, 2) + ), + ARCH, + None, + _header_refusal("global_atomic_add_pk_bf16"), + ), + "global_atomic_add_pk_f16": ( + F16, + lambda b, x: b.global_atomic_add_pk_f16(x, b.const_i32(0), b.zero_vec(F16, 2)), + ARCH, + None, + _header_refusal("global_atomic_add_pk_f16"), + ), + "global_flag_store": ( + I32, + lambda b, x: b.global_flag_store(x, b.const_i32(0), b.const_i32(1)), + ARCH, + None, + _header_refusal("global_flag_store"), + ), + "global_flag_wait_eq": ( + I32, + lambda b, x: b.global_flag_wait_eq(x, b.const_i32(0), b.const_i32(1)), + ARCH, + None, + _header_refusal("global_flag_wait_eq"), + ), + "cooperative_global_store": ( + F32, + lambda b, x: b.cooperative_global_store( + x, b.zero_vec(I32, 4), b.zero_vec(F32, 4) + ), + ARCH, + None, + _header_refusal("cooperative_global_store"), + ), + "global_load_lds": ( + F32, + lambda b, x: b.global_load_lds(x, b.const_i32(0), _lds_address(b, F32), 4), + ARCH, + None, + _header_refusal("global_load_lds"), + ), + "global_load_async_to_lds": ( + I32, + lambda b, x: b.global_load_async_to_lds( + x, + b.const_i32(0), + b.smem_alloc(I32, [64], name_hint="stage"), + [b.const_i32(0)], + width_bytes=4, + ), + "gfx1250", + "llvm23", + _header_refusal("global_load_async_to_lds"), + ), + "global_store_async_from_lds": ( + I16, + lambda b, x: b.global_store_async_from_lds( + x, _lds_address(b, I16), width_bytes=16 + ), + "gfx1250", + "llvm23", + "global_store_async_from_lds dst_ptr must be a global pointer", + ), + "global_load_tr16_b128": ( + I16, + lambda b, x: b.global_load_tr16_b128(x, dtype=I16), + "gfx1250", + "llvm23", + "global_load_tr16_b128 src_ptr must be a global pointer", + ), +} + + +@pytest.mark.parametrize("op_name", sorted(REFUSALS)) +def test_both_llvm_lowerings_refuse_a_constant_space_parameter(op_name): + dtype, apply_op, arch, flavor, message = REFUSALS[op_name] + b = IRBuilder(f"{op_name}_constant") + apply_op(b, _constant_param(b, dtype)) + b.ret() + assert_both_engines_reject(b.kernel, arch, TypeError, message, llvm_flavor=flavor) diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/test_rocke.py b/dnn-providers/hip-kernel-provider/rocke/platform/tests/test_rocke.py index f0ad818e8483..5d38b0aa8e09 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/test_rocke.py +++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/test_rocke.py @@ -4693,7 +4693,7 @@ def test_av_b128_rejects_lds_pointer(self): with self.assertRaisesRegex( error_type, r"av_load_b128: pointer operand is ptr addrspace\(3\), " - r"but the intrinsic accepts only ptr, ptr addrspace\(1\)", + r"but the op accepts only ptr, ptr addrspace\(1\)", ): lower_kernel_to_llvm(b.kernel, arch="gfx950") @@ -4798,7 +4798,7 @@ def test_s_prefetch_inst_rejects_lds_pointer(self): with self.assertRaisesRegex( error_type, r"s_prefetch_inst: pointer operand is ptr addrspace\(3\), " - r"but the intrinsic accepts only ptr, ptr addrspace\(1\), ptr addrspace\(4\)", + r"but the op accepts only ptr, ptr addrspace\(1\), ptr addrspace\(4\)", ): lower_kernel_to_llvm(b.kernel, arch="gfx950") From eabc26ad9443c430708dcb0d9dfc501f2e322ffd Mon Sep 17 00:00:00 2001 From: Chris Millette Date: Tue, 3 Nov 2026 01:31:50 -0500 Subject: [PATCH 4/7] fix(rocke): report the global-pointer header refusal as a type error The Python engine refuses a global-pointer op on a parameter whose header type is not ptr addrspace(1) with a TypeError. The C++ engine failed the same refusal with ROCKE_ERR_VALUE, and so did its global_store_async_from_lds and global_load_tr16_b128 pointer checks. It now fails them with ROCKE_ERR_TYPE, the status that matches TypeError. Co-Authored-By: Claude --- .../rocke/platform/cpp/core/lower_llvm/mem.cpp | 14 +++++++------- .../tests/core/future_intrinsic_lowering.cpp | 2 +- 2 files changed, 8 insertions(+), 8 deletions(-) diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mem.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mem.cpp index 79bfccce0a17..c16fcf7eb7cc 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mem.cpp +++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mem.cpp @@ -108,17 +108,17 @@ static void op_tile_smem_alloc(rocke_lower_t* L, const rocke_op_t* op) (void)op; } -/* Python _require_global_header_ptr; the text is the Python engine's. The - * global-pointer ops emit against ptr addrspace(1); the IR rules check the IR - * type, but a parameter's addr_space override (such as "constant") moves its - * header type away from it. */ +/* Python _require_global_header_ptr; the text is the Python engine's, and the + * status is its TypeError's. The global-pointer ops emit against + * ptr addrspace(1); the IR rules check the IR type, but a parameter's + * addr_space override (such as "constant") moves its header type away from it. */ static void ll_require_global_header_ptr(rocke_lower_t* L, const char* op_name, const rocke_value_t* ptr) { const char* header = rocke_ll_value_ptr_type(L, ptr); if(strcmp(header, "ptr addrspace(1)") != 0) rocke_ll_fail(L, - ROCKE_ERR_VALUE, + ROCKE_ERR_TYPE, "%s ptr must be a global pointer, got %s (its parameter's addr_space)", op_name, header); @@ -1793,7 +1793,7 @@ static void op_tile_global_store_async_from_lds(rocke_lower_t* L, const rocke_op lds = op->operands[1]; if(strcmp(rocke_ll_value_ptr_type(L, dst), "ptr addrspace(1)") != 0) rocke_ll_fail( - L, ROCKE_ERR_VALUE, "global_store_async_from_lds dst_ptr must be a global pointer"); + L, ROCKE_ERR_TYPE, "global_store_async_from_lds dst_ptr must be a global pointer"); width = ll_attr_int(op, "width_bytes", 0); if(width == 1) suffix = "b8"; @@ -1840,7 +1840,7 @@ static void op_tile_global_load_tr16_b128(rocke_lower_t* L, const rocke_op_t* op rocke_ll_fail(L, ROCKE_ERR_VALUE, "global_load_tr16_b128 expects one operand and result"); src = op->operands[0]; if(strcmp(rocke_ll_value_ptr_type(L, src), "ptr addrspace(1)") != 0) - rocke_ll_fail(L, ROCKE_ERR_VALUE, "global_load_tr16_b128 src_ptr must be a global pointer"); + rocke_ll_fail(L, ROCKE_ERR_TYPE, "global_load_tr16_b128 src_ptr must be a global pointer"); dtype = ll_attr_str(op, "dtype", ""); if(strcmp(dtype, "f16") == 0) suffix = "v8f16"; diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/future_intrinsic_lowering.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/future_intrinsic_lowering.cpp index c8b1c6c23510..04dc02e6d136 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/future_intrinsic_lowering.cpp +++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/future_intrinsic_lowering.cpp @@ -3356,7 +3356,7 @@ void case_buffer_rsrc_refusals() "gfx942", flavor, &err); - if(st != ROCKE_ERR_VALUE + if(st != ROCKE_ERR_TYPE || err.find("buffer_rsrc ptr must be a global pointer, got ptr addrspace(4) " "(its parameter's addr_space)") == std::string::npos) From 6196bec36a07c8fc8cb9930d915833c6e9184f03 Mon Sep 17 00:00:00 2001 From: Chris Millette Date: Tue, 3 Nov 2026 16:43:39 -0500 Subject: [PATCH 5/7] fix(rocke): refuse a pointer in the wrong space through one check per engine Every LLVM-lowering op that names its pointer in one fixed space now refuses any other header type through one helper per engine (_require_header_ptr / rocke_ll_require_header_ptr, with a global-pointer form), so the store, atomic, flag, prefetch, transposed-load, cluster-load and LDS-copy refusals share one text and one error type: TypeError in Python, ROCKE_ERR_TYPE in C++. The multi-space resolver the global loads, av_*_b128 and s_prefetch_* use refuses with the same type. The refusal names the op's builder (store_f16 rather than global_store, global_store rather than global_store_typed) and the operand, and blames the parameter's addr_space only when the override moved its header type. The engine-parity test helpers gain per-engine forms, which also check the C++ status against the expected exception type. The constant-space and buffer_rsrc tests run per engine, so their Python cases run in lanes that do not build the rocke_engine binding, and a new future_intrinsic_lowering CTest case checks a sample of the C++ side there. The comments on the HIP builtin now say that clang after ROCm 7.1 takes a 64-bit num_records (int64_t in LLVM 22, uint64_t after). Co-Authored-By: Claude --- .../cpp/core/lower_hip/lower_hip_mem.cpp | 4 +- .../platform/cpp/core/lower_llvm/control.cpp | 11 +- .../platform/cpp/core/lower_llvm/core.cpp | 60 +++- .../platform/cpp/core/lower_llvm/mem.cpp | 68 ++--- .../cpp/include/rocke/lower_llvm_internal.h | 33 ++- .../dsl_docs/reference/op_vocabulary.md | 5 +- .../platform/python/rocke/core/lower_hip.py | 6 +- .../platform/python/rocke/core/lower_llvm.py | 101 ++++--- .../platform/python/rocke/tests/engines.py | 256 ++++++++++++++---- .../tests/core/future_intrinsic_lowering.cpp | 172 +++++++++++- .../tests/core/test_buffer_rsrc_operands.py | 77 ++++-- .../core/test_constant_space_global_ops.py | 170 +++++++++--- .../tests/core/test_gfx1250_prefetch.py | 2 +- .../rocke/platform/tests/test_rocke.py | 6 +- 14 files changed, 722 insertions(+), 249 deletions(-) diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp index 611a250f491f..afa2bcb60d55 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp +++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp @@ -1074,8 +1074,8 @@ static rocke_status_t _op_tile_buffer_rsrc(rocke_h_lowerer_t* lw, const rocke_op } ptr = op->operands[0]; num_bytes = op->operands[1]; - /* num_records is cast to (int), which ROCm 7's builtin (int num_records) - * and later clang's (uint64_t num_records) both accept; the cast would + /* num_records is cast to (int), which ROCm 7.1's builtin (int num_records) + * and later clang's (a 64-bit num_records) both accept; the cast would * truncate an i64 byte count. A constant-space ptr is kept: the HIP * signature declares every pointer parameter the same way, and the builtin * takes a void*. */ diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/control.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/control.cpp index 936a7a1e3d36..5ab8a87183c2 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/control.cpp +++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/control.cpp @@ -1374,16 +1374,12 @@ static void _op_tile_s_buffer_prefetch_data(rocke_lower_t* L, const rocke_op_t* static void _op_tile_global_prefetch(rocke_lower_t* L, const rocke_op_t* op) { - const char* ty; int64_t cp; if(!rocke_ll_live(L)) return; rocke_ll_require_llvm23_target_fact(L, "global_prefetch", ROCKE_LL_FACT_GLOBAL_PREFETCH); ll_expect_operands(L, op, "global_prefetch", 1); - ty = rocke_ll_value_ptr_type(L, op->operands[0]); - if(strcmp(ty, "ptr addrspace(1)") != 0) - rocke_ll_fail( - L, ROCKE_ERR_VALUE, "global_prefetch ptr must be a global pointer, got %s", ty); + rocke_ll_require_global_header_ptr(L, "global_prefetch", op->operands[0], "ptr"); cp = ll_check_cachepolicy(L, op, "global_prefetch"); rocke_ll_need(L, "global.prefetch"); rocke_ll_emitf(L, @@ -1394,15 +1390,12 @@ static void _op_tile_global_prefetch(rocke_lower_t* L, const rocke_op_t* op) static void _op_tile_flat_prefetch(rocke_lower_t* L, const rocke_op_t* op) { - const char* ty; int64_t cp; if(!rocke_ll_live(L)) return; rocke_ll_require_llvm23_target_fact(L, "flat_prefetch", ROCKE_LL_FACT_GLOBAL_PREFETCH); ll_expect_operands(L, op, "flat_prefetch", 1); - ty = rocke_ll_value_ptr_type(L, op->operands[0]); - if(strcmp(ty, "ptr") != 0) - rocke_ll_fail(L, ROCKE_ERR_VALUE, "flat_prefetch ptr must be a flat pointer, got %s", ty); + rocke_ll_require_header_ptr(L, "flat_prefetch", op->operands[0], "ptr", "flat", "ptr"); cp = ll_check_cachepolicy(L, op, "flat_prefetch"); rocke_ll_need(L, "flat.prefetch"); rocke_ll_emitf(L, diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp index 7fa183c7c800..f3408f2366d7 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp +++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp @@ -653,24 +653,62 @@ const char* rocke_ll_param_llvm_type(rocke_lower_t* L, const rocke_param_t* p) return rocke_ll_llvm_type(L, p->type); } +const rocke_param_t* rocke_ll_value_param(const rocke_lower_t* L, const rocke_value_t* v) +{ + if(!L || !L->kernel || !v || !v->name || v->name[0] != '%') + return NULL; + for(int i = 0; i < L->kernel->num_params; i++) + { + const rocke_param_t* p = L->kernel->params[i]; + if(p && p->name && strcmp(v->name + 1, p->name) == 0) + return p; + } + return NULL; +} + const char* rocke_ll_value_ptr_type(rocke_lower_t* L, const rocke_value_t* v) { - int i; if(!v) return ""; /* Python _Lowerer._ptr_llvm_type: kernel params carry the header's type. */ - if(L && L->kernel && v->name) - { - for(i = 0; i < L->kernel->num_params; i++) - { - const rocke_param_t* p = L->kernel->params[i]; - if(p && p->name && v->name[0] == '%' && strcmp(v->name + 1, p->name) == 0) - return rocke_ll_param_llvm_type(L, p); - } - } + const rocke_param_t* param = rocke_ll_value_param(L, v); + if(param) + return rocke_ll_param_llvm_type(L, param); return rocke_ll_llvm_type(L, v->type); } +void rocke_ll_require_header_ptr(rocke_lower_t* L, + const char* op_name, + const rocke_value_t* ptr, + const char* operand, + const char* space, + const char* expected) +{ + const char* header = rocke_ll_value_ptr_type(L, ptr); + if(strcmp(header, expected) == 0) + return; + const rocke_param_t* param = rocke_ll_value_param(L, ptr); + const bool moved + = param + && strcmp(rocke_ll_param_llvm_type(L, param), rocke_ll_llvm_type(L, param->type)) != 0; + rocke_ll_fail(L, + ROCKE_ERR_TYPE, + "%s %s must be a %s pointer, got %s%s", + op_name, + operand, + space, + header, + moved ? " (its parameter's addr_space)" : ""); +} + +void rocke_ll_require_global_header_ptr(rocke_lower_t* L, + const char* op_name, + const rocke_value_t* ptr, + const char* operand) +{ + rocke_ll_require_header_ptr(L, op_name, ptr, operand, "global", "ptr addrspace(1)"); +} + int rocke_ll_anyptr_space(rocke_lower_t* L, const char* op, const rocke_value_t* ptr, @@ -701,7 +739,7 @@ int rocke_ll_anyptr_space(rocke_lower_t* L, used += (size_t)n; } rocke_ll_fail( - L, ROCKE_ERR_VALUE, "%s: pointer operand is %s, but the op accepts only %s", op, ty, list); + L, ROCKE_ERR_TYPE, "%s: pointer operand is %s, but the op accepts only %s", op, ty, list); } const char* rocke_ll_smem_storage_type(rocke_lower_t* L, const rocke_type_t* smem) diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mem.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mem.cpp index c16fcf7eb7cc..d178acc0767c 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mem.cpp +++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mem.cpp @@ -108,22 +108,6 @@ static void op_tile_smem_alloc(rocke_lower_t* L, const rocke_op_t* op) (void)op; } -/* Python _require_global_header_ptr; the text is the Python engine's, and the - * status is its TypeError's. The global-pointer ops emit against - * ptr addrspace(1); the IR rules check the IR type, but a parameter's - * addr_space override (such as "constant") moves its header type away from it. */ -static void - ll_require_global_header_ptr(rocke_lower_t* L, const char* op_name, const rocke_value_t* ptr) -{ - const char* header = rocke_ll_value_ptr_type(L, ptr); - if(strcmp(header, "ptr addrspace(1)") != 0) - rocke_ll_fail(L, - ROCKE_ERR_TYPE, - "%s ptr must be a global pointer, got %s (its parameter's addr_space)", - op_name, - header); -} - /* Python _anyptr_space over _GLOBAL_LOAD_PTR_TYPES: the pointer type a global * load names, a global pointer or a parameter declared addr_space="constant". */ static const char* @@ -143,7 +127,7 @@ static void op_memref_global_load(rocke_lower_t* L, const rocke_op_t* op) { const rocke_value_t* ptr = op->operands[0]; const rocke_value_t* idx = op->operands[1]; - const char* ptr_ty = ll_global_load_ptr_type(L, "global_load", ptr); + const char* ptr_ty = ll_global_load_ptr_type(L, "global_load_f16", ptr); const char* gep = rocke_ll_fresh(L, "gep"); int64_t align = ll_attr_int(op, "align", 2); rocke_ll_emitf(L, @@ -160,7 +144,7 @@ static void op_memref_global_load_typed(rocke_lower_t* L, const rocke_op_t* op) { const rocke_value_t* ptr = op->operands[0]; const rocke_value_t* idx = op->operands[1]; - const char* ptr_ty = ll_global_load_ptr_type(L, "global_load_typed", ptr); + const char* ptr_ty = ll_global_load_ptr_type(L, "global_load", ptr); const char* elem_ty = rocke_ll_llvm_type(L, op->results[0]->type); const char* gep = rocke_ll_fresh(L, "gep"); int64_t align = ll_attr_int(op, "align", 1); @@ -180,7 +164,7 @@ static void op_memref_global_store(rocke_lower_t* L, const rocke_op_t* op) const rocke_value_t* ptr = op->operands[0]; const rocke_value_t* idx = op->operands[1]; const rocke_value_t* val = op->operands[2]; - ll_require_global_header_ptr(L, "global_store", ptr); + rocke_ll_require_global_header_ptr(L, "store_f16", ptr, "ptr"); const char* gep = rocke_ll_fresh(L, "gep"); int64_t align = ll_attr_int(op, "align", 2); rocke_ll_emitf(L, @@ -200,7 +184,7 @@ static void op_memref_global_store_typed(rocke_lower_t* L, const rocke_op_t* op) const rocke_value_t* ptr = op->operands[0]; const rocke_value_t* idx = op->operands[1]; const rocke_value_t* val = op->operands[2]; - ll_require_global_header_ptr(L, "global_store_typed", ptr); + rocke_ll_require_global_header_ptr(L, "global_store", ptr, "ptr"); const char* elem_ty = rocke_ll_llvm_type(L, val->type); const char* gep = rocke_ll_fresh(L, "gep"); int64_t align = ll_attr_int(op, "align", 1); @@ -266,7 +250,7 @@ static void op_memref_global_flag_store(rocke_lower_t* L, const rocke_op_t* op) const rocke_value_t* ptr = op->operands[0]; const rocke_value_t* idx = op->operands[1]; const rocke_value_t* val = op->operands[2]; - ll_require_global_header_ptr(L, "global_flag_store", ptr); + rocke_ll_require_global_header_ptr(L, "global_flag_store", ptr, "ptr"); const char* scope = ll_sync_attr(L, op, "scope", k_ll_sync_scopes); const char* ordering = ll_sync_attr(L, op, "ordering", k_ll_flag_store_orderings); const char* gep = rocke_ll_fresh(L, "gep"); @@ -292,7 +276,7 @@ static void op_memref_global_flag_wait_eq(rocke_lower_t* L, const rocke_op_t* op const rocke_value_t* ptr = op->operands[0]; const rocke_value_t* idx = op->operands[1]; const rocke_value_t* expect = op->operands[2]; - ll_require_global_header_ptr(L, "global_flag_wait_eq", ptr); + rocke_ll_require_global_header_ptr(L, "global_flag_wait_eq", ptr, "ptr"); const char* scope = ll_sync_attr(L, op, "scope", k_ll_sync_scopes); const char* gep = rocke_ll_fresh(L, "gep"); rocke_ll_block_t* cur = rocke_ll_current(L); @@ -340,7 +324,7 @@ static void op_memref_global_atomic_add(rocke_lower_t* L, const rocke_op_t* op) bool is_f32; const char* rmw_op; const char* md = ""; - ll_require_global_header_ptr(L, "global_atomic_add", ptr); + rocke_ll_require_global_header_ptr(L, "global_atomic_add", ptr, "ptr"); rocke_ll_emitf(L, " %s = getelementptr inbounds %s, ptr addrspace(1) %s, i32 %s", gep, @@ -373,7 +357,7 @@ static void op_memref_global_atomic_add_f32(rocke_lower_t* L, const rocke_op_t* const rocke_value_t* ptr = op->operands[0]; const rocke_value_t* idx = op->operands[1]; const rocke_value_t* val = op->operands[2]; - ll_require_global_header_ptr(L, "global_atomic_add_f32", ptr); + rocke_ll_require_global_header_ptr(L, "global_atomic_add_f32", ptr, "ptr"); const char* gep = rocke_ll_fresh(L, "gep"); rocke_ll_emitf(L, " %s = getelementptr inbounds float, ptr addrspace(1) %s, i32 %s", @@ -410,7 +394,7 @@ static void op_memref_global_atomic_add_pk_bf16(rocke_lower_t* L, const rocke_op const char* idx_ty = rocke_ll_llvm_type(L, idx->type); const char* gep = rocke_ll_fresh(L, "gep"); const char* ordering = ll_attr_str(op, "ordering", "monotonic"); - ll_require_global_header_ptr(L, "global_atomic_add_pk_bf16", ptr); + rocke_ll_require_global_header_ptr(L, "global_atomic_add_pk_bf16", ptr, "ptr"); rocke_ll_emitf(L, " %s = getelementptr inbounds bfloat, ptr addrspace(1) %s, %s %s", gep, @@ -435,7 +419,7 @@ static void op_memref_global_atomic_add_pk_f16(rocke_lower_t* L, const rocke_op_ const rocke_value_t* idx = op->operands[1]; const rocke_value_t* val = op->operands[2]; const char* gep; - ll_require_global_header_ptr(L, "global_atomic_add_pk_f16", ptr); + rocke_ll_require_global_header_ptr(L, "global_atomic_add_pk_f16", ptr, "ptr"); rocke_ll_need(L, "global.atomic.fadd.v2f16"); gep = rocke_ll_fresh(L, "gep"); rocke_ll_emitf(L, @@ -521,7 +505,7 @@ static void op_memref_global_store_vN(rocke_lower_t* L, const rocke_op_t* op) const rocke_value_t* ptr = op->operands[0]; const rocke_value_t* idx = op->operands[1]; const rocke_value_t* val = op->operands[2]; - ll_require_global_header_ptr(L, "global_store_vN", ptr); + rocke_ll_require_global_header_ptr(L, "global_store_vN", ptr, "ptr"); int64_t vec = ll_attr_int(op, "vec", 0); const char* gep = rocke_ll_fresh(L, "gep"); const char* elem_ty = ll_is_vec(val->type) ? rocke_ll_llvm_type(L, val->type->elem) @@ -559,7 +543,7 @@ static void op_memref_cooperative_global_store(rocke_lower_t* L, const rocke_op_ const char* addrs_ty; const char* values_ty; int64_t i; - ll_require_global_header_ptr(L, "cooperative_global_store", ptr); + rocke_ll_require_global_header_ptr(L, "cooperative_global_store", ptr, "ptr"); if(!ll_is_vec(values->type)) { rocke_ll_fail(L, ROCKE_ERR_NOTIMPL, "cooperative_global_store requires vector values"); @@ -967,7 +951,7 @@ static void op_tile_global_ptr_add(rocke_lower_t* L, const rocke_op_t* op) { const rocke_value_t* ptr = op->operands[0]; const rocke_value_t* off = op->operands[1]; - ll_require_global_header_ptr(L, "global_ptr_add", ptr); + rocke_ll_require_global_header_ptr(L, "global_ptr_add", ptr, "ptr"); const char* off_ty = rocke_ll_llvm_type(L, off->type); const char* off64; if(strcmp(off_ty, "i32") == 0) @@ -1003,7 +987,7 @@ static void op_tile_global_ptr_offset(rocke_lower_t* L, const rocke_op_t* op) rocke_ll_fail(L, ROCKE_ERR_VALUE, "%s", error); return; } - ll_require_global_header_ptr(L, "global_ptr_offset", op->operands[0]); + rocke_ll_require_global_header_ptr(L, "global_ptr_offset", op->operands[0], "ptr"); rocke_ll_emitf(L, " %s = getelementptr i8, ptr addrspace(1) %s, i64 %s", ll_res(op), @@ -1023,7 +1007,7 @@ static void op_tile_global_ptr_to_i64(rocke_lower_t* L, const rocke_op_t* op) return; } ptr = op->operands[0]; - ll_require_global_header_ptr(L, "global_ptr_to_i64", ptr); + rocke_ll_require_global_header_ptr(L, "global_ptr_to_i64", ptr, "ptr"); rocke_ll_emitf( L, " %s = ptrtoint ptr addrspace(1) %s to i64", ll_res(op), rocke_ll_operand(L, ptr)); } @@ -1041,7 +1025,7 @@ static void op_tile_buffer_rsrc(rocke_lower_t* L, const rocke_op_t* op) int word3; /* Python _op_tile_buffer_rsrc. The p1 overload names ptr addrspace(1), so * a constant-space parameter (ptr addrspace(4) in the header) is refused. */ - ll_require_global_header_ptr(L, "buffer_rsrc", ptr); + rocke_ll_require_global_header_ptr(L, "buffer_rsrc", ptr, "ptr"); rocke_ll_need(L, "make.buffer.rsrc.p1"); if(rocke_ll_flavor_is_modern(L->flavor)) { @@ -1706,7 +1690,7 @@ static void op_tile_global_load_async_to_lds(rocke_lower_t* L, const rocke_op_t* const char* idx_ty; rocke_strbuf_t gep; int i; - ll_require_global_header_ptr(L, "global_load_async_to_lds", src_ptr); + rocke_ll_require_global_header_ptr(L, "global_load_async_to_lds", src_ptr, "ptr"); if(width == 1) suffix = "b8"; @@ -1791,9 +1775,7 @@ static void op_tile_global_store_async_from_lds(rocke_lower_t* L, const rocke_op rocke_ll_fail(L, ROCKE_ERR_VALUE, "global_store_async_from_lds expects two operands"); dst = op->operands[0]; lds = op->operands[1]; - if(strcmp(rocke_ll_value_ptr_type(L, dst), "ptr addrspace(1)") != 0) - rocke_ll_fail( - L, ROCKE_ERR_TYPE, "global_store_async_from_lds dst_ptr must be a global pointer"); + rocke_ll_require_global_header_ptr(L, "global_store_async_from_lds", dst, "dst_ptr"); width = ll_attr_int(op, "width_bytes", 0); if(width == 1) suffix = "b8"; @@ -1839,8 +1821,7 @@ static void op_tile_global_load_tr16_b128(rocke_lower_t* L, const rocke_op_t* op if(op->num_operands != 1 || op->num_results != 1) rocke_ll_fail(L, ROCKE_ERR_VALUE, "global_load_tr16_b128 expects one operand and result"); src = op->operands[0]; - if(strcmp(rocke_ll_value_ptr_type(L, src), "ptr addrspace(1)") != 0) - rocke_ll_fail(L, ROCKE_ERR_TYPE, "global_load_tr16_b128 src_ptr must be a global pointer"); + rocke_ll_require_global_header_ptr(L, "global_load_tr16_b128", src, "src_ptr"); dtype = ll_attr_str(op, "dtype", ""); if(strcmp(dtype, "f16") == 0) suffix = "v8f16"; @@ -1877,7 +1858,6 @@ static void op_tile_cluster_load(rocke_lower_t* L, const rocke_op_t* op) const rocke_value_t* ptr; const rocke_value_t* mask; const rocke_type_t* rty; - const char* ptr_ty; const char* bits; const char* suffix; const char* expected; @@ -1892,9 +1872,7 @@ static void op_tile_cluster_load(rocke_lower_t* L, const rocke_op_t* op) rocke_ll_fail(L, ROCKE_ERR_VALUE, "cluster_load expects two operands and one result"); ptr = op->operands[0]; mask = op->operands[1]; - ptr_ty = rocke_ll_value_ptr_type(L, ptr); - if(strcmp(ptr_ty, "ptr addrspace(1)") != 0) - rocke_ll_fail(L, ROCKE_ERR_VALUE, "cluster_load ptr must be a global pointer"); + rocke_ll_require_global_header_ptr(L, "cluster_load", ptr, "ptr"); if(!ll_is_i32_value(mask)) rocke_ll_fail(L, ROCKE_ERR_VALUE, "cluster_load mask must be i32"); width = ll_attr_int(op, "width_bytes", 0); @@ -1974,9 +1952,7 @@ static void op_tile_cluster_load_async_to_lds(rocke_lower_t* L, const rocke_op_t src = op->operands[0]; lds = op->operands[1]; mask = op->operands[2]; - if(strcmp(rocke_ll_value_ptr_type(L, src), "ptr addrspace(1)") != 0) - rocke_ll_fail( - L, ROCKE_ERR_VALUE, "cluster_load_async_to_lds src_ptr must be a global pointer"); + rocke_ll_require_global_header_ptr(L, "cluster_load_async_to_lds", src, "src_ptr"); if(!ll_is_i32_value(mask)) rocke_ll_fail(L, ROCKE_ERR_VALUE, "cluster_load_async_to_lds mask must be i32"); width = ll_attr_int(op, "width_bytes", 0); @@ -2073,7 +2049,7 @@ static void op_tile_global_load_lds(rocke_lower_t* L, const rocke_op_t* op) int64_t aux = ll_attr_int(op, "aux", 0); const char* gep; const char* ptr_name; - ll_require_global_header_ptr(L, "global_load_lds", src_ptr); + rocke_ll_require_global_header_ptr(L, "global_load_lds", src_ptr, "ptr"); rocke_ll_need(L, "global.load.lds"); gep = rocke_ll_fresh(L, "gld_src"); rocke_ll_emitf(L, diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_llvm_internal.h b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_llvm_internal.h index 9e621a2bf7ec..4a25da9ef4c1 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_llvm_internal.h +++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_llvm_internal.h @@ -588,11 +588,38 @@ const char* rocke_ll_param_llvm_type(rocke_lower_t* L, const rocke_param_t* p); * _Lowerer._ptr_llvm_type). */ const char* rocke_ll_value_ptr_type(rocke_lower_t* L, const rocke_value_t* v); +/* The kernel parameter `v` names, or NULL for any other value (Python + * _Lowerer._param_of). */ +const rocke_param_t* rocke_ll_value_param(const rocke_lower_t* L, const rocke_value_t* v); + +/* Fail (ROCKE_ERR_TYPE, does not return) unless the module types `ptr` as + * `expected`, the LLVM pointer text of the `space` the op names (Python + * _Lowerer._require_header_ptr, same text). An op that names its pointer in + * one fixed space checks it here: a parameter's addr_space override (such as + * "constant") moves its header type away from the IR type the IR rules check, + * and the message blames the override only when it moved the header. + * `op_name` is the op's builder; `operand` names the pointer operand ("ptr", + * "src_ptr", ...). */ +void rocke_ll_require_header_ptr(rocke_lower_t* L, + const char* op_name, + const rocke_value_t* ptr, + const char* operand, + const char* space, + const char* expected); + +/* rocke_ll_require_header_ptr for an op that names its pointer + * ptr addrspace(1) (Python _Lowerer._require_global_header_ptr). */ +void rocke_ll_require_global_header_ptr(rocke_lower_t* L, + const char* op_name, + const rocke_value_t* ptr, + const char* operand); + /* Address space of a pointer operand, validated against `allowed` (Python * _Lowerer._anyptr_space). Also writes the matching pointer text to - * *out_ptr_ty when non-NULL. Fails (does not return) for a space the op does - * not accept: the space is part of an llvm_anyptr_ty intrinsic's overload and - * of a plain load, so the call site has to name the pointer's real type. */ + * *out_ptr_ty when non-NULL. Fails (ROCKE_ERR_TYPE, does not return) for a + * space the op does not accept: the space is part of an llvm_anyptr_ty + * intrinsic's overload and of a plain load, so the call site has to name the + * pointer's real type. */ int rocke_ll_anyptr_space(rocke_lower_t* L, const char* op, const rocke_value_t* ptr, diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/op_vocabulary.md b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/op_vocabulary.md index 304b7d5bf3ff..0b86ce2f1d34 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/op_vocabulary.md +++ b/dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/reference/op_vocabulary.md @@ -109,8 +109,9 @@ Complete reference of operations recognized by `core/ir.py` and lowered to AMDGP A pointer parameter declared `addr_space="constant"` is `ptr addrspace(4)` in the kernel header. The three global loads above read through it as `load addrspace(4)`; they refuse a flat or LDS pointer. Every other op that addresses global memory through it (the stores, atomics and flag ops above, -`cooperative_global_store`, `global_load_tr16_b128`, and the copies between global memory and LDS) is -refused by the LLVM lowering of both engines. +`cooperative_global_store`, `global_prefetch`, `global_load_tr16_b128`, the cluster loads, and the +copies between global memory and LDS) is refused by the LLVM lowering of both engines. Each of these +refusals is a `TypeError` (`ROCKE_ERR_TYPE` in the C++ engine) that names the op's builder. ### Global pointers diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py index 9db4a7f7d073..9554282b273b 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py +++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py @@ -2402,10 +2402,10 @@ def _op_tile_buffer_rsrc(self, op: Op) -> None: # ``raw_buffer_load/store`` family takes as its first argument. # Using the opaque type keeps the bitcast pun out of the user code. # - # Signature on ROCm 7 amdclang: + # Signature on ROCm 7.1 amdclang: # ``__builtin_amdgcn_make_buffer_rsrc(void*, short stride, - # int num_records, int flags)``; later clang takes a ``uint64_t`` - # ``num_records``. + # int num_records, int flags)``; later clang takes a 64-bit + # ``num_records`` (``int64_t`` in LLVM 22, ``uint64_t`` after). # # The flags word is the rsrc DWORD3. Its format / out-of-bounds # encoding is ISA-specific (the gfx9 value makes every raw buffer diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py index 1cbfd67dd772..9927d4ba5da2 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py +++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py @@ -1882,13 +1882,15 @@ def _anyptr_space( the call site has to name the pointer's real type. Naming a bare ``ptr`` for an ``addrspace(N)`` value is not a lax spelling -- LLVM rejects the module outright with "defined with type - 'ptr addrspace(N)' but expected 'ptr'". + 'ptr addrspace(N)' but expected 'ptr'". A pointer in another space is + refused with a ``TypeError``, as :meth:`_require_global_header_ptr` + refuses one. """ ty = self._ptr_llvm_type(ptr) for space, text in allowed.items(): if text == ty: return space, text - raise ValueError( + raise TypeError( f"{op}: pointer operand is {ty}, but the op accepts only " f"{', '.join(allowed[s] for s in sorted(allowed))}" ) @@ -1924,27 +1926,48 @@ def _ptr_llvm_type(self, v: Value) -> str: IR type at a call site instead would emit a type mismatch against the signature. """ - for p in self.kernel.params: - if f"%{p.name}" == v.name: - return _param_llvm_type(p) + param = self._param_of(v) + if param is not None: + return _param_llvm_type(param) return _llvm_type(v.type) - def _require_global_header_ptr(self, op_name: str, ptr: Value) -> None: - """Refuse ``op_name`` when the module types its ``ptr`` operand other - than ``ptr addrspace(1)``. + def _param_of(self, v: Value) -> Param | None: + """The kernel parameter ``v`` names, or ``None`` for any other value.""" + for p in self.kernel.params: + if f"%{p.name}" == v.name: + return p + return None - The global-pointer ops emit against ``ptr addrspace(1)``. The IR rules - check the IR type, but a parameter's ``addr_space`` override (such as - ``"constant"``) moves its header type away from it, and the emitted - instruction would then not match the header. The C++ engine - (``ll_require_global_header_ptr``) refuses with the same text. + def _require_header_ptr( + self, op_name: str, ptr: Value, operand: str, space: str, expected: str + ) -> None: + """Refuse ``op_name`` unless the module types its ``operand`` as + ``expected``, the LLVM pointer text of the ``space`` the op names. + + An op that names its pointer in one fixed space checks it here. The IR + rules check the IR type, but a parameter's ``addr_space`` override + (such as ``"constant"``) moves its header type away from it, and the + emitted instruction would then not match the header; the message + blames the override only when it moved the header. ``op_name`` is the + op's builder. The C++ engine (``rocke_ll_require_header_ptr``) refuses + with the same text. """ header = self._ptr_llvm_type(ptr) - if header != "ptr addrspace(1)": - raise TypeError( - f"{op_name} ptr must be a global pointer, got {header} " - "(its parameter's addr_space)" - ) + if header == expected: + return + param = self._param_of(ptr) + moved = param is not None and _param_llvm_type(param) != _llvm_type(param.type) + cause = " (its parameter's addr_space)" if moved else "" + raise TypeError( + f"{op_name} {operand} must be a {space} pointer, got {header}{cause}" + ) + + def _require_global_header_ptr( + self, op_name: str, ptr: Value, operand: str = "ptr" + ) -> None: + """:meth:`_require_header_ptr` for an op that names its pointer + ``ptr addrspace(1)``.""" + self._require_header_ptr(op_name, ptr, operand, "global", "ptr addrspace(1)") def _need(self, key: str) -> None: """Register a demand for an intrinsic declaration, and vet it. @@ -3438,7 +3461,7 @@ def _smem_global_name(self, smem_value: Value) -> Tuple[str, SmemType]: def _op_memref_global_load(self, op: Op) -> None: ptr, idx = op.operands - _, ptr_ty = self._anyptr_space("global_load", ptr, _GLOBAL_LOAD_PTR_TYPES) + _, ptr_ty = self._anyptr_space("global_load_f16", ptr, _GLOBAL_LOAD_PTR_TYPES) gep = self._fresh("gep") align = int(op.attrs.get("align", 2)) self._current().emit( @@ -3451,7 +3474,7 @@ def _op_memref_global_load(self, op: Op) -> None: def _op_memref_global_load_typed(self, op: Op) -> None: ptr, idx = op.operands - _, ptr_ty = self._anyptr_space("global_load_typed", ptr, _GLOBAL_LOAD_PTR_TYPES) + _, ptr_ty = self._anyptr_space("global_load", ptr, _GLOBAL_LOAD_PTR_TYPES) elem_ty = _llvm_type(op.result.type) gep = self._fresh("gep") align = int(op.attrs.get("align", 1)) @@ -3465,7 +3488,7 @@ def _op_memref_global_load_typed(self, op: Op) -> None: def _op_memref_global_store(self, op: Op) -> None: ptr, idx, val = op.operands - self._require_global_header_ptr("global_store", ptr) + self._require_global_header_ptr("store_f16", ptr) gep = self._fresh("gep") align = int(op.attrs.get("align", 2)) self._current().emit( @@ -3478,7 +3501,7 @@ def _op_memref_global_store(self, op: Op) -> None: def _op_memref_global_store_typed(self, op: Op) -> None: ptr, idx, val = op.operands - self._require_global_header_ptr("global_store_typed", ptr) + self._require_global_header_ptr("global_store", ptr) elem_ty = _llvm_type(val.type) gep = self._fresh("gep") align = int(op.attrs.get("align", 1)) @@ -5086,11 +5109,7 @@ def _op_tile_global_prefetch(self, op: Op) -> None: "global_prefetch", _OpTargetFact.GLOBAL_PREFETCH ) (ptr,) = op.operands - if self._ptr_llvm_type(ptr) != "ptr addrspace(1)": - raise TypeError( - f"global_prefetch ptr must be a global pointer, " - f"got {self._ptr_llvm_type(ptr)}" - ) + self._require_global_header_ptr("global_prefetch", ptr) cp = self._check_cachepolicy("global_prefetch", op.attrs.get("cachepolicy", 0)) self._need("global.prefetch") self._current().emit( @@ -5101,11 +5120,7 @@ def _op_tile_global_prefetch(self, op: Op) -> None: def _op_tile_flat_prefetch(self, op: Op) -> None: self._require_llvm23_target_fact("flat_prefetch", _OpTargetFact.GLOBAL_PREFETCH) (ptr,) = op.operands - if self._ptr_llvm_type(ptr) != "ptr": - raise TypeError( - f"flat_prefetch ptr must be a flat pointer, " - f"got {self._ptr_llvm_type(ptr)}" - ) + self._require_header_ptr("flat_prefetch", ptr, "ptr", "flat", "ptr") cp = self._check_cachepolicy("flat_prefetch", op.attrs.get("cachepolicy", 0)) self._need("flat.prefetch") self._current().emit( @@ -5612,10 +5627,9 @@ def _op_tile_global_store_async_from_lds(self, op: Op) -> None: if len(op.operands) != 2: raise ValueError("global_store_async_from_lds expects two operands") dst_ptr, lds_ptr = op.operands - if self._ptr_llvm_type(dst_ptr) != "ptr addrspace(1)": - raise TypeError( - "global_store_async_from_lds dst_ptr must be a global pointer" - ) + self._require_global_header_ptr( + "global_store_async_from_lds", dst_ptr, operand="dst_ptr" + ) width = int(op.attrs["width_bytes"]) if width not in (1, 4, 8, 16): raise ValueError( @@ -5648,8 +5662,9 @@ def _op_tile_global_load_tr16_b128(self, op: Op) -> None: if len(op.operands) != 1: raise ValueError("global_load_tr16_b128 expects one operand") src_ptr = op.operands[0] - if self._ptr_llvm_type(src_ptr) != "ptr addrspace(1)": - raise TypeError("global_load_tr16_b128 src_ptr must be a global pointer") + self._require_global_header_ptr( + "global_load_tr16_b128", src_ptr, operand="src_ptr" + ) dtype = str(op.attrs.get("dtype", "")) suffixes = {"f16": "v8f16", "bf16": "v8bf16", "i16": "v8i16"} if dtype not in suffixes: @@ -5684,8 +5699,7 @@ def _op_tile_cluster_load(self, op: Op) -> None: if len(op.operands) != 2 or len(op.results) != 1: raise ValueError("cluster_load expects two operands and one result") ptr, mask = op.operands - if self._ptr_llvm_type(ptr) != "ptr addrspace(1)": - raise TypeError("cluster_load ptr must be a global pointer") + self._require_global_header_ptr("cluster_load", ptr) if mask.type != I32: raise TypeError("cluster_load mask must be i32") width = int(op.attrs.get("width_bytes", 0)) @@ -5721,10 +5735,9 @@ def _op_tile_cluster_load_async_to_lds(self, op: Op) -> None: if len(op.operands) != 3: raise ValueError("cluster_load_async_to_lds expects three operands") src_ptr, lds_ptr, mask = op.operands - if self._ptr_llvm_type(src_ptr) != "ptr addrspace(1)": - raise TypeError( - "cluster_load_async_to_lds src_ptr must be a global pointer" - ) + self._require_global_header_ptr( + "cluster_load_async_to_lds", src_ptr, operand="src_ptr" + ) if mask.type != I32: raise TypeError("cluster_load_async_to_lds mask must be i32") width = int(op.attrs.get("width_bytes", 0)) diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/tests/engines.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/tests/engines.py index 9f070620372d..57043ae29df9 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/tests/engines.py +++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/tests/engines.py @@ -1,7 +1,7 @@ # Copyright (c) Advanced Micro Devices, Inc., or its affiliates. # SPDX-License-Identifier: MIT -"""Test helper: lower one kernel with both engines and require the same outcome. +"""Test helper: lower one kernel with each engine and require the same outcome. rocKE's Python and C++ engines must emit the same LLVM IR for the same kernel, and refuse the same kernel with the same message. ``ROCKE_BACKEND=both`` checks @@ -10,25 +10,45 @@ each engine directly, so a test states the parity it relies on whatever backend the run selects. -- ``lower_with_both_engines`` returns the LLVM IR both engines emit, after - asserting that the texts are identical. -- ``assert_both_engines_reject`` asserts that both engines refuse the kernel - with exactly ``message``. The C++ engine reports through the ``rocke_engine`` - binding as a ``RuntimeError`` whose text ends with ``": " + message``. -- ``lower_hip_with_both_engines`` and ``assert_both_engines_reject_hip`` are - the same two checks for the HIP-source lowering. - -All four skip, naming the reason, when the ``rocke_engine`` binding is not built. -``available_engines`` names the engines a test that drives each engine itself can -use: the installed wheel has no ``rocke_engine`` binding, so such a test runs the -Python engine alone there and both engines wherever the binding is supplied. -``all_engines_backend`` is the ``backend`` argument that runs those same engines -through ``compile_kernel`` and similar entry points. +Per engine (``engine`` is one of :data:`ENGINES`): + +- ``lower_with_engine`` returns the LLVM IR ``engine`` emits. For ``"cpp"`` it + also asserts that the text is identical to the Python engine's: IR parity is + part of what the C++ engine owes. A refusal has no output to compare, so + ``assert_engine_rejects`` checks the one engine named. +- ``assert_engine_rejects`` asserts that ``engine`` refuses the kernel with + exactly ``message``. The Python engine raises ``error_type``; the C++ engine + reports through the ``rocke_engine`` binding as a ``RuntimeError`` whose text + ends with ``"(status N): " + message``, where ``N`` is the + ``rocke_status_t`` that ``error_type`` maps to. +- ``lower_hip_with_engine`` and ``assert_engine_rejects_hip`` are the same two + checks for the HIP-source lowering. + +A test parametrized over :data:`ENGINES` checks the Python engine in every +lane, including the lanes that do not build the binding; only its ``"cpp"`` +cases skip there, naming the reason. + +Both engines at once: ``lower_with_both_engines``, +``assert_both_engines_reject``, ``lower_hip_with_both_engines`` and +``assert_both_engines_reject_hip`` run the same checks for the Python and the +C++ engine, and skip as a whole when the binding is not built. Their refusal +checks compare the C++ engine's message but not its status, which some older +refusals do not yet share with the Python engine's exception type. + +A test that loops over the engines itself, rather than through ``pytest`` +parameters, iterates ``available_engines()``: the part of :data:`ENGINES` this +run can drive. The installed wheel has no ``rocke_engine`` binding, so such a +test runs the Python engine alone there and both engines wherever the binding +is supplied. ``all_engines_backend`` is the ``backend`` argument that runs those +same engines through ``compile_kernel`` and similar entry points. """ from __future__ import annotations import importlib.util +from collections.abc import Callable +from types import ModuleType +from typing import Literal import pytest @@ -39,21 +59,38 @@ from rocke.core.lower_llvm import _lower_kernel_to_llvm_python __all__ = [ + "ENGINES", + "Engine", "all_engines_backend", "assert_both_engines_reject", "assert_both_engines_reject_hip", + "assert_engine_rejects", + "assert_engine_rejects_hip", "available_engines", "lower_hip_with_both_engines", + "lower_hip_with_engine", "lower_with_both_engines", + "lower_with_engine", ] +Engine = Literal["python", "cpp"] +ENGINES: tuple[Engine, ...] = ("python", "cpp") + +# The rocke_status_t (rocke/ir.h) each Python exception type maps to. +_NATIVE_STATUS: dict[type[Exception], int] = { + ValueError: 1, + TypeError: 2, + KeyError: 3, + NotImplementedError: 5, +} -def available_engines() -> tuple[str, ...]: - """``("python", "cpp")`` when the ``rocke_engine`` binding is importable, else + +def available_engines() -> tuple[Engine, ...]: + """:data:`ENGINES` when the ``rocke_engine`` binding is importable, else ``("python",)``.""" if importlib.util.find_spec("rocke_engine") is None: return ("python",) - return ("python", "cpp") + return ENGINES def all_engines_backend() -> str: @@ -62,51 +99,176 @@ def all_engines_backend() -> str: return "both" if "cpp" in available_engines() else "python" -def lower_with_both_engines( - kernel: KernelDef, arch: str, llvm_flavor: str | None = None +def _require_native_engine() -> ModuleType: + """The ``rocke_engine`` binding; skip the calling test when it is not built.""" + return pytest.importorskip("rocke_engine") + + +def _require_engine_name(engine: Engine) -> None: + if engine not in ENGINES: + raise ValueError(f"engine must be one of {ENGINES}, got {engine!r}") + + +def _native_status(error_type: type[Exception]) -> int: + if error_type not in _NATIVE_STATUS: + names = sorted(t.__name__ for t in _NATIVE_STATUS) + raise ValueError( + f"error_type must be one of {names}, got {error_type.__name__}" + ) + return _NATIVE_STATUS[error_type] + + +def _assert_python_rejects( + lower: Callable[[], str], error_type: type[Exception], message: str +) -> None: + with pytest.raises(error_type) as python_error: + lower() + assert str(python_error.value) == message + + +def _assert_native_rejects( + lower: Callable[[], str], message: str, status: int | None +) -> None: + """Assert the binding refuses with ``message``, after ``(status N): `` when + ``status`` is given.""" + with pytest.raises(RuntimeError) as native_error: + lower() + text = str(native_error.value) + suffix = f"(status {status}): {message}" if status is not None else f": {message}" + assert text.endswith(suffix), text + + +def _python_llvm(kernel: KernelDef, arch: str, llvm_flavor: str | None) -> str: + return _lower_kernel_to_llvm_python(kernel, arch=arch, llvm_flavor=llvm_flavor) + + +def _native_llvm(kernel: KernelDef, arch: str, llvm_flavor: str | None) -> str: + _require_native_engine() + return _lower_via_cpp_engine(serialize(kernel), arch, llvm_flavor) + + +def _native_hip(kernel: KernelDef, arch: str) -> str: + return _require_native_engine().lower_serialized_ir_hip(serialize(kernel), arch) + + +def lower_with_engine( + kernel: KernelDef, arch: str, *, engine: Engine, llvm_flavor: str | None = None ) -> str: - """Lower ``kernel`` for ``arch`` with each engine; return the shared IR text.""" - pytest.importorskip("rocke_engine") - python_ir = _lower_kernel_to_llvm_python(kernel, arch=arch, llvm_flavor=llvm_flavor) - native_ir = _lower_via_cpp_engine(serialize(kernel), arch, llvm_flavor) + """Lower ``kernel`` for ``arch`` with ``engine``; return its LLVM IR text.""" + _require_engine_name(engine) + python_ir = _python_llvm(kernel, arch, llvm_flavor) + if engine == "python": + return python_ir + native_ir = _native_llvm(kernel, arch, llvm_flavor) assert native_ir == python_ir, ( f"the C++ engine lowered kernel '{kernel.name}' for {arch} differently " "from the Python engine" ) - return python_ir + return native_ir -def assert_both_engines_reject( +def assert_engine_rejects( kernel: KernelDef, arch: str, + *, + engine: Engine, error_type: type[Exception], message: str, llvm_flavor: str | None = None, ) -> None: - """Assert that both engines refuse ``kernel`` for ``arch`` with ``message``. + """Assert that ``engine`` refuses ``kernel`` for ``arch`` with ``message``. The Python engine must raise ``error_type`` with exactly ``message``; the - C++ engine must fail with the same message after its binding prefix. + C++ engine must fail with ``error_type``'s status and the same message + after its binding prefix. """ - pytest.importorskip("rocke_engine") - with pytest.raises(error_type) as python_error: - _lower_kernel_to_llvm_python(kernel, arch=arch, llvm_flavor=llvm_flavor) - assert str(python_error.value) == message - with pytest.raises(RuntimeError) as native_error: - _lower_via_cpp_engine(serialize(kernel), arch, llvm_flavor) - assert str(native_error.value).endswith(": " + message), str(native_error.value) + _require_engine_name(engine) + if engine == "python": + _assert_python_rejects( + lambda: _python_llvm(kernel, arch, llvm_flavor), error_type, message + ) + else: + _assert_native_rejects( + lambda: _native_llvm(kernel, arch, llvm_flavor), + message, + _native_status(error_type), + ) -def lower_hip_with_both_engines(kernel: KernelDef, arch: str) -> str: - """Lower ``kernel`` to HIP source for ``arch`` with each engine; return the shared text.""" - engine = pytest.importorskip("rocke_engine") +def lower_hip_with_engine(kernel: KernelDef, arch: str, *, engine: Engine) -> str: + """Lower ``kernel`` to HIP source for ``arch`` with ``engine``; return the text. + + For ``"cpp"`` it also asserts that the text is identical to the Python + engine's. + """ + _require_engine_name(engine) python_hip = lower_kernel_to_hip(kernel, arch=arch) - native_hip = engine.lower_serialized_ir_hip(serialize(kernel), arch) + if engine == "python": + return python_hip + native_hip = _native_hip(kernel, arch) assert native_hip == python_hip, ( f"the C++ engine lowered kernel '{kernel.name}' to HIP for {arch} " "differently from the Python engine" ) - return python_hip + return native_hip + + +def assert_engine_rejects_hip( + kernel: KernelDef, + arch: str, + *, + engine: Engine, + error_type: type[Exception], + message: str, +) -> None: + """Assert that ``engine``'s HIP lowering refuses ``kernel`` with ``message``. + + Same contract as :func:`assert_engine_rejects`, for the HIP source. + """ + _require_engine_name(engine) + if engine == "python": + _assert_python_rejects( + lambda: lower_kernel_to_hip(kernel, arch=arch), error_type, message + ) + else: + _assert_native_rejects( + lambda: _native_hip(kernel, arch), message, _native_status(error_type) + ) + + +def lower_with_both_engines( + kernel: KernelDef, arch: str, llvm_flavor: str | None = None +) -> str: + """Lower ``kernel`` for ``arch`` with each engine; return the shared IR text.""" + _require_native_engine() + return lower_with_engine(kernel, arch, engine="cpp", llvm_flavor=llvm_flavor) + + +def assert_both_engines_reject( + kernel: KernelDef, + arch: str, + error_type: type[Exception], + message: str, + llvm_flavor: str | None = None, +) -> None: + """Assert that both engines refuse ``kernel`` for ``arch`` with ``message``. + + The C++ engine's status is not compared; see the module docstring. + """ + _require_native_engine() + _assert_python_rejects( + lambda: _python_llvm(kernel, arch, llvm_flavor), error_type, message + ) + _assert_native_rejects( + lambda: _native_llvm(kernel, arch, llvm_flavor), message, None + ) + + +def lower_hip_with_both_engines(kernel: KernelDef, arch: str) -> str: + """Lower ``kernel`` to HIP source for ``arch`` with each engine; return the + shared text.""" + _require_native_engine() + return lower_hip_with_engine(kernel, arch, engine="cpp") def assert_both_engines_reject_hip( @@ -117,12 +279,10 @@ def assert_both_engines_reject_hip( ) -> None: """Assert that both HIP lowerings refuse ``kernel`` for ``arch`` with ``message``. - Same contract as :func:`assert_both_engines_reject`, for the HIP source. + The C++ engine's status is not compared; see the module docstring. """ - engine = pytest.importorskip("rocke_engine") - with pytest.raises(error_type) as python_error: - lower_kernel_to_hip(kernel, arch=arch) - assert str(python_error.value) == message - with pytest.raises(RuntimeError) as native_error: - engine.lower_serialized_ir_hip(serialize(kernel), arch) - assert str(native_error.value).endswith(": " + message), str(native_error.value) + _require_native_engine() + _assert_python_rejects( + lambda: lower_kernel_to_hip(kernel, arch=arch), error_type, message + ) + _assert_native_rejects(lambda: _native_hip(kernel, arch), message, None) diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/future_intrinsic_lowering.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/future_intrinsic_lowering.cpp index 04dc02e6d136..9d42e63722a7 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/future_intrinsic_lowering.cpp +++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/future_intrinsic_lowering.cpp @@ -1731,11 +1731,13 @@ void case_gfx1250_data_prefetch_lowerer_rechecks() { struct Recheck { + rocke_status_t status; const char* want; void (*build)(rocke_ir_builder_t*); }; static const Recheck rechecks[] = { - {"global_prefetch cachepolicy must be in 0..31, got 40", + {ROCKE_ERR_VALUE, + "global_prefetch cachepolicy must be in 0..31, got 40", [](rocke_ir_builder_t* b) { emit_raw_op(b, ROCKE_OP_TILE_GLOBAL_PREFETCH, @@ -1743,7 +1745,8 @@ void case_gfx1250_data_prefetch_lowerer_rechecks() "cachepolicy", 40); }}, - {"flat_prefetch cachepolicy must be in 0..31, got -1", + {ROCKE_ERR_VALUE, + "flat_prefetch cachepolicy must be in 0..31, got -1", [](rocke_ir_builder_t* b) { emit_raw_op(b, ROCKE_OP_TILE_FLAT_PREFETCH, @@ -1751,14 +1754,18 @@ void case_gfx1250_data_prefetch_lowerer_rechecks() "cachepolicy", -1); }}, - {"s_setreg simm16 must fit an unsigned i16", + {ROCKE_ERR_VALUE, + "s_setreg simm16 must fit an unsigned i16", [](rocke_ir_builder_t* b) { emit_raw_op(b, ROCKE_OP_TILE_S_SETREG, rocke_b_const_i32(b, 1), "simm16", 70000); }}, /* The builder sees a global PtrType, but the param ABI moved it to * constant space: the header says ptr addrspace(4), which - * global.prefetch cannot take. */ - {"global_prefetch ptr must be a global pointer, got ptr addrspace(4)", + * global.prefetch cannot take: a type error, as for every op that + * names its pointer ptr addrspace(1). */ + {ROCKE_ERR_TYPE, + "global_prefetch ptr must be a global pointer, got ptr addrspace(4) " + "(its parameter's addr_space)", [](rocke_ir_builder_t* b) { rocke_param_opts_t constant{}; constant.addr_space = "constant"; @@ -1773,7 +1780,7 @@ void case_gfx1250_data_prefetch_lowerer_rechecks() std::string err; const rocke_status_t st = lower_expect_error("recheck", r.build, "gfx1250", ROCKE_LLVM_FLAVOR_LLVM23, &err); - if(st != ROCKE_ERR_VALUE || err.find(r.want) == std::string::npos) + if(st != r.status || err.find(r.want) == std::string::npos) { char msg[ROCKE_ERR_MSG_CAP + 128]; snprintf(msg, @@ -3418,6 +3425,158 @@ void case_buffer_rsrc_refusals() "/*num_records=*/(int)n, "); } +/* ---- global-memory ops on a parameter declared addr_space="constant" ---- */ + +/* A C++ sample of tests/core/test_constant_space_global_ops.py, for the lanes + * that do not build the Python binding. Such a parameter keeps its global IR + * type, but the kernel header types it ptr addrspace(4). */ +static rocke_value_t* constant_space_param(rocke_ir_builder_t* b, const rocke_type_t* elem) +{ + rocke_param_opts_t constant{}; + constant.addr_space = "constant"; + return rocke_b_param(b, "X", rocke_ptr_type(b, elem, "global"), &constant); +} + +void case_constant_space_global_ops() +{ + /* The global loads name the pointer's own type, so they read through it. */ + struct Load + { + const char* name; + void (*build)(rocke_ir_builder_t*); + const char* gep; + const char* load; + }; + static const Load loads[] = { + {"global_load_f16_constant", + [](rocke_ir_builder_t* b) { + rocke_value_t* x = constant_space_param(b, rocke_f16()); + rocke_value_t* y = global_ptr_param(b, "Y", rocke_f16()); + rocke_b_store_f16(b, + y, + rocke_b_const_i32(b, 0), + rocke_b_global_load_f16(b, x, rocke_b_const_i32(b, 1), 2)); + }, + "getelementptr inbounds half, ptr addrspace(4) %X, i32 ", + "= load half, ptr addrspace(4) %gep"}, + {"global_load_constant", + [](rocke_ir_builder_t* b) { + rocke_value_t* x = constant_space_param(b, rocke_f32()); + rocke_value_t* y = global_ptr_param(b, "Y", rocke_f32()); + rocke_b_global_store( + b, + y, + rocke_b_const_i32(b, 0), + rocke_b_global_load(b, x, rocke_b_const_i32(b, 1), rocke_f32(), 4), + 4); + }, + "getelementptr inbounds float, ptr addrspace(4) %X, i32 ", + "= load float, ptr addrspace(4) %gep"}, + {"global_load_vN_constant", + [](rocke_ir_builder_t* b) { + rocke_value_t* x = constant_space_param(b, rocke_f16()); + rocke_value_t* y = global_ptr_param(b, "Y", rocke_f16()); + rocke_b_global_store_vN( + b, + y, + rocke_b_const_i32(b, 0), + rocke_b_global_load_vN(b, x, rocke_b_const_i32(b, 4), rocke_f16(), 4, 0), + 4, + 0); + }, + "getelementptr inbounds half, ptr addrspace(4) %X, i32 ", + "= load <4 x half>, ptr addrspace(4) %gep"}, + }; + for(const Load& load : loads) + { + const std::string ir = lower_one(load.name, load.build, "gfx942", ROCKE_LLVM_FLAVOR_LLVM20); + expect_contains(ir, "(ptr addrspace(4) nocapture %X, ", __LINE__); + expect_contains(ir, load.gep, __LINE__); + expect_contains(ir, load.load, __LINE__); + EXPECT_NO_IR(ir, "ptr addrspace(1) %X"); + } + + /* Every refusal is a type error. A load names the spaces it accepts; an + * op that names ptr addrspace(1) blames the parameter's addr_space only + * when the parameter declares one. */ + struct Refusal + { + void (*build)(rocke_ir_builder_t*); + const char* want; + }; + static const Refusal refusals[] = { + {[](rocke_ir_builder_t* b) { + rocke_value_t* x = constant_space_param(b, rocke_f16()); + rocke_b_store_f16(b, + x, + rocke_b_const_i32(b, 0), + rocke_b_global_load_f16(b, x, rocke_b_const_i32(b, 1), 2)); + }, + "store_f16 ptr must be a global pointer, got ptr addrspace(4) " + "(its parameter's addr_space)"}, + {[](rocke_ir_builder_t* b) { + rocke_b_global_store(b, + constant_space_param(b, rocke_f32()), + rocke_b_const_i32(b, 0), + rocke_b_const_f32(b, 1.0), + 1); + }, + "global_store ptr must be a global pointer, got ptr addrspace(4) " + "(its parameter's addr_space)"}, + {[](rocke_ir_builder_t* b) { + rocke_b_global_atomic_add(b, + constant_space_param(b, rocke_i32()), + rocke_b_const_i32(b, 0), + rocke_b_const_i32(b, 1), + nullptr); + }, + "global_atomic_add ptr must be a global pointer, got ptr addrspace(4) " + "(its parameter's addr_space)"}, + {[](rocke_ir_builder_t* b) { + rocke_value_t* x + = rocke_b_param(b, "X", rocke_ptr_type(b, rocke_f32(), "lds"), nullptr); + rocke_b_global_store(b, x, rocke_b_const_i32(b, 0), rocke_b_const_f32(b, 1.0), 1); + }, + "global_store ptr must be a global pointer, got ptr addrspace(3)"}, + {[](rocke_ir_builder_t* b) { + rocke_value_t* x + = rocke_b_param(b, "X", rocke_ptr_type(b, rocke_f32(), "lds"), nullptr); + rocke_value_t* y = global_ptr_param(b, "Y", rocke_f32()); + rocke_b_global_store( + b, + y, + rocke_b_const_i32(b, 0), + rocke_b_global_load(b, x, rocke_b_const_i32(b, 1), rocke_f32(), 4), + 4); + }, + "global_load: pointer operand is ptr addrspace(3), but the op accepts only " + "ptr addrspace(1), ptr addrspace(4)"}, + }; + for(const Refusal& refusal : refusals) + { + std::string err; + const rocke_status_t st = lower_expect_error( + "refusal", refusal.build, "gfx942", ROCKE_LLVM_FLAVOR_LLVM20, &err); + /* The message must end with the text: no " (its parameter's + * addr_space)" may follow a refusal that does not expect it. */ + const std::string want(refusal.want); + const bool ends_with_want + = err.size() >= want.size() + && err.compare(err.size() - want.size(), want.size(), want) == 0; + if(st != ROCKE_ERR_TYPE || !ends_with_want) + { + char msg[ROCKE_ERR_MSG_CAP + 256]; + snprintf(msg, + sizeof(msg), + "want ROCKE_ERR_TYPE \"%s\", status %d, err \"%s\"", + refusal.want, + (int)st, + err.c_str()); + fail(msg, __LINE__); + } + } +} + /* ---- dtype-generic buffer_load_vN / buffer_store_vN (Python IRBuilder) ---- */ /* Run `build` on a fresh builder; return the refusal text, or "" when the @@ -4200,6 +4359,7 @@ const TestCase k_cases[] = { {"vmcnt_bits", case_vmcnt_bits}, {"buffer_load_lds_async", case_buffer_load_lds_async}, {"buffer_rsrc_refusals", case_buffer_rsrc_refusals}, + {"constant_space_global_ops", case_constant_space_global_ops}, {"buffer_load_vN_generic", case_buffer_load_vN_generic}, {"buffer_store_vN_generic", case_buffer_store_vN_generic}, {"buffer_vN_verify", case_buffer_vN_verify}, diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_operands.py b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_operands.py index 7cda00da13e2..cee9e46758df 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_operands.py +++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_operands.py @@ -21,24 +21,32 @@ 'i32'"). The HIP lowering calls ``__builtin_amdgcn_make_buffer_rsrc`` with -``num_records`` cast to ``(int)`` (ROCm 7's builtin takes an ``int`` there, -later clang a ``uint64_t``). The cast would truncate an i64 ``num_bytes``, so -the HIP lowering refuses any width but i32. It keeps a constant-space pointer: the HIP signature declares every pointer -parameter the same way, and the builtin takes a ``void*``. +``num_records`` cast to ``(int)`` (ROCm 7.1's builtin takes an ``int`` there, +later clang a 64-bit integer: ``int64_t`` in LLVM 22, ``uint64_t`` after). The +cast would truncate an i64 ``num_bytes``, so the HIP lowering refuses any width +but i32. It keeps a constant-space pointer: the HIP signature declares every +pointer parameter the same way, and the builtin takes a ``void*``. + +Each case runs per engine, so the Python engine is checked in every lane; the +C++ cases skip where the ``rocke_engine`` binding is not built. There, +``future_intrinsic_lowering``'s ``buffer_rsrc_refusals`` CTest case checks the +C++ engine's refusals and its i32 / i64 lowerings. """ from __future__ import annotations import pytest from rocke.core.ir import F32, I32, I64, IRBuilder, PtrType +from rocke.core.lower_llvm import LLVM_FLAVORS from rocke.tests.engines import ( - assert_both_engines_reject, - assert_both_engines_reject_hip, - lower_hip_with_both_engines, - lower_with_both_engines, + ENGINES, + assert_engine_rejects, + assert_engine_rejects_hip, + lower_hip_with_engine, + lower_with_engine, ) -FLAVORS = ("llvm20", "llvm22", "llvm23", "llvm24") +FLAVORS_AFTER_LLVM20 = tuple(f for f in LLVM_FLAVORS if f != "llvm20") ARCH = "gfx942" @@ -58,35 +66,40 @@ def _rsrc_kernel(name, *, constant, num_bytes_type): return b.kernel -@pytest.mark.parametrize("flavor", FLAVORS) -def test_both_llvm_lowerings_refuse_a_constant_space_pointer(flavor): +@pytest.mark.parametrize("engine", ENGINES) +@pytest.mark.parametrize("flavor", LLVM_FLAVORS) +def test_the_llvm_lowering_refuses_a_constant_space_pointer(flavor, engine): kernel = _rsrc_kernel("rsrc_constant", constant=True, num_bytes_type=I32) - assert_both_engines_reject( + assert_engine_rejects( kernel, ARCH, - TypeError, - "buffer_rsrc ptr must be a global pointer, got ptr addrspace(4) " + engine=engine, + error_type=TypeError, + message="buffer_rsrc ptr must be a global pointer, got ptr addrspace(4) " "(its parameter's addr_space)", llvm_flavor=flavor, ) -def test_both_llvm_lowerings_refuse_an_i64_num_bytes_at_llvm20(): +@pytest.mark.parametrize("engine", ENGINES) +def test_the_llvm_lowering_refuses_an_i64_num_bytes_at_llvm20(engine): kernel = _rsrc_kernel("rsrc_i64", constant=False, num_bytes_type=I64) - assert_both_engines_reject( + assert_engine_rejects( kernel, ARCH, - ValueError, - "tile.buffer_rsrc num_bytes must be i32 at llvm20, got i64; the llvm20 " + engine=engine, + error_type=ValueError, + message="tile.buffer_rsrc num_bytes must be i32 at llvm20, got i64; the llvm20 " "make.buffer.rsrc takes an i32 num_records, so pass the byte count as i32", llvm_flavor="llvm20", ) -@pytest.mark.parametrize("flavor", FLAVORS) -def test_an_i32_num_bytes_lowers_at_every_flavor(flavor): +@pytest.mark.parametrize("engine", ENGINES) +@pytest.mark.parametrize("flavor", LLVM_FLAVORS) +def test_an_i32_num_bytes_lowers_at_every_flavor(flavor, engine): kernel = _rsrc_kernel("rsrc_i32", constant=False, num_bytes_type=I32) - ir = lower_with_both_engines(kernel, ARCH, llvm_flavor=flavor) + ir = lower_with_engine(kernel, ARCH, engine=engine, llvm_flavor=flavor) if flavor == "llvm20": assert ( "@llvm.amdgcn.make.buffer.rsrc.p1(ptr addrspace(1) %X, i16 0, i32 %n," in ir @@ -99,33 +112,37 @@ def test_an_i32_num_bytes_lowers_at_every_flavor(flavor): ) -@pytest.mark.parametrize("flavor", FLAVORS[1:]) -def test_an_i64_num_bytes_lowers_after_llvm20(flavor): +@pytest.mark.parametrize("engine", ENGINES) +@pytest.mark.parametrize("flavor", FLAVORS_AFTER_LLVM20) +def test_an_i64_num_bytes_lowers_after_llvm20(flavor, engine): kernel = _rsrc_kernel("rsrc_i64", constant=False, num_bytes_type=I64) - ir = lower_with_both_engines(kernel, ARCH, llvm_flavor=flavor) + ir = lower_with_engine(kernel, ARCH, engine=engine, llvm_flavor=flavor) assert ( "@llvm.amdgcn.make.buffer.rsrc.p8.p1(ptr addrspace(1) %X, i16 0, i64 %n," in ir ) -def test_both_hip_lowerings_refuse_an_i64_num_bytes(): +@pytest.mark.parametrize("engine", ENGINES) +def test_the_hip_lowering_refuses_an_i64_num_bytes(engine): kernel = _rsrc_kernel("rsrc_i64", constant=False, num_bytes_type=I64) - assert_both_engines_reject_hip( + assert_engine_rejects_hip( kernel, ARCH, - ValueError, - "tile.buffer_rsrc num_bytes must be i32 in the HIP lowering, got i64; " + engine=engine, + error_type=ValueError, + message="tile.buffer_rsrc num_bytes must be i32 in the HIP lowering, got i64; " "the HIP lowering casts num_records to (int), so a wider byte count " "would be truncated; pass it as i32", ) +@pytest.mark.parametrize("engine", ENGINES) @pytest.mark.parametrize("constant", (False, True)) def test_the_hip_lowering_keeps_an_i32_num_bytes_and_a_constant_space_pointer( - constant, + constant, engine ): kernel = _rsrc_kernel("rsrc_hip", constant=constant, num_bytes_type=I32) - hip = lower_hip_with_both_engines(kernel, ARCH) + hip = lower_hip_with_engine(kernel, ARCH, engine=engine) assert "void rsrc_hip(float* __restrict__ X, float* __restrict__ Y, int n)" in hip assert ( "__builtin_amdgcn_make_buffer_rsrc((void*)X, /*stride=*/(short)0, " diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_constant_space_global_ops.py b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_constant_space_global_ops.py index 8b899fb27435..cd3c513c3745 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_constant_space_global_ops.py +++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_constant_space_global_ops.py @@ -12,10 +12,20 @@ and the module compiles. They refuse a pointer in any other space (flat or LDS), which they used to name ``ptr addrspace(1)`` as well. - Every other op that addresses global memory through the pointer -- stores, - atomics, flag ops, the cooperative store, the transposed load and the copies - between global memory and LDS -- is refused by the LLVM lowering of both - engines. Most use the text the ``global_ptr_*`` ops use; the two ops that - already refused a non-global pointer keep their own text. + atomics, flag ops, the cooperative store, prefetch, the transposed and + cluster loads, and the copies between global memory and LDS -- is refused by + the LLVM lowering of both engines, with one text that names the op's builder + and the operand. ``flat_prefetch`` refuses a parameter whose override moved + it out of the flat space the same way. +- Every one of these refusals is a ``TypeError`` (``ROCKE_ERR_TYPE`` in the + C++ engine). The text blames the parameter's ``addr_space`` only when the + override moved its header type. + +Each case runs per engine, so the Python engine is checked in every lane; the +C++ cases skip where the ``rocke_engine`` binding is not built. There, +``future_intrinsic_lowering``'s ``constant_space_global_ops`` CTest case +checks a sample of the C++ engine's side: the three loads and the +``store_f16``, ``global_store`` and ``global_atomic_add`` refusals. """ from __future__ import annotations @@ -35,10 +45,10 @@ Type, Value, ) -from rocke.tests.engines import assert_both_engines_reject, lower_with_both_engines +from rocke.core.lower_llvm import LLVM_FLAVORS +from rocke.tests.engines import ENGINES, assert_engine_rejects, lower_with_engine from rocke.tests.ir_validity import assert_target_accepts -FLAVORS = ("llvm20", "llvm22", "llvm23", "llvm24") ARCH = "gfx942" @@ -76,14 +86,15 @@ def _load_kernel( return b.kernel +# builder -> (the load applied to X, element type of X, GEP element, loaded type) LOADS = { - "global_load": ( + "global_load_f16": ( lambda b, x: b.global_load_f16(x, b.const_i32(1)), F16, "half", "half", ), - "global_load_typed": ( + "global_load": ( lambda b, x: b.global_load(x, b.const_i32(1), F32, align=4), F32, "float", @@ -98,40 +109,50 @@ def _load_kernel( } -@pytest.mark.parametrize("flavor", FLAVORS) -@pytest.mark.parametrize("op_name", sorted(LOADS)) -def test_both_llvm_lowerings_load_through_a_constant_space_parameter(op_name, flavor): - load, dtype, elem_ty, loaded_ty = LOADS[op_name] - ir = lower_with_both_engines( - _load_kernel(f"{op_name}_constant", load, dtype), ARCH, llvm_flavor=flavor +@pytest.mark.parametrize("engine", ENGINES) +@pytest.mark.parametrize("flavor", LLVM_FLAVORS) +@pytest.mark.parametrize("builder", sorted(LOADS)) +def test_a_load_reads_through_a_constant_space_parameter(builder, flavor, engine): + load, dtype, elem_ty, loaded_ty = LOADS[builder] + ir = lower_with_engine( + _load_kernel(f"{builder}_constant", load, dtype), + ARCH, + engine=engine, + llvm_flavor=flavor, ) - assert f"@{op_name}_constant(ptr addrspace(4) nocapture %X, " in ir + assert f"@{builder}_constant(ptr addrspace(4) nocapture %X, " in ir assert f"getelementptr inbounds {elem_ty}, ptr addrspace(4) %X, i32 " in ir assert f"= load {loaded_ty}, ptr addrspace(4) %gep" in ir assert "ptr addrspace(1) %X" not in ir -@pytest.mark.parametrize("op_name", sorted(LOADS)) -def test_a_load_through_a_constant_space_parameter_compiles(op_name): - load, dtype, _, _ = LOADS[op_name] - ir = lower_with_both_engines(_load_kernel(f"{op_name}_compiles", load, dtype), ARCH) +@pytest.mark.parametrize("builder", sorted(LOADS)) +def test_a_load_through_a_constant_space_parameter_compiles(builder): + """The IR is the same for both engines (the test above), so one compile + judges it.""" + load, dtype, _, _ = LOADS[builder] + ir = lower_with_engine( + _load_kernel(f"{builder}_compiles", load, dtype), ARCH, engine="python" + ) assert_target_accepts(ir, ARCH) +@pytest.mark.parametrize("engine", ENGINES) @pytest.mark.parametrize( "x_space, header_type", [("private", "ptr"), ("lds", "ptr addrspace(3)")] ) -@pytest.mark.parametrize("op_name", sorted(LOADS)) -def test_both_llvm_lowerings_refuse_a_load_through_another_space( - op_name, x_space, header_type +@pytest.mark.parametrize("builder", sorted(LOADS)) +def test_a_load_refuses_a_pointer_in_another_space( + builder, x_space, header_type, engine ): - load, dtype, _, _ = LOADS[op_name] - kernel = _load_kernel(f"{op_name}_{x_space}", load, dtype, x_space) - assert_both_engines_reject( + load, dtype, _, _ = LOADS[builder] + kernel = _load_kernel(f"{builder}_{x_space}", load, dtype, x_space) + assert_engine_rejects( kernel, ARCH, - ValueError, - f"{op_name}: pointer operand is {header_type}, but the op accepts only " + engine=engine, + error_type=TypeError, + message=f"{builder}: pointer operand is {header_type}, but the op accepts only " "ptr addrspace(1), ptr addrspace(4)", ) @@ -141,9 +162,9 @@ def test_both_llvm_lowerings_refuse_a_load_through_another_space( # -------------------------------------------------------------------------- -def _header_refusal(op_name: str) -> str: +def _header_refusal(builder: str, operand: str = "ptr") -> str: return ( - f"{op_name} ptr must be a global pointer, got ptr addrspace(4) " + f"{builder} {operand} must be a global pointer, got ptr addrspace(4) " "(its parameter's addr_space)" ) @@ -152,23 +173,23 @@ def _lds_address(b: IRBuilder, dtype: Type) -> Value: return b.smem_addr_of(b.smem_alloc(dtype, [64], name_hint="stage")) -# op name -> (element type of X, the op applied to X, arch, flavor, message) +# builder -> (element type of X, the op applied to X, arch, flavor, message) REFUSALS = { - "global_store": ( + "store_f16": ( F16, lambda b, x: b.store_f16( x, b.const_i32(0), b.global_load_f16(x, b.const_i32(1)) ), ARCH, None, - _header_refusal("global_store"), + _header_refusal("store_f16"), ), - "global_store_typed": ( + "global_store": ( F32, lambda b, x: b.global_store(x, b.const_i32(0), b.const_f32(1.0)), ARCH, None, - _header_refusal("global_store_typed"), + _header_refusal("global_store"), ), "global_store_vN": ( F16, @@ -257,22 +278,89 @@ def _lds_address(b: IRBuilder, dtype: Type) -> Value: ), "gfx1250", "llvm23", - "global_store_async_from_lds dst_ptr must be a global pointer", + _header_refusal("global_store_async_from_lds", "dst_ptr"), ), "global_load_tr16_b128": ( I16, lambda b, x: b.global_load_tr16_b128(x, dtype=I16), "gfx1250", "llvm23", - "global_load_tr16_b128 src_ptr must be a global pointer", + _header_refusal("global_load_tr16_b128", "src_ptr"), + ), + "global_prefetch": ( + I32, + lambda b, x: b.global_prefetch(x), + "gfx1250", + "llvm23", + _header_refusal("global_prefetch"), + ), + "cluster_load": ( + I32, + lambda b, x: b.cluster_load(x, b.const_i32(1)), + "gfx1250", + "llvm23", + _header_refusal("cluster_load"), + ), + "cluster_load_async_to_lds": ( + I32, + lambda b, x: b.cluster_load_async_to_lds( + x, _lds_address(b, I32), b.const_i32(1), width_bytes=4 + ), + "gfx1250", + "llvm23", + _header_refusal("cluster_load_async_to_lds", "src_ptr"), ), } -@pytest.mark.parametrize("op_name", sorted(REFUSALS)) -def test_both_llvm_lowerings_refuse_a_constant_space_parameter(op_name): - dtype, apply_op, arch, flavor, message = REFUSALS[op_name] - b = IRBuilder(f"{op_name}_constant") +@pytest.mark.parametrize("engine", ENGINES) +@pytest.mark.parametrize("builder", sorted(REFUSALS)) +def test_an_op_that_names_a_global_pointer_refuses_a_constant_space_parameter( + builder, engine +): + dtype, apply_op, arch, flavor, message = REFUSALS[builder] + b = IRBuilder(f"{builder}_constant") apply_op(b, _constant_param(b, dtype)) b.ret() - assert_both_engines_reject(b.kernel, arch, TypeError, message, llvm_flavor=flavor) + assert_engine_rejects( + b.kernel, + arch, + engine=engine, + error_type=TypeError, + message=message, + llvm_flavor=flavor, + ) + + +@pytest.mark.parametrize("engine", ENGINES) +def test_the_refusal_blames_addr_space_only_when_the_override_moved_it(engine): + """An LDS parameter without an ``addr_space`` override reaches the store as + ``ptr addrspace(3)``; its own type is the cause, not an override.""" + b = IRBuilder("store_lds_param") + x = b.param("X", PtrType(F32, "lds")) + b.global_store(x, b.const_i32(0), b.const_f32(1.0)) + b.ret() + assert_engine_rejects( + b.kernel, + ARCH, + engine=engine, + error_type=TypeError, + message="global_store ptr must be a global pointer, got ptr addrspace(3)", + ) + + +@pytest.mark.parametrize("engine", ENGINES) +def test_flat_prefetch_refuses_a_parameter_moved_to_constant_space(engine): + b = IRBuilder("flat_prefetch_constant") + x = b.param("X", PtrType(I32, "private"), addr_space="constant") + b.flat_prefetch(x) + b.ret() + assert_engine_rejects( + b.kernel, + "gfx1250", + engine=engine, + error_type=TypeError, + message="flat_prefetch ptr must be a flat pointer, got ptr addrspace(4) " + "(its parameter's addr_space)", + llvm_flavor="llvm23", + ) diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_gfx1250_prefetch.py b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_gfx1250_prefetch.py index c147071374ee..2b4c84cc377a 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_gfx1250_prefetch.py +++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_gfx1250_prefetch.py @@ -239,7 +239,7 @@ def test_s_prefetch_data_rejects_lds_pointer(self): p = b.param("p", PtrType(I64, "lds")) b.s_prefetch_data(p, b.const_i32(1)) b.ret() - with self.assertRaisesRegex(ValueError, "s_prefetch_data: pointer operand"): + with self.assertRaisesRegex(TypeError, "s_prefetch_data: pointer operand"): _lower(b.kernel) def test_distinct_from_instruction_prefetch(self): diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/test_rocke.py b/dnn-providers/hip-kernel-provider/rocke/platform/tests/test_rocke.py index 5d38b0aa8e09..860cfc1239bd 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/test_rocke.py +++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/test_rocke.py @@ -4688,7 +4688,7 @@ def test_av_b128_rejects_lds_pointer(self): b.av_load_b128(p) # Permissive C++ mode falls back even when the native engine rejects IR. error_type = ( - RuntimeError if resolve_backend() == "cpp" and _cpp_strict() else ValueError + RuntimeError if resolve_backend() == "cpp" and _cpp_strict() else TypeError ) with self.assertRaisesRegex( error_type, @@ -4793,7 +4793,7 @@ def test_s_prefetch_inst_rejects_lds_pointer(self): p = b.param("p", PtrType(I32, "lds"), align=4) b.s_prefetch_inst(p, b.const_i32(64)) error_type = ( - RuntimeError if resolve_backend() == "cpp" and _cpp_strict() else ValueError + RuntimeError if resolve_backend() == "cpp" and _cpp_strict() else TypeError ) with self.assertRaisesRegex( error_type, @@ -4907,7 +4907,7 @@ def test_lds_rejection_without_cpp_binding(monkeypatch, strict, op): error_type, message = BackendError, "rocke_engine.*not importable" else: error_type, message = ( - ValueError, + TypeError, rf"{op}: pointer operand is ptr addrspace\(3\)", ) with pytest.raises(error_type, match=message): From f99650e091fcfe1b9f412bca1d6e0cddd6aebdb1 Mon Sep 17 00:00:00 2001 From: Chris Millette Date: Wed, 4 Nov 2026 11:39:18 -0500 Subject: [PATCH 6/7] fix(rocke): name the builder and operand in buffer_rsrc and LDS-copy refusals The buffer_rsrc byte-count refusals now name the builder (buffer_rsrc) rather than the op (tile.buffer_rsrc), and global_load_lds and global_load_async_to_lds name their pointer operand src_ptr, as their builders do. Both engines change together. The HIP lowering's comment now says the builtin takes an int64_t num_records in LLVM 22 and 23. Co-Authored-By: Claude --- .../platform/cpp/core/lower_hip/lower_hip_mem.cpp | 2 +- .../rocke/platform/cpp/core/lower_llvm/mem.cpp | 8 ++++---- .../rocke/platform/python/rocke/core/lower_hip.py | 4 ++-- .../rocke/platform/python/rocke/core/lower_llvm.py | 10 ++++++---- .../platform/tests/core/future_intrinsic_lowering.cpp | 4 ++-- .../platform/tests/core/test_buffer_rsrc_operands.py | 4 ++-- .../tests/core/test_constant_space_global_ops.py | 4 ++-- 7 files changed, 19 insertions(+), 17 deletions(-) diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp index afa2bcb60d55..47ad62883b59 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp +++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp @@ -1086,7 +1086,7 @@ static rocke_status_t _op_tile_buffer_rsrc(rocke_h_lowerer_t* lw, const rocke_op { return rocke_h_fail(lw, ROCKE_ERR_VALUE, - "tile.buffer_rsrc num_bytes must be i32 in the HIP lowering, got " + "buffer_rsrc num_bytes must be i32 in the HIP lowering, got " "%s; the HIP lowering casts num_records to (int), so a wider byte " "count would be truncated; pass it as i32", nb_ty); diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mem.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mem.cpp index d178acc0767c..f419165b5b72 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mem.cpp +++ b/dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mem.cpp @@ -1044,7 +1044,7 @@ static void op_tile_buffer_rsrc(rocke_lower_t* L, const rocke_op_t* op) else { rocke_ll_fail( - L, ROCKE_ERR_VALUE, "tile.buffer_rsrc num_bytes must be i32 or i64, got %s", nb_ty); + L, ROCKE_ERR_VALUE, "buffer_rsrc num_bytes must be i32 or i64, got %s", nb_ty); } nb_text = rocke_arena_printf(&L->arena, "i64 %s", nb_arg); } @@ -1059,7 +1059,7 @@ static void op_tile_buffer_rsrc(rocke_lower_t* L, const rocke_op_t* op) rocke_ll_fail( L, ROCKE_ERR_VALUE, - "tile.buffer_rsrc num_bytes must be i32 at llvm20, got %s; the llvm20 " + "buffer_rsrc num_bytes must be i32 at llvm20, got %s; the llvm20 " "make.buffer.rsrc takes an i32 num_records, so pass the byte count as i32", nb_ty); } @@ -1690,7 +1690,7 @@ static void op_tile_global_load_async_to_lds(rocke_lower_t* L, const rocke_op_t* const char* idx_ty; rocke_strbuf_t gep; int i; - rocke_ll_require_global_header_ptr(L, "global_load_async_to_lds", src_ptr, "ptr"); + rocke_ll_require_global_header_ptr(L, "global_load_async_to_lds", src_ptr, "src_ptr"); if(width == 1) suffix = "b8"; @@ -2049,7 +2049,7 @@ static void op_tile_global_load_lds(rocke_lower_t* L, const rocke_op_t* op) int64_t aux = ll_attr_int(op, "aux", 0); const char* gep; const char* ptr_name; - rocke_ll_require_global_header_ptr(L, "global_load_lds", src_ptr, "ptr"); + rocke_ll_require_global_header_ptr(L, "global_load_lds", src_ptr, "src_ptr"); rocke_ll_need(L, "global.load.lds"); gep = rocke_ll_fresh(L, "gld_src"); rocke_ll_emitf(L, diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py index 9554282b273b..1009727d5349 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py +++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py @@ -2405,7 +2405,7 @@ def _op_tile_buffer_rsrc(self, op: Op) -> None: # Signature on ROCm 7.1 amdclang: # ``__builtin_amdgcn_make_buffer_rsrc(void*, short stride, # int num_records, int flags)``; later clang takes a 64-bit - # ``num_records`` (``int64_t`` in LLVM 22, ``uint64_t`` after). + # ``num_records``: ``int64_t`` in LLVM 22 and 23, ``uint64_t`` after. # # The flags word is the rsrc DWORD3. Its format / out-of-bounds # encoding is ISA-specific (the gfx9 value makes every raw buffer @@ -2425,7 +2425,7 @@ def _op_tile_buffer_rsrc(self, op: Op) -> None: nb_ty = op.operands[1].type.name if nb_ty != "i32": raise ValueError( - f"tile.buffer_rsrc num_bytes must be i32 in the HIP lowering, got " + f"buffer_rsrc num_bytes must be i32 in the HIP lowering, got " f"{nb_ty}; the HIP lowering casts num_records to (int), so a wider " "byte count would be truncated; pass it as i32" ) diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py index 9927d4ba5da2..6d89cd65a1db 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py +++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py @@ -5590,7 +5590,9 @@ def _op_tile_global_load_async_to_lds(self, op: Op) -> None: src_index = op.operands[1] lds_smem = op.operands[2] lds_indices = op.operands[3:] - self._require_global_header_ptr("global_load_async_to_lds", src_ptr) + self._require_global_header_ptr( + "global_load_async_to_lds", src_ptr, operand="src_ptr" + ) width = int(op.attrs["width_bytes"]) cpol = int(op.attrs.get("cpol", 0)) ioff = int(op.attrs.get("offset_bytes", 0)) @@ -6041,7 +6043,7 @@ def _op_tile_buffer_rsrc(self, op: Op) -> None: ) else: raise ValueError( - f"tile.buffer_rsrc num_bytes must be i32 or i64, got {nb_ty}" + f"buffer_rsrc num_bytes must be i32 or i64, got {nb_ty}" ) nb_text = f"i64 {nb_arg}" else: @@ -6049,7 +6051,7 @@ def _op_tile_buffer_rsrc(self, op: Op) -> None: nb_ty = _llvm_type(num_bytes.type) if nb_ty != "i32": raise ValueError( - f"tile.buffer_rsrc num_bytes must be i32 at llvm20, got {nb_ty}; " + f"buffer_rsrc num_bytes must be i32 at llvm20, got {nb_ty}; " "the llvm20 make.buffer.rsrc takes an i32 num_records, so pass " "the byte count as i32" ) @@ -6466,7 +6468,7 @@ def _op_tile_buffer_load_lds_async(self, op: Op) -> None: def _op_tile_global_load_lds(self, op: Op) -> None: src_ptr, byte_off, lds_addr = op.operands - self._require_global_header_ptr("global_load_lds", src_ptr) + self._require_global_header_ptr("global_load_lds", src_ptr, operand="src_ptr") size_bytes = int(op.attrs["size_bytes"]) aux = int(op.attrs.get("aux", 0)) self._need("global.load.lds") diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/future_intrinsic_lowering.cpp b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/future_intrinsic_lowering.cpp index 9d42e63722a7..b8bf14886d06 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/future_intrinsic_lowering.cpp +++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/future_intrinsic_lowering.cpp @@ -3379,7 +3379,7 @@ void case_buffer_rsrc_refusals() ROCKE_LLVM_FLAVOR_LLVM20, &err); if(st != ROCKE_ERR_VALUE - || err.find("tile.buffer_rsrc num_bytes must be i32 at llvm20, got i64; the llvm20 " + || err.find("buffer_rsrc num_bytes must be i32 at llvm20, got i64; the llvm20 " "make.buffer.rsrc takes an i32 num_records, so pass the byte count as i32") == std::string::npos) fail(("i64 num_bytes at llvm20: " + err).c_str(), __LINE__); @@ -3407,7 +3407,7 @@ void case_buffer_rsrc_refusals() "gfx942", &hip_err); if(hip_st != ROCKE_ERR_VALUE - || hip_err.find("tile.buffer_rsrc num_bytes must be i32 in the HIP lowering, got i64; " + || hip_err.find("buffer_rsrc num_bytes must be i32 in the HIP lowering, got i64; " "the HIP lowering casts num_records to (int), so a wider byte count " "would be truncated; pass it as i32") == std::string::npos) diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_operands.py b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_operands.py index cee9e46758df..c3b535445e71 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_operands.py +++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_operands.py @@ -89,7 +89,7 @@ def test_the_llvm_lowering_refuses_an_i64_num_bytes_at_llvm20(engine): ARCH, engine=engine, error_type=ValueError, - message="tile.buffer_rsrc num_bytes must be i32 at llvm20, got i64; the llvm20 " + message="buffer_rsrc num_bytes must be i32 at llvm20, got i64; the llvm20 " "make.buffer.rsrc takes an i32 num_records, so pass the byte count as i32", llvm_flavor="llvm20", ) @@ -130,7 +130,7 @@ def test_the_hip_lowering_refuses_an_i64_num_bytes(engine): ARCH, engine=engine, error_type=ValueError, - message="tile.buffer_rsrc num_bytes must be i32 in the HIP lowering, got i64; " + message="buffer_rsrc num_bytes must be i32 in the HIP lowering, got i64; " "the HIP lowering casts num_records to (int), so a wider byte count " "would be truncated; pass it as i32", ) diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_constant_space_global_ops.py b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_constant_space_global_ops.py index cd3c513c3745..ee0d2ecfeef3 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_constant_space_global_ops.py +++ b/dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_constant_space_global_ops.py @@ -256,7 +256,7 @@ def _lds_address(b: IRBuilder, dtype: Type) -> Value: lambda b, x: b.global_load_lds(x, b.const_i32(0), _lds_address(b, F32), 4), ARCH, None, - _header_refusal("global_load_lds"), + _header_refusal("global_load_lds", "src_ptr"), ), "global_load_async_to_lds": ( I32, @@ -269,7 +269,7 @@ def _lds_address(b: IRBuilder, dtype: Type) -> Value: ), "gfx1250", "llvm23", - _header_refusal("global_load_async_to_lds"), + _header_refusal("global_load_async_to_lds", "src_ptr"), ), "global_store_async_from_lds": ( I16, From 6efc6001215ed9c90cf671f0bfb7e0649889ba82 Mon Sep 17 00:00:00 2001 From: Chris Millette Date: Wed, 4 Nov 2026 11:54:59 -0500 Subject: [PATCH 7/7] test(rocke): skip a C++ engine case before the Python engine runs The per-engine helpers now check for the rocke_engine binding as soon as they are asked for the C++ engine, so a "cpp" case skips with the binding's reason in a lane that does not build it, rather than lowering with the Python engine first. available_engines() says when to use it, and the helpers' own argument errors name the bad value and what they expected. Co-Authored-By: Claude --- .../platform/python/rocke/tests/engines.py | 32 +++++++++++-------- 1 file changed, 19 insertions(+), 13 deletions(-) diff --git a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/tests/engines.py b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/tests/engines.py index 57043ae29df9..b7ce0a86be02 100644 --- a/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/tests/engines.py +++ b/dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/tests/engines.py @@ -76,7 +76,9 @@ Engine = Literal["python", "cpp"] ENGINES: tuple[Engine, ...] = ("python", "cpp") -# The rocke_status_t (rocke/ir.h) each Python exception type maps to. +# The rocke_status_t (rocke/ir.h is the source of truth) each Python exception +# type maps to. Matched by exact type: ArchDomainError (status 6) is left out, as +# no refusal checked here raises it. _NATIVE_STATUS: dict[type[Exception], int] = { ValueError: 1, TypeError: 2, @@ -87,7 +89,11 @@ def available_engines() -> tuple[Engine, ...]: """:data:`ENGINES` when the ``rocke_engine`` binding is importable, else - ``("python",)``.""" + ``("python",)``. + + For a test that loops over the engines itself; a ``pytest`` test + parametrizes over :data:`ENGINES` instead. + """ if importlib.util.find_spec("rocke_engine") is None: return ("python",) return ENGINES @@ -104,17 +110,19 @@ def _require_native_engine() -> ModuleType: return pytest.importorskip("rocke_engine") -def _require_engine_name(engine: Engine) -> None: +def _require_engine(engine: Engine) -> None: + """Validate ``engine``; for ``"cpp"``, skip the calling test when the + binding is not built, before any engine runs.""" if engine not in ENGINES: - raise ValueError(f"engine must be one of {ENGINES}, got {engine!r}") + raise ValueError(f"engine={engine!r}, expected one of {ENGINES}") + if engine == "cpp": + _require_native_engine() def _native_status(error_type: type[Exception]) -> int: if error_type not in _NATIVE_STATUS: names = sorted(t.__name__ for t in _NATIVE_STATUS) - raise ValueError( - f"error_type must be one of {names}, got {error_type.__name__}" - ) + raise ValueError(f"error_type={error_type.__name__}, expected one of {names}") return _NATIVE_STATUS[error_type] @@ -155,7 +163,7 @@ def lower_with_engine( kernel: KernelDef, arch: str, *, engine: Engine, llvm_flavor: str | None = None ) -> str: """Lower ``kernel`` for ``arch`` with ``engine``; return its LLVM IR text.""" - _require_engine_name(engine) + _require_engine(engine) python_ir = _python_llvm(kernel, arch, llvm_flavor) if engine == "python": return python_ir @@ -182,7 +190,7 @@ def assert_engine_rejects( C++ engine must fail with ``error_type``'s status and the same message after its binding prefix. """ - _require_engine_name(engine) + _require_engine(engine) if engine == "python": _assert_python_rejects( lambda: _python_llvm(kernel, arch, llvm_flavor), error_type, message @@ -201,7 +209,7 @@ def lower_hip_with_engine(kernel: KernelDef, arch: str, *, engine: Engine) -> st For ``"cpp"`` it also asserts that the text is identical to the Python engine's. """ - _require_engine_name(engine) + _require_engine(engine) python_hip = lower_kernel_to_hip(kernel, arch=arch) if engine == "python": return python_hip @@ -225,7 +233,7 @@ def assert_engine_rejects_hip( Same contract as :func:`assert_engine_rejects`, for the HIP source. """ - _require_engine_name(engine) + _require_engine(engine) if engine == "python": _assert_python_rejects( lambda: lower_kernel_to_hip(kernel, arch=arch), error_type, message @@ -240,7 +248,6 @@ def lower_with_both_engines( kernel: KernelDef, arch: str, llvm_flavor: str | None = None ) -> str: """Lower ``kernel`` for ``arch`` with each engine; return the shared IR text.""" - _require_native_engine() return lower_with_engine(kernel, arch, engine="cpp", llvm_flavor=llvm_flavor) @@ -267,7 +274,6 @@ def assert_both_engines_reject( def lower_hip_with_both_engines(kernel: KernelDef, arch: str) -> str: """Lower ``kernel`` to HIP source for ``arch`` with each engine; return the shared text.""" - _require_native_engine() return lower_hip_with_engine(kernel, arch, engine="cpp")