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
Flags	—

Signature

Kind	Name	Type	Cardinality	Description
Attribute	callee	symbol	required	—
Attribute	target	symbol (target)	optional	—
Attribute	export_symbol	string	optional	—
Attribute	export_linkage	enum ExportLinkage	optional	—
Attribute	predicates	predicate_list	optional	—
Attribute	retain	enum Retain	optional	—
Region	config	region	required	Launch configuration region. The region owns host launch inputs and must terminate with kernel.launch.config. (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
}