Struct melior::dialect::ods::gpu::LaunchOperation

pub struct LaunchOperation<'c> { /* private fields */ }

A launch operation. 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 an optional list of async dependencies followed by six operands and an optional operand.

The async keyword indicates the kernel should be launched asynchronously; the operation returns a new !gpu.async.token when the keyword is specified. The kernel launched does not start executing until the ops producing its async dependencies (optional operands) have completed.

The first three operands (following any async dependencies) are grid sizes along the x,y,z dimensions and the following three are block sizes along the x,y,z dimensions. When a lower-dimensional kernel is required, unused sizes must be explicitly set to 1. The last operand is optional and corresponds to the amount of dynamic shared memory a kernel’s workgroup should be allocated; when this operand is not present, a zero size is assumed.

The body region has at least 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).
• a variadic number of Workgroup memory attributions.
• a variadic number of Private memory attributions.

Syntax:

operation ::= `gpu.launch` (`async` (`[` ssa-id-list `]`)? )?
                         `block` `(` ssa-id-list `)` `in` ssa-reassignment
                         `threads` `(` ssa-id-list `)` `in` ssa-reassignment
                         (dynamic_shared_memory_size ssa-use)?
                         memory-attribution
                         region attr-dict?
ssa-reassignment ::= `(` ssa-id `=` ssa-use (`,` ssa-id `=` ssa-use)* `)`
memory-attribution ::= (`workgroup` `(` ssa-id-and-type-list `)`)?
                       (`private` `(` ssa-id-and-type-list `)`)?

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
}

// Launch with memory attributions.
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)
           workgroup(%workgroup: memref<32xf32, 3>)
           private(%private: memref<1xf32, 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 %workgroup[%bx] : memref<32xf32, 3>
}

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.

Implementations

impl<'c> LaunchOperation<'c>

pub fn name() -> &'static str

Returns a name.

pub fn as_operation(&self) -> &Operation<'c>

Returns a generic operation.

pub fn builder( context: &'c Context, location: Location<'c> ) -> LaunchOperationBuilder<'c, Unset, Unset, Unset, Unset, Unset, Unset, Unset, Unset>

Creates a builder.

pub fn async_token(&self) -> Result<OperationResult<'c, '_>, Error>

pub fn async_dependencies( &self ) -> Result<impl Iterator<Item = Value<'c, '_>>, Error>

pub fn grid_size_x(&self) -> Result<Value<'c, '_>, Error>

pub fn grid_size_y(&self) -> Result<Value<'c, '_>, Error>

pub fn grid_size_z(&self) -> Result<Value<'c, '_>, Error>

pub fn block_size_x(&self) -> Result<Value<'c, '_>, Error>

pub fn block_size_y(&self) -> Result<Value<'c, '_>, Error>

pub fn block_size_z(&self) -> Result<Value<'c, '_>, Error>

pub fn dynamic_shared_memory_size(&self) -> Result<Value<'c, '_>, Error>

pub fn body(&self) -> Result<RegionRef<'c, '_>, Error>

Trait Implementations

impl<'c> From<LaunchOperation<'c>> for Operation<'c>

fn from(operation: LaunchOperation<'c>) -> Self

Converts to this type from the input type.

impl<'c> TryFrom<Operation<'c>> for LaunchOperation<'c>

type Error = Error

The type returned in the event of a conversion error.

fn try_from(operation: Operation<'c>) -> Result<Self, Self::Error>

Performs the conversion.