low.kernel.def
← low dialect
Target-bound low kernel entry with register-typed launch ABI values. Helper calls stay in low.func.def; kernel launch/export contracts live on this entry op.
Operation contract
| Property |
Value |
| Semantic phase |
executable |
| Traits |
SymbolDefine, IsolatedFromAbove |
| Interfaces |
FuncLike |
Symbol contract
| Property |
Value |
| Kind |
function |
| Defining field |
callee |
| Interfaces |
func_like, kernel_entry |
| Flags |
— |
Signature
| Kind |
Name |
Type |
Cardinality |
Description |
| Attribute |
callee |
symbol |
required |
— |
| Attribute |
target |
symbol (target) |
optional |
— |
| Attribute |
descriptor_set |
string |
required |
Canonical descriptor-set key governing the low representation. |
| Attribute |
abi_layout |
dict |
optional |
— |
| Attribute |
export_symbol |
string |
optional |
— |
| Attribute |
export_linkage |
enum ExportLinkage |
optional |
— |
| Attribute |
workgroup_size_x |
i64 |
optional |
— |
| Attribute |
workgroup_size_y |
i64 |
optional |
— |
| Attribute |
workgroup_size_z |
i64 |
optional |
— |
| Attribute |
workgroup_count_x |
i64 |
optional |
— |
| Attribute |
workgroup_count_y |
i64 |
optional |
— |
| Attribute |
workgroup_count_z |
i64 |
optional |
— |
| Attribute |
workgroup_cluster_size_x |
i64 |
optional |
— |
| Attribute |
workgroup_cluster_size_y |
i64 |
optional |
— |
| Attribute |
workgroup_cluster_size_z |
i64 |
optional |
— |
| Attribute |
allocation |
enum LowAllocationMode |
optional |
— |
| Attribute |
schedule |
enum LowScheduleMode |
optional |
— |
| Attribute |
predicates |
predicate_list |
optional |
— |
| Attribute |
retain |
enum Retain |
optional |
— |
| Region |
body |
region |
required |
Low kernel body. (terminator low.return.) |
Verification constraints
Examples
low.kernel.def target<amdgpu.gfx11.generic.core>(@gfx11_generic) export("matmul") workgroup_size(16, 4, 1) @matmul(%lhs: reg<amdgpu.sgpr x4>, %rhs: reg<amdgpu.sgpr x4>, %out: reg<amdgpu.sgpr x4>) {
low.return
}