'gpu' Dialect

Note: this dialect is more likely to change than others in the near future; use with caution.

This dialect provides middle-level abstractions for launching GPU kernels following a programming model similar to that of CUDA or OpenCL. It provides abstractions for kernel invocations (and may eventually provide those for device management) that are not present at the lower level (e.g., as LLVM IR intrinsics for GPUs). Its goal is to abstract away device- and driver-specific manipulations to launch a GPU kernel and provide a simple path towards GPU execution from MLIR. It may be targeted, for example, by DSLs using MLIR. The dialect uses gpu as its canonical prefix.

Memory attribution ¶

Memory buffers are defined at the function level, either in “gpu.launch” or in “gpu.func” ops. This encoding makes it clear where the memory belongs and makes the lifetime of the memory visible. The memory is only accessible while the kernel is launched/the function is currently invoked. The latter is more strict than actual GPU implementations but using static memory at the function level is just for convenience. It is also always possible to pass pointers to the workgroup memory into other functions, provided they expect the correct memory space.

The buffers are considered live throughout the execution of the GPU function body. The absence of memory attribution syntax means that the function does not require special buffers. Rationale: although the underlying models declare memory buffers at the module level, we chose to do it at the function level to provide some structuring for the lifetime of those buffers; this avoids the incentive to use the buffers for communicating between different kernels or launches of the same kernel, which should be done through function arguments instead; we chose not to use alloca-style approach that would require more complex lifetime analysis following the principles of MLIR that promote structure and representing analysis results in the IR.

Operations ¶

gpu.all_reduce (::mlir::gpu::AllReduceOp) ¶

Reduce values among workgroup.

The all_reduce op reduces the value of every work item across a local workgroup. The result is equal for all work items of a workgroup.

For example, both

%1 = "gpu.all_reduce"(%0) ({}) { op = "add" } : (f32) -> (f32)
%2 = "gpu.all_reduce"(%0) ({
^bb(%lhs : f32, %rhs : f32):
  %sum = addf %lhs, %rhs : f32
  "gpu.yield"(%sum) : (f32) -> ()
}) : (f32) -> (f32)

compute the sum of each work item’s %0 value. The first version specifies the accumulation as operation, whereas the second version specifies the accumulation as code region. The accumulation operation must be one of: add, and, max, min, mul, or, xor.

Either none or all work items of a workgroup need to execute this op in convergence.

Attributes: ¶

Operands: ¶

Results: ¶

gpu.alloc (::mlir::gpu::AllocOp) ¶

GPU memory allocation operation.

Syntax:

operation ::= `gpu.alloc` custom<AsyncDependencies>(type($asyncToken), $asyncDependencies) ` `
              `(` $dynamicSizes `)` (`` `[` $symbolOperands^ `]`)? attr-dict `:` type($memref)

The gpu.alloc operation allocates a region of memory on the GPU. It is similar to the memref.alloc op, but supports asynchronous GPU execution.

The op does not execute before all async dependencies have finished executing.

If the async keyword is present, the op is executed asynchronously (i.e. it does not block until the execution has finished on the device). In that case, it also returns a !gpu.async.token.

Example:

%memref, %token = gpu.alloc async [%dep] (%width) : memref<64x?xf32, 1>

Operands: ¶

Results: ¶

gpu.barrier (::mlir::gpu::BarrierOp) ¶

Synchronizes all work items of a workgroup.

The “barrier” op synchronizes all work items of a workgroup. It is used to coordinate communication between the work items of the workgroup.

gpu.barrier

waits until all work items in the workgroup have reached this point and all memory accesses made by these work items prior to the op are visible to all work items in the workgroup. Data hazards between work items accessing the same memory can be avoided by synchronizing work items in-between these accesses.

Either none or all work items of a workgroup need to execute this op in convergence.

gpu.block_dim (::mlir::gpu::BlockDimOp) ¶

Returns the number of threads in the thread block (aka the block size) along the x, y, or z dimension.

Example:

%bDimX = "gpu.block_dim"() {dimension = "x"} : () -> (index)

Attributes: ¶

Results: ¶

gpu.block_id (::mlir::gpu::BlockIdOp) ¶

Returns the block id, i.e. the index of the current block within the grid along the x, y, or z dimension.

Example:

%bIdY = "gpu.block_id"() {dimension = "y"} : () -> (index)

Attributes: ¶

Results: ¶

gpu.dealloc (::mlir::gpu::DeallocOp) ¶

GPU memory deallocation operation

Syntax:

operation ::= `gpu.dealloc` custom<AsyncDependencies>(type($asyncToken), $asyncDependencies)
              $memref attr-dict `:` type($memref)

The gpu.dealloc operation frees the region of memory referenced by a memref which was originally created by the gpu.alloc operation. It is similar to the memref.dealloc op, but supports asynchronous GPU execution.

The op does not execute before all async dependencies have finished executing.

If the async keyword is present, the op is executed asynchronously (i.e. it does not block until the execution has finished on the device). In that case, it returns a !gpu.async.token.

Example:

%token = gpu.dealloc async [%dep] %memref : memref<8x64xf32, 1>

Operands: ¶

Results: ¶

gpu.func (::mlir::gpu::GPUFuncOp) ¶

Function executable on a GPU

Defines a function that can be executed on a GPU. This supports memory attribution and its body has a particular execution model.

GPU functions are either kernels (as indicated by the kernel attribute) or regular functions. The former can be launched from the host side, while the latter are device side only.

The memory attribution defines SSA values that correspond to memory buffers allocated in the memory hierarchy of the GPU (see below).

The operation has one attached region that corresponds to the body of the function. The region arguments consist of the function arguments without modification, followed by buffers defined in memory annotations. The body of a GPU function, when launched, is executed by multiple work items. There are no guarantees on the order in which work items execute, or on the connection between them. In particular, work items are not necessarily executed in lock-step. Synchronization ops such as “gpu.barrier” should be used to coordinate work items. Declarations of GPU functions, i.e. not having the body region, are not supported.

Syntax:

op ::= `gpu.func` symbol-ref-id `(` argument-list `)` (`->`
function-result-list)?
       memory-attribution `kernel`? function-attributes? region

memory-attribution ::= (`workgroup` `(` ssa-id-and-type-list `)`)?
                       (`private` `(` ssa-id-and-type-list `)`)?

Example:

gpu.func @foo(%arg0: index)
    workgroup(%workgroup: memref<32xf32, 3>)
    private(%private: memref<1xf32, 5>)
    kernel
    attributes {qux: "quux"} {
  gpu.return
}

The generic form illustrates the concept

"gpu.func"(%arg: index) {sym_name: "foo", kernel, qux: "quux"} ({
^bb0(%arg0: index, %workgroup: memref<32xf32, 3>,
     %private: memref<1xf32, 5>):
  "gpu.return"() : () -> ()
}) : (index) -> ()

Note the non-default memory spaces used in memref types in memory attribution.

gpu.module (::mlir::gpu::GPUModuleOp) ¶

A top level compilation unit containing code to be run on a GPU.

GPU module contains code that is intended to be run on a GPU. A host device can launch this code through a gpu.launc_func that creates a fully qualified symbol through the gpu.module’s symbol and a gpu.func symbol contained in the gpu.module.

The module’s top-level scope is modeled by a single region with a single block. GPU modules are required to have a name that is used for symbol resolution by the gpu.launch_func operation.

Using an op with a region to define a GPU module enables “embedding” GPU modules with SIMT execution models in other dialects in a clean manner and allows filtering of code regions to execute passes on only code intended to or not intended to be run on the separate device.

  gpu.module @symbol_name {
  gpu.func {}
    ...
  gpu.module_end
}

gpu.grid_dim (::mlir::gpu::GridDimOp) ¶

Returns the number of thread blocks in the grid along the x, y, or z dimension.

Example:

%gDimZ = "gpu.grid_dim"() {dimension = "z"} : () -> (index)

Attributes: ¶

Results: ¶

gpu.host_register (::mlir::gpu::HostRegisterOp) ¶

Registers a memref for access from device.

Syntax:

operation ::= `gpu.host_register` $value attr-dict `:` type($value)

This op maps the provided host buffer into the device address space.

This operation may not be supported in every environment, there is not yet a way to check at runtime whether this feature is supported.

Writes from the host are guaranteed to be visible to device kernels that are launched afterwards. Writes from the device are guaranteed to be visible on the host after synchronizing with the device kernel completion.

Operands: ¶

gpu.launch_func (::mlir::gpu::LaunchFuncOp) ¶

Launches a function as a GPU kernel

Syntax:

operation ::= `gpu.launch_func` custom<AsyncDependencies>(type($asyncToken), $asyncDependencies)
              $kernel
              `blocks` `in` ` ` `(`$gridSizeX`,` $gridSizeY`,` $gridSizeZ`)`
              `threads` `in` ` ` `(`$blockSizeX`,` $blockSizeY`,` $blockSizeZ`)`
              custom<LaunchFuncOperands>($operands, type($operands))
              attr-dict

Launch a kernel function on the specified grid of thread blocks. gpu.launch operations are lowered to gpu.launch_func operations by outlining the kernel body into a function in a dedicated module, which reflects the separate compilation process. The kernel function is required to have the gpu.kernel attribute. The module containing the kernel function is required to be a gpu.module. And finally, the module containing the kernel module (which thus cannot be the top-level module) is required to have the gpu.container_module attribute. The gpu.launch_func operation has a symbol attribute named kernel to identify the fully specified kernel function to launch (both the gpu.module and func).

The gpu.launch_func supports async dependencies: the kernel does not start executing until the ops producing those async dependencies have completed.

By the default, the host implicitly blocks until kernel execution has completed. If the async keyword is present, the host does not block but instead a !gpu.async.token is returned. Other async GPU ops can take this token as dependency.

The operation requires at least the grid and block sizes along the x,y,z dimensions as arguments. When a lower-dimensional kernel is required, unused sizes must be explicitly set to 1.

The remaining operands are passed as arguments to the kernel function.

Example:

module attributes {gpu.container_module} {

  // This module creates a separate compilation unit for the GPU compiler.
  gpu.module @kernels {
    func @kernel_1(%arg0 : f32, %arg1 : memref<?xf32, 1>)
        attributes { nvvm.kernel = true } {

      // Operations that produce block/thread IDs and dimensions are
      // injected when outlining the `gpu.launch` body to a function called
      // by `gpu.launch_func`.
      %tIdX = "gpu.thread_id"() {dimension = "x"} : () -> (index)
      %tIdY = "gpu.thread_id"() {dimension = "y"} : () -> (index)
      %tIdZ = "gpu.thread_id"() {dimension = "z"} : () -> (index)

      %bDimX = "gpu.block_dim"() {dimension = "x"} : () -> (index)
      %bDimY = "gpu.block_dim"() {dimension = "y"} : () -> (index)
      %bDimZ = "gpu.block_dim"() {dimension = "z"} : () -> (index)

      %bIdX = "gpu.block_id"() {dimension = "x"} : () -> (index)
      %bIdY = "gpu.block_id"() {dimension = "y"} : () -> (index)
      %bIdZ = "gpu.block_id"() {dimension = "z"} : () -> (index)

      %gDimX = "gpu.grid_dim"() {dimension = "x"} : () -> (index)
      %gDimY = "gpu.grid_dim"() {dimension = "y"} : () -> (index)
      %gDimZ = "gpu.grid_dim"() {dimension = "z"} : () -> (index)

      "some_op"(%bx, %tx) : (index, index) -> ()
      %42 = load %arg1[%bx] : memref<?xf32, 1>
    }
  }

  %t0 = gpu.wait async
  gpu.launch_func
      async                           // (Optional) Don't block host, return token.
      [%t0]                           // (Optional) Execute only after %t0 has completed.
      @kernels::@kernel_1             // Kernel function.
      blocks in (%cst, %cst, %cst)    // Grid size.
      threads in (%cst, %cst, %cst)   // Block size.
      args(%arg0 : f32,               // (Optional) Kernel arguments.
           %arg1 : memref<?xf32, 1>)
}

Attributes: ¶

Operands: ¶

Results: ¶

gpu.launch (::mlir::gpu::LaunchOp) ¶

GPU kernel launch operation

Launch a kernel on the specified grid of thread blocks. The body of the kernel is defined by the single region that this operation contains. The operation takes six operands, with first three operands being grid sizes along x,y,z dimensions and the following three arguments being block sizes along x,y,z dimension. When a lower-dimensional kernel is required, unused sizes must be explicitly set to 1.

The body region has twelve arguments, grouped as follows:

• three arguments that contain block identifiers along x,y,z dimensions;
• three arguments that contain thread identifiers along x,y,z dimensions;
• operands of the gpu.launch operation as is (i.e. the operands for grid and block sizes).

Syntax:

operation ::= `gpu.launch` `block` `(` ssa-id-list `)` `in` ssa-reassignment
                         `threads` `(` ssa-id-list `)` `in` ssa-reassignment
                           region attr-dict?
ssa-reassignment ::= `(` ssa-id `=` ssa-use (`,` ssa-id `=` ssa-use)* `)`

Example:

gpu.launch blocks(%bx, %by, %bz) in (%sz_bx = %0, %sz_by = %1, %sz_bz = %2)
           threads(%tx, %ty, %tz) in (%sz_tx = %3, %sz_ty = %4, %sz_tz = %5) {
  // Block and thread identifiers, as well as block/grid sizes are
  // immediately usable inside body region.
  "some_op"(%bx, %tx) : (index, index) -> ()
  // Assuming %val1 is defined outside the gpu.launch region.
  %42 = load %val1[%bx] : memref<?xf32, 1>
}

// Generic syntax explains how the pretty syntax maps to the IR structure.
"gpu.launch"(%cst, %cst, %c1,  // Grid sizes.
             %cst, %c1, %c1)   // Block sizes.

    {/*attributes*/}
    // All sizes and identifiers have "index" size.
    : (index, index, index, index, index, index) -> () {
// The operation passes block and thread identifiers, followed by grid and
// block sizes.
^bb0(%bx : index, %by : index, %bz : index,
     %tx : index, %ty : index, %tz : index,
     %num_bx : index, %num_by : index, %num_bz : index,
     %num_tx : index, %num_ty : index, %num_tz : index)
  "some_op"(%bx, %tx) : (index, index) -> ()
  %3 = "memref.load"(%val1, %bx) : (memref<?xf32, 1>, index) -> f32
}

Rationale: using operation/block arguments gives analyses a clear way of understanding that a value has additional semantics (e.g., we will need to know what value corresponds to threadIdx.x for coalescing). We can recover these properties by analyzing the operations producing values, but it is easier just to have that information by construction.

Operands: ¶

gpu.memcpy (::mlir::gpu::MemcpyOp) ¶

GPU memcpy operation

Syntax:

operation ::= `gpu.memcpy` custom<AsyncDependencies>(type($asyncToken), $asyncDependencies)
              $dst`,` $src `:` type($dst)`,` type($src) attr-dict

The gpu.memcpy operation copies the content of one memref to another.

The op does not execute before all async dependencies have finished executing.

If the async keyword is present, the op is executed asynchronously (i.e. it does not block until the execution has finished on the device). In that case, it returns a !gpu.async.token.

Example:

%token = gpu.memcpy async [%dep] %dst, %src : memref<?xf32, 1>, memref<?xf32>

Operands: ¶

Results: ¶

gpu.module_end (::mlir::gpu::ModuleEndOp) ¶

A pseudo op that marks the end of a gpu.module.

This op terminates the only block inside the only region of a gpu.module.

gpu.num_subgroups (::mlir::gpu::NumSubgroupsOp) ¶

Syntax:

operation ::= `gpu.num_subgroups` attr-dict `:` type($result)

Returns the number of subgroups within a workgroup.

Example:

%numSg = gpu.num_subgroups : index

Results: ¶

gpu.return (::mlir::gpu::ReturnOp) ¶

Terminator for GPU functions.

A terminator operation for regions that appear in the body of gpu.func functions. The operands to the gpu.return are the result values returned by an invocation of the gpu.func.

Operands: ¶

gpu.shuffle (::mlir::gpu::ShuffleOp) ¶

Shuffles values within a subgroup.

The “shuffle” op moves values to a different invocation within the same subgroup.

Example:

%1, %2 = gpu.shuffle %0, %offset, %width xor : f32

For lane k returns the value from lane k ^ offset and true if that lane is smaller than %width. Otherwise it returns an unspecified value and false. A lane is the index of an invocation relative to its subgroup.

The width specifies the number of invocations that participate in the shuffle. The width needs to be the same for all invocations that participate in the shuffle. Exactly the first width invocations of a subgroup need to execute this op in convergence.

Attributes: ¶

Operands: ¶

Results: ¶

gpu.subgroup_id (::mlir::gpu::SubgroupIdOp) ¶

Syntax:

operation ::= `gpu.subgroup_id` attr-dict `:` type($result)

Returns the subgroup id, i.e. the index of the current subgroup within the workgroup.

Example:

%sgId = gpu.subgroup_id : index

Results: ¶

gpu.subgroup_mma_compute (::mlir::gpu::SubgroupMmaComputeOp) ¶

GPU warp synchronous matrix multiply accumulate

Syntax:

operation ::= `gpu.subgroup_mma_compute` $opA`,` $opB`,` $opC attr-dict `:` type($opA)`,` type($opB)`,` type($opC) `->` type($res)

The gpu.subgroup_mma_compute operation performs a matrix-multiply accumulate(mma) operation using all the threads in a subgroup.

This operation takes three !gpu.mma_matrixs as arguments. All of them hold A, B and Coperands for the mma operation. The operation performed is represented as D = A * B + C. The op returns a !gpu.mma_matrix which contains the result of the operation held by the current thread.

This op is meant to be used along with gpu.subgroup_mma_store_matrix and gpu.subgroup_mma_load_matrix.

Example:

%D = gpu.subgroup_mma_compute_matrix %A, %B, %C :
!gpu.mma_matrix<16x16xf16, "AOp">, !gpu.mma_matrix<16x16xf16, "BOp">,
!gpu.mma_matrix<16x16xf16, "COp"> -> !gpu.mma_matrix<16x16xf16, "DOp">

Operands: ¶

Results: ¶

gpu.subgroup_mma_load_matrix (::mlir::gpu::SubgroupMmaLoadMatrixOp) ¶

GPU warp synchronous matrix load

Syntax:

operation ::= `gpu.subgroup_mma_load_matrix` $srcMemref`[`$indices`]` attr-dict `:` type($srcMemref) `->` type($res)

The gpu.subgroup_mma_load_matrix operation loads a matrix collectively using all the threads in a subgroup.

This operation takes a memref as argument. It is the source matrix from which data is to be loaded. The op returns a !gpu.mma_matrix. The source memref can be in the global or shared memory space. The starting of the load address is determined using indices provided. The matrix being loaded is specified in the result type. This attribute is necessary because there exists a different LLVM intrinsic for loading each operand, This is probably because all operands need to be laid out in a specific/different way for the operation in the registers. leadDimension attribute specifies the leading dimension of the source matrix.

This op is meant to be used along with gpu.subgroup_mma_store_matrix and gpu.subgroup_mma_compute.

Example:

 %0 = gpu.subgroup_mma_load_matrix src[%i,%j] : {leadDimension = 32
: i32} : memref<32x32xf16, 3>, !gpu.mma_matrix<16x16xf16, "AOp">

Attributes: ¶

Operands: ¶

Results: ¶

gpu.subgroup_mma_store_matrix (::mlir::gpu::SubgroupMmaStoreMatrixOp) ¶

GPU warp synchronous matrix store

Syntax:

operation ::= `gpu.subgroup_mma_store_matrix` $src`,` $dstMemref`[`$indices`]` attr-dict `:` type($src)`,` type($dstMemref)

The gpu.subgroup_mma_store_matrix operation stores a matrix collectively using all the threads in a subgroup.

This operation takes a !gpu.mma_matrix and a memref as arguments. !gpu.mma_matrix is the source which contains the data to be stored. The destination can be in the global or shared memory space. The starting of store address is determined using indices provided. The leadDimension attribute specifies the leading dimension of the destination matrix.

This op is meant to be used along with gpu.subgroup_mma_load_matrix and gpu.subgroup_mma_compute.

Example:

gpu.subgroup_mma_store_matrix %D, %sg[%i,%j] : { leadDimension = 32 : i32} :
!gpu.mma_matrix<16x16xf16, "DOp">, memref<32x32xf16, 3>

Attributes: ¶

Operands: ¶

gpu.subgroup_size (::mlir::gpu::SubgroupSizeOp) ¶

Syntax:

operation ::= `gpu.subgroup_size` attr-dict `:` type($result)

Returns the number of threads within a subgroup.

Example:

%sgSz = gpu.subgroup_size : index

Results: ¶

gpu.terminator (::mlir::gpu::TerminatorOp) ¶

Terminator for GPU launch regions.

A terminator operation for regions that appear in the body of gpu.launch operation. These regions are not expected to return any value so the terminator takes no operands.

gpu.thread_id (::mlir::gpu::ThreadIdOp) ¶

Returns the thread id, i.e. the index of the current thread within the block along the x, y, or z dimension.

Example:

%tIdX = "gpu.thread_id"() {dimension = "x"} : () -> (index)

Attributes: ¶

Results: ¶

gpu.wait (::mlir::gpu::WaitOp) ¶

Wait for async gpu ops to complete.

Syntax:

operation ::= `gpu.wait` custom<AsyncDependencies>(type($asyncToken), $asyncDependencies) attr-dict

This op synchronizes the host or the device with a list of dependent ops.

If the op contains the async keyword, it returns a new async token which is synchronized with the op arguments. This new token is merely a shortcut to the argument list, and one could replace the uses of the result with the arguments for the same effect. The async version of this op is primarily used to make each async token have a single use during lowering and thereby make forks in async execution explicit. Example usage:

%t0 = gpu.foo async : !gpu.async.token
%t1 = gpu.bar async : !gpu.async.token
%t2 = gpu.wait async [%t0, %t1]
// gpu.baz doesn't run until gpu.foo and gpu.bar have both completed, just
// as if the async dependencies were [%t0, %t1].
%t3 = gpu.baz async [%t2]

If the op does not contain the async keyword, it does not return a new async token but blocks until all ops producing the async dependency tokens finished execution. All dependent memory operations are visible to the host once this op completes. Example usage:

%t0 = gpu.foo async : !gpu.async.token
%t1 = gpu.bar async : !gpu.async.token
// The gpu.wait op blocks until gpu.foo and gpu.bar have completed.
gpu.wait [%t0, %t1]

Operands: ¶

Results: ¶

gpu.yield (::mlir::gpu::YieldOp) ¶

GPU yield operation

gpu.yield` is a special terminator operation for blocks inside regions in gpu ops. It returns values to the immediately enclosing gpu op.

Example:

gpu.yield %f0, %f1 : f32, f32

Operands: ¶