kernel.dispatch
← kernel dialect
Dispatch an executable kernel entry with exact workgroup counts and device-ABI operands. Counts are one to three index values or one dense view<3xi32> for indirect dispatch. An optional XYZ workgroup size overrides the executable entry default for targets supporting dispatch-time workgroup sizes.
Operation contract
| Property |
Value |
| Semantic phase |
executable |
| Traits |
UnknownEffects, CommandEffect, NoAncestor(kernel.def) |
| Interfaces |
CallLike |
Signature
| Kind |
Name |
Type |
Cardinality |
Description |
| Operand |
workgroup_counts |
any |
variadic |
— |
| Operand |
arguments |
any |
variadic |
— |
| Operand |
workgroup_size |
index |
variadic |
— |
| Attribute |
callee |
symbol (kernel_entry) |
required |
— |
Examples
kernel.dispatch @fill[%count](%count, %output) : [index](index, buffer)
kernel.dispatch @fill[%x, %y, %z](%count, %output) : [index, index, index](index, buffer)
kernel.dispatch @fill[%counts](%count, %output) : [view<3xi32>](index, buffer)
kernel.dispatch @fill[%count](%count, %output) workgroup_size(%x, %y, %z) : [index](index, buffer)