kernel.def

← kernel dialect

Dispatchable source-level kernel entry. Kernel entries own launch and export contracts; ordinary func.def bodies remain helper/callable code.

Operation contract

Property	Value
Semantic phase	executable
Traits	SymbolDefine, IsolatedFromAbove
Interfaces	FuncLike

Symbol contract

Property	Value
Kind	kernel
Defining field	callee
Interfaces	func_like, kernel, kernel_entry
Flags	—

Signature

Kind	Name	Type	Cardinality	Description
Attribute	callee	symbol	required	—
Attribute	linkage	enum Linkage	optional	—
Attribute	target	symbol (target)	optional	—
Attribute	export_symbol	string	optional	Artifact export name override; omission exports the kernel's symbol name.
Attribute	export_linkage	enum ExportLinkage	optional	—
Attribute	predicates	predicate_list	optional	—
Attribute	retain	enum Retain	optional	—
Region	config	region	required	Deterministic launch configuration region. The region computes launch parameters from its workload arguments and immutable compilation inputs and must terminate with kernel.launch.config. Workgroup counts may use deterministic loads from explicit buffer or view workload arguments. Workgroup and cluster sizes remain independent of loaded workload data. (single block, terminator kernel.launch.config.)
Region	body	region	required	Kernel body. (terminator kernel.return.)

Examples

kernel.def @entry() {
  %one = index.constant 1 : index
  kernel.launch.config workgroups(%one, %one, %one) workgroup_size(%one, %one, %one) : index
} launch(%buffer: buffer) {
  kernel.return
}

kernel.def target(@gfx11_generic) export("matmul") @matmul(%m: index, %n: index) {
  %one = index.constant 1 : index
  %threads = index.constant 256 : index
  kernel.launch.config workgroups(%m, %n, %one) workgroup_size(%threads, %one, %one) : index
} launch(%lhs: buffer, %rhs: buffer, %out: buffer) {
  kernel.return
}