Skip to content

refactor(rocke): read gfx1250-only checks from target facts in both engines #14040

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
cgmillette merged 3 commits into from
Oct 21, 2026
Merged

refactor(rocke): read gfx1250-only checks from target facts in both engines #14040

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
ec1e78b
refactor(rocke): admit LLVM-23-only memory and control ops by target …
cgmillette Oct 20, 2026
b92495a
refactor(rocke): resolve scaled WMMA atoms from the MMA catalogs, not…
cgmillette Oct 20, 2026
8b75396
refactor(rocke): admit scaled WMMA in HIP lowering by the target catalog
cgmillette Oct 20, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
15 changes: 15 additions & 0 deletions dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/arch/query.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -446,6 +446,21 @@ const char* rocke_arch_mma_op_id_family(const char* op_id)
return index.lookup(op_id);
}

const rocke_mma_op_t* rocke_arch_mma_op_by_id(const char* op_id)
{
if(!op_id)
return NULL;
/* Mirrors target._mma_op_by_id: the registry is in known-arch order. */
for(int i = 0; i < rocke_ati_arch_registry_len; ++i)
{
const rocke_arch_target_t* t = rocke_ati_arch_registry[i].target;
const rocke_mma_op_t* op = t ? rocke_mma_catalog_by_op_id(&t->mma, op_id) : NULL;
if(op)
return op;
}
return NULL;
}

const char* rocke_arch_mma_op_id_c_dtype(const char* op_id)
{
int i;
Expand Down
23 changes: 20 additions & 3 deletions dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mma.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -28,6 +28,8 @@
* getattr-dispatch raises immediately and we reproduce that same rejection.
* Every opcode is registered in the dispatch table; none is dropped.
*/
#include <ctype.h>

#include "rocke/arch_target.h"
#include "rocke/ir.h"
#include "rocke/ir_internal.h"
Expand Down Expand Up @@ -119,13 +121,28 @@ static rocke_status_t h_emit_gfx1250_scaled_wmma(rocke_h_lowerer_t* lw, const ro
const bool scale16 = spec->scales.block_k == 16;
const char* builtin = scale16 ? "__builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4"
: "__builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4";
if(!lw->arch.gfx || __builtin_strcmp(lw->arch.gfx, "gfx1250") != 0)
/* Admitted on any target whose MMA catalog lists the atom, never by the
* target's name; the rule and message of Python _require_wmma_arch. */
const rocke_arch_target_t* target
= lw->arch.gfx ? rocke_arch_target_from_gfx(lw->arch.gfx) : NULL;
if(!target || !rocke_mma_catalog_by_op_id(&target->mma, op_id))
{
char family[16] = "(UNKNOWN)";
if(target)
{
size_t i = 0;
for(; target->family[i] && i + 1 < sizeof(family); ++i)
family[i] = (char)toupper((unsigned char)target->family[i]);
family[i] = '\0';
}
return rocke_h_fail(lw,
ROCKE_ERR_NOTIMPL,
"WMMA op '%s' is not available on %s",
"WMMA op '%s' is not available on %s "
"(WMMA is an RDNA/gfx11 instruction; this is a %s target). "
"The MFMA atoms are the matrix path on CDNA.",
op_id,
lw->arch.gfx ? lw->arch.gfx : "(unknown)");
lw->arch.gfx ? lw->arch.gfx : "(unknown)",
family);
}
if(op->num_operands != 5 || op->num_results != 1)
{
Expand Down
47 changes: 14 additions & 33 deletions dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/control.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1086,26 +1086,6 @@ static void ll_check_u16(rocke_lower_t* L, const char* op, const char* field, in
}
}

static void ll_require_gfx1250_llvm23(rocke_lower_t* L, const char* op)
{
if(!L->backend || !L->backend->gfx || strcmp(L->backend->gfx, "gfx1250") != 0)
{
rocke_ll_fail(L,
ROCKE_ERR_VALUE,
"%s requires gfx1250, got %s",
op,
(L->backend && L->backend->gfx) ? L->backend->gfx : "(unknown)");
}
if(L->flavor != ROCKE_LLVM_FLAVOR_LLVM23)
{
rocke_ll_fail(L,
ROCKE_ERR_VALUE,
"%s requires LLVM flavor llvm23, got %s",
op,
rocke_llvm_flavor_name(L->flavor));
}
}

static const char*
ll_named_barrier_ptr(rocke_lower_t* L, const char* op, const rocke_value_t* value)
{
Expand Down Expand Up @@ -1143,7 +1123,7 @@ static void _op_tile_s_wait_asynccnt(rocke_lower_t* L, const rocke_op_t* op)
static void _op_tile_s_wait_tensorcnt(rocke_lower_t* L, const rocke_op_t* op)
{
int64_t n = 0;
ll_require_gfx1250_llvm23(L, "s_wait_tensorcnt");
rocke_ll_require_llvm23_target_fact(L, "s_wait_tensorcnt", ROCKE_LL_FACT_TDM);
rocke_attr_get_int(&op->attrs, "n", &n);
ll_check_u16(L, "s_wait_tensorcnt", "n", n);
rocke_ll_need(L, "s.wait.tensorcnt");
Expand All @@ -1153,7 +1133,7 @@ static void _op_tile_s_wait_tensorcnt(rocke_lower_t* L, const rocke_op_t* op)
static void _op_tile_s_barrier_signal(rocke_lower_t* L, const rocke_op_t* op)
{
int64_t barrier_type = 0;
ll_require_gfx1250_llvm23(L, "s_barrier_signal");
rocke_ll_require_llvm23_target_fact(L, "s_barrier_signal", ROCKE_LL_FACT_GFX12_CDNA_FAMILY);
rocke_attr_get_int(&op->attrs, "barrier_type", &barrier_type);
if(barrier_type < 0 || (uint64_t)barrier_type > UINT32_MAX)
rocke_ll_fail(L, ROCKE_ERR_VALUE, "s_barrier_signal barrier_type must fit unsigned i32");
Expand All @@ -1166,7 +1146,7 @@ static void _op_tile_s_barrier_signal(rocke_lower_t* L, const rocke_op_t* op)
static void _op_tile_s_barrier_wait(rocke_lower_t* L, const rocke_op_t* op)
{
int64_t barrier_type = 0;
ll_require_gfx1250_llvm23(L, "s_barrier_wait");
rocke_ll_require_llvm23_target_fact(L, "s_barrier_wait", ROCKE_LL_FACT_GFX12_CDNA_FAMILY);
rocke_attr_get_int(&op->attrs, "barrier_type", &barrier_type);
ll_check_u16(L, "s_barrier_wait", "barrier_type", barrier_type);
rocke_ll_need(L, "s.barrier.wait");
Expand All @@ -1185,7 +1165,7 @@ static void ll_named_barrier_count(rocke_lower_t* L,
if(!count->type || count->type->kind != ROCKE_TYPE_SCALAR
|| count->type->scalar != ROCKE_SCALAR_I32)
rocke_ll_fail(L, ROCKE_ERR_VALUE, "%s member_count must be i32", short_name);
ll_require_gfx1250_llvm23(L, short_name);
rocke_ll_require_llvm23_target_fact(L, short_name, ROCKE_LL_FACT_GFX12_CDNA_FAMILY);
const char* ptr = ll_named_barrier_ptr(L, short_name, barrier);
rocke_ll_need(L, intrinsic);
rocke_ll_emitf(L,
Expand All @@ -1212,7 +1192,7 @@ static void ll_named_barrier_one(rocke_lower_t* L,
{
if(op->num_operands != 1)
rocke_ll_fail(L, ROCKE_ERR_VALUE, "%s expects one barrier pointer", short_name);
ll_require_gfx1250_llvm23(L, short_name);
rocke_ll_require_llvm23_target_fact(L, short_name, ROCKE_LL_FACT_GFX12_CDNA_FAMILY);
const char* ptr = ll_named_barrier_ptr(L, short_name, op->operands[0]);
rocke_ll_need(L, intrinsic);
rocke_ll_emitf(L, " call void @llvm.amdgcn.%s(ptr addrspace(3) %s)", intrinsic, ptr);
Expand All @@ -1231,7 +1211,7 @@ static void _op_tile_s_wakeup_barrier(rocke_lower_t* L, const rocke_op_t* op)
static void _op_tile_s_barrier_leave(rocke_lower_t* L, const rocke_op_t* op)
{
int64_t barrier_type = 0;
ll_require_gfx1250_llvm23(L, "s_barrier_leave");
rocke_ll_require_llvm23_target_fact(L, "s_barrier_leave", ROCKE_LL_FACT_GFX12_CDNA_FAMILY);
rocke_attr_get_int(&op->attrs, "barrier_type", &barrier_type);
ll_check_u16(L, "s_barrier_leave", "barrier_type", barrier_type);
rocke_ll_need(L, "s.barrier.leave");
Expand Down Expand Up @@ -1316,7 +1296,7 @@ static void _op_tile_s_setreg(rocke_lower_t* L, const rocke_op_t* op)
int64_t simm16 = 0;
if(!rocke_ll_live(L))
return;
ll_require_gfx1250_llvm23(L, "s_setreg");
rocke_ll_require_llvm23_target_fact(L, "s_setreg", ROCKE_LL_FACT_GFX12_CDNA_FAMILY);
ll_expect_operands(L, op, "s_setreg", 1);
rocke_attr_get_int(&op->attrs, "simm16", &simm16);
ll_check_u16(L, "s_setreg", "simm16", simm16);
Expand All @@ -1333,7 +1313,7 @@ static void _op_tile_s_prefetch_data(rocke_lower_t* L, const rocke_op_t* op)
int space;
if(!rocke_ll_live(L))
return;
ll_require_gfx1250_llvm23(L, "s_prefetch_data");
rocke_ll_require_llvm23_target_fact(L, "s_prefetch_data", ROCKE_LL_FACT_SCALAR_DATA_PREFETCH);
ll_expect_operands(L, op, "s_prefetch_data", 2);
space = rocke_ll_anyptr_space(L,
"s_prefetch_data",
Expand All @@ -1355,7 +1335,8 @@ static void _op_tile_s_buffer_prefetch_data(rocke_lower_t* L, const rocke_op_t*
int64_t offset = 0;
if(!rocke_ll_live(L))
return;
ll_require_gfx1250_llvm23(L, "s_buffer_prefetch_data");
rocke_ll_require_llvm23_target_fact(
L, "s_buffer_prefetch_data", ROCKE_LL_FACT_SCALAR_DATA_PREFETCH);
ll_expect_operands(L, op, "s_buffer_prefetch_data", 2);
rocke_attr_get_int(&op->attrs, "offset", &offset);
if(offset < INT32_MIN || offset > INT32_MAX)
Expand All @@ -1378,7 +1359,7 @@ static void _op_tile_global_prefetch(rocke_lower_t* L, const rocke_op_t* op)
int64_t cp;
if(!rocke_ll_live(L))
return;
ll_require_gfx1250_llvm23(L, "global_prefetch");
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)
Expand All @@ -1398,7 +1379,7 @@ static void _op_tile_flat_prefetch(rocke_lower_t* L, const rocke_op_t* op)
int64_t cp;
if(!rocke_ll_live(L))
return;
ll_require_gfx1250_llvm23(L, "flat_prefetch");
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)
Expand All @@ -1425,7 +1406,7 @@ static void ll_lower_cluster_read(rocke_lower_t* L,
char key[64];
if(!rocke_ll_live(L) || !res)
return;
ll_require_gfx1250_llvm23(L, short_name);
rocke_ll_require_llvm23_target_fact(L, short_name, ROCKE_LL_FACT_CLUSTER_LAUNCH);
if(op->num_operands != 0)
rocke_ll_fail(L, ROCKE_ERR_VALUE, "%s takes no operands", short_name);
if(with_axis)
Expand Down Expand Up @@ -1483,7 +1464,7 @@ static void _op_tile_cluster_barrier(rocke_lower_t* L, const rocke_op_t* op)
(void)op;
if(!rocke_ll_live(L))
return;
ll_require_gfx1250_llvm23(L, "cluster_barrier");
rocke_ll_require_llvm23_target_fact(L, "cluster_barrier", ROCKE_LL_FACT_CLUSTER_LAUNCH);
rocke_ll_need(L, "s.cluster.barrier");
rocke_ll_emit(L, " fence syncscope(\"cluster\") release");
rocke_ll_emit(L, " call void @llvm.amdgcn.s.cluster.barrier()");
Expand Down
82 changes: 67 additions & 15 deletions dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -323,6 +323,71 @@ bool rocke_ll_live(const rocke_lower_t* L)
return L != NULL;
}

static const char* ll_target_fact_name(rocke_ll_target_fact_t fact)
{
switch(fact)
{
case ROCKE_LL_FACT_TDM:
return "memory.has_tdm";
case ROCKE_LL_FACT_CLUSTER_LAUNCH:
return "memory.has_cluster_launch";
case ROCKE_LL_FACT_MULTICAST_LOAD:
return "memory.has_multicast_load";
case ROCKE_LL_FACT_SCALAR_DATA_PREFETCH:
return "memory.has_scalar_data_prefetch";
case ROCKE_LL_FACT_GLOBAL_PREFETCH:
return "memory.has_global_prefetch";
case ROCKE_LL_FACT_GFX12_CDNA_FAMILY:
return "target_family gfx12_cdna";
}
return "(unknown fact)";
}

static bool ll_target_fact_holds(const rocke_arch_target_t* t, rocke_ll_target_fact_t fact)
{
switch(fact)
{
case ROCKE_LL_FACT_TDM:
return t->memory.has_tdm;
case ROCKE_LL_FACT_CLUSTER_LAUNCH:
return t->memory.has_cluster_launch;
case ROCKE_LL_FACT_MULTICAST_LOAD:
return t->memory.has_multicast_load;
case ROCKE_LL_FACT_SCALAR_DATA_PREFETCH:
return t->memory.has_scalar_data_prefetch;
case ROCKE_LL_FACT_GLOBAL_PREFETCH:
return t->memory.has_global_prefetch;
case ROCKE_LL_FACT_GFX12_CDNA_FAMILY:
return strcmp(t->target_family, "gfx12_cdna") == 0;
}
return false;
}

void rocke_ll_require_llvm23_target_fact(rocke_lower_t* L,
const char* op,
rocke_ll_target_fact_t fact)
{
const char* gfx = (L->backend && L->backend->gfx) ? L->backend->gfx : NULL;
const rocke_arch_target_t* target = gfx ? rocke_arch_target_from_gfx(gfx) : NULL;
if(!target || !ll_target_fact_holds(target, fact))
{
rocke_ll_fail(L,
ROCKE_ERR_VALUE,
"%s requires a target with %s, got %s",
op,
ll_target_fact_name(fact),
gfx ? gfx : "(unknown)");
}
if(L->flavor != ROCKE_LLVM_FLAVOR_LLVM23)
{
rocke_ll_fail(L,
ROCKE_ERR_VALUE,
"%s requires LLVM flavor llvm23, got %s",
op,
rocke_llvm_flavor_name(L->flavor));
}
}

/* ====================================================================== */
/* Block / CFG model (Python _Block) */
/* ====================================================================== */
Expand Down Expand Up @@ -1799,27 +1864,14 @@ const char* rocke_ll_format_agpr_alloc(rocke_lower_t* L, const rocke_attr_value_
}

/* gfx1250 workgroup-cluster shape; mirrors Python ir.check_cluster_dims (same
* limits and messages) behind the gfx1250 + llvm23 gate. The backend copies
* limits and messages) behind the cluster-launch + llvm23 gate. The backend copies
* the attribute to the code-object .cluster_dims metadata unchecked. */
#define ROCKE_LL_MAX_CLUSTER_DIM 15
#define ROCKE_LL_MAX_CLUSTER_WORKGROUPS 16

static const char* rocke_ll_format_cluster_dims(rocke_lower_t* L, const rocke_attr_value_t* v)
{
if(!L->backend || !L->backend->gfx || strcmp(L->backend->gfx, "gfx1250") != 0)
{
rocke_ll_fail(L,
ROCKE_ERR_VALUE,
"cluster_dims requires gfx1250, got %s",
(L->backend && L->backend->gfx) ? L->backend->gfx : "(unknown)");
}
if(L->flavor != ROCKE_LLVM_FLAVOR_LLVM23)
{
rocke_ll_fail(L,
ROCKE_ERR_VALUE,
"cluster_dims requires LLVM flavor llvm23, got %s",
rocke_llvm_flavor_name(L->flavor));
}
rocke_ll_require_llvm23_target_fact(L, "cluster_dims", ROCKE_LL_FACT_CLUSTER_LAUNCH);
if(!v || v->kind != ROCKE_ATTR_INT_LIST || v->u.ilist.count != 3)
{
rocke_ll_fail(L, ROCKE_ERR_VALUE, "cluster_dims must be three integers (x, y, z)");
Expand Down
33 changes: 8 additions & 25 deletions dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mem.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1450,26 +1450,6 @@ static const char* ll_lds_ptr_operand(rocke_lower_t* L, const char* op_name, con
return name;
}

static void ll_require_gfx1250_llvm23(rocke_lower_t* L, const char* op)
{
if(!L->backend || !L->backend->gfx || strcmp(L->backend->gfx, "gfx1250") != 0)
{
rocke_ll_fail(L,
ROCKE_ERR_VALUE,
"%s requires gfx1250, got %s",
op,
(L->backend && L->backend->gfx) ? L->backend->gfx : "(unknown)");
}
if(L->flavor != ROCKE_LLVM_FLAVOR_LLVM23)
{
rocke_ll_fail(L,
ROCKE_ERR_VALUE,
"%s requires LLVM flavor llvm23, got %s",
op,
rocke_llvm_flavor_name(L->flavor));
}
}

static void op_tile_async_buffer_load_lds_addr(rocke_lower_t* L, const rocke_op_t* op)
{
const rocke_value_t* rsrc = op->operands[0];
Expand Down Expand Up @@ -1643,7 +1623,8 @@ static void op_tile_global_store_async_from_lds(rocke_lower_t* L, const rocke_op
int64_t cachepolicy;
const char* suffix;
const char* key;
ll_require_gfx1250_llvm23(L, "global_store_async_from_lds");
rocke_ll_require_llvm23_target_fact(
L, "global_store_async_from_lds", ROCKE_LL_FACT_GFX12_CDNA_FAMILY);
if(op->num_operands != 2)
rocke_ll_fail(L, ROCKE_ERR_VALUE, "global_store_async_from_lds expects two operands");
dst = op->operands[0];
Expand Down Expand Up @@ -1692,7 +1673,8 @@ static void op_tile_global_load_tr16_b128(rocke_lower_t* L, const rocke_op_t* op
const char* suffix;
const char* key;
const char* llvm_type;
ll_require_gfx1250_llvm23(L, "global_load_tr16_b128");
rocke_ll_require_llvm23_target_fact(
L, "global_load_tr16_b128", ROCKE_LL_FACT_GFX12_CDNA_FAMILY);
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];
Expand Down Expand Up @@ -1745,7 +1727,7 @@ static void op_tile_cluster_load(rocke_lower_t* L, const rocke_op_t* op)
int lanes;
if(!rocke_ll_live(L))
return;
ll_require_gfx1250_llvm23(L, "cluster_load");
rocke_ll_require_llvm23_target_fact(L, "cluster_load", ROCKE_LL_FACT_MULTICAST_LOAD);
if(op->num_operands != 2 || op->num_results != 1)
rocke_ll_fail(L, ROCKE_ERR_VALUE, "cluster_load expects two operands and one result");
ptr = op->operands[0];
Expand Down Expand Up @@ -1825,7 +1807,8 @@ static void op_tile_cluster_load_async_to_lds(rocke_lower_t* L, const rocke_op_t
int64_t cachepolicy;
if(!rocke_ll_live(L))
return;
ll_require_gfx1250_llvm23(L, "cluster_load_async_to_lds");
rocke_ll_require_llvm23_target_fact(
L, "cluster_load_async_to_lds", ROCKE_LL_FACT_MULTICAST_LOAD);
if(op->num_operands != 3)
rocke_ll_fail(L, ROCKE_ERR_VALUE, "cluster_load_async_to_lds expects three operands");
src = op->operands[0];
Expand Down Expand Up @@ -1884,7 +1867,7 @@ static void op_tile_tensor_lds_transfer(rocke_lower_t* L,
static const int lanes[] = {4, 8, 4, 4, 8};
int64_t cachepolicy;
int i;
ll_require_gfx1250_llvm23(L, short_name);
rocke_ll_require_llvm23_target_fact(L, short_name, ROCKE_LL_FACT_TDM);
if(op->num_operands != 5)
rocke_ll_fail(L, ROCKE_ERR_VALUE, "%s expects five descriptor groups", short_name);
for(i = 0; i < 5; ++i)
Expand Down
Loading
Loading