MLIR

https://mlir.llvm.org
Admin
MLIR (Multi-Level Intermediate Representation) is a novel and extensible compiler infrastructure...
Tokens:12,478,394
Snippets:116,650
Trust Score:9.5
Update:1 week ago
Show doc for...
Tokens:
Raw
### nvgpu.mbarrier.init example

Source: https://mlir.llvm.org/docs/Dialects/NVGPU

Example of the nvgpu.mbarrier.init operation.

```mlir
  %num_threads = gpu.block_dim x
  %barrier = nvgpu.mbarrier.create -> !nvgpu.mbarrier.barrier<memorySpace = #gpu.address_space<workgroup>>
  nvgpu.mbarrier.init %barrier, %num_threads : !nvgpu.mbarrier.barrier<memorySpace = #gpu.address_space<workgroup>>

```

--------------------------------

### gpu.grid_dim example

Source: https://mlir.llvm.org/python-bindings/autoapi/mlir/dialects/_gpu_ops_gen/index.html

Example of using the gpu.grid_dim operation to get the number of thread blocks in the z dimension.

```mlir
%gDimZ = gpu.grid_dim z
```

--------------------------------

### initialize

Source: https://mlir.llvm.org/doxygen/mlir_2IR_2Remarks_8h_source.html

Setup the remark engine with the given output path and format.

```APIDOC
## initialize

### Description
Setup the remark engine with the given output path and format.

### Signature
```cpp
LogicalResult initialize(std::unique_ptr<MLIRRemarkStreamerBase> streamer, std::unique_ptr<RemarkEmittingPolicyBase> remarkEmittingPolicy, std::string *errMsg);
```

### Parameters
- **streamer** (std::unique_ptr<MLIRRemarkStreamerBase>) - Required - The streamer to use for emitting remarks.
- **remarkEmittingPolicy** (std::unique_ptr<RemarkEmittingPolicyBase>) - Required - The policy to use for emitting remarks.
- **errMsg** (std::string*) - Optional - If provided, will be set to an error message on failure.

### Returns
- **LogicalResult** - Indicates success or failure.
```

--------------------------------

### pdl_interp.get_operands syntax and examples

Source: https://mlir.llvm.org/docs/Dialects/PDLInterpOps

Shows the syntax and usage of the pdl_interp.get_operands operation, which tries to get a specific operand group from an operation. If no index is provided, it returns all operands. The examples show getting a single element, a range, and all operands.

```mlir
operation ::= `pdl_interp.get_operands` ($index^)? `of` $inputOp `:` type($value) attr-dict
```

```mlir
// Get the first group of operands from an operation, and expect a single
// element.
%operand = pdl_interp.get_operands 0 of %op : !pdl.value

// Get the first group of operands from an operation.
%operands = pdl_interp.get_operands 0 of %op : !pdl.range<value>

// Get all of the operands from an operation.
%operands = pdl_interp.get_operands of %op : !pdl.range<value>
```

--------------------------------

### wasmssa.local_get example

Source: https://mlir.llvm.org/docs/Dialects/WasmSSAOps

Gets the value of a local variable and returns a reference to it. The example retrieves a reference to `%a`, a float32 local.

```mlir
// Retrieves a reference to `%a`, a float32 local
%b = wasmssa.local_get %a : ref to f32
```

--------------------------------

### nvgpu.mbarrier.create example

Source: https://mlir.llvm.org/docs/Dialects/NVGPU

Example of the nvgpu.mbarrier.create operation.

```mlir
  %barrier = nvgpu.mbarrier.create -> !nvgpu.mbarrier.barrier<memorySpace = #gpu.address_space<workgroup>>

```

--------------------------------

### gpu.subgroup_id MLIR example

Source: https://mlir.llvm.org/python-bindings/_sources/autoapi/mlir/dialects/_gpu_ops_gen/index.rst.txt

Example showing how to get the subgroup id using the gpu.subgroup_id operation. The result is an index value.

```mlir
%sgId = gpu.subgroup_id : index
```

--------------------------------

### acc.init operation examples

Source: https://mlir.llvm.org/python-bindings/autoapi/mlir/dialects/_acc_ops_gen/index.html

Shows two usage examples of the acc.init operation: one without arguments and one with a device_num argument.

```mlir
acc.init
```

```mlir
acc.init device_num(%dev1 : i32)
```

--------------------------------

### spirv.mlir.addressof example

Source: https://mlir.llvm.org/docs/Dialects/SPIR-V

Example of getting the address of a global variable @global_var, returning a pointer to f32 in Input storage class.

```mlir
%0 = spirv.mlir.addressof @global_var : !spirv.ptr<f32, Input>

```

--------------------------------

### mlir.dialects._nvgpu_ops_gen.mbarrier_init

Source: https://mlir.llvm.org/python-bindings/_sources/autoapi/mlir/dialects/_nvgpu_ops_gen/index.rst.txt

Creates an MBarrierInitOp for initializing an mbarrier.

```APIDOC
## Function: mbarrier_init

### Description
Creates an `MBarrierInitOp` for initializing an mbarrier.

### Signature
`mbarrier_init(...)`

### Parameters
- (Parameters not explicitly documented in source.)

### Returns
- **MBarrierInitOp** - The created operation.
```

--------------------------------

### llvm.mlir.addressof usage example

Source: https://mlir.llvm.org/docs/Dialects/LLVM

Example showing how to get addresses of globals, functions, and aliases, and use them for loads and indirect calls.

```mlir
func @foo() {
  // Get the address of a global variable.
  %0 = llvm.mlir.addressof @const : !llvm.ptr

  // Use it as a regular pointer.
  %1 = llvm.load %0 : !llvm.ptr -> i32

  // Get the address of a function.
  %2 = llvm.mlir.addressof @foo : !llvm.ptr

  // The function address can be used for indirect calls.
  llvm.call %2() : !llvm.ptr, () -> ()

  // Get the address of an aliased global.
  %3 = llvm.mlir.addressof @const_alias : !llvm.ptr
}

// Define the global.
llvm.mlir.global @const(42 : i32) : i32

// Define an alias.
llvm.mlir.alias @const_alias : i32 {
  %0 = llvm.mlir.addressof @const : !llvm.ptr
  llvm.return %0 : !llvm.ptr
}
```

--------------------------------

### Example using tile_sizes

Source: https://mlir.llvm.org/python-bindings/autoapi/mlir/dialects/_structured_transform_ops_gen/index.html

Shows how to use tile_using_forall with tile_sizes, including a dynamic size from a matched operation.

```mlir
%0 = transform.structured.match ops{["linalg.matmul"]} in %arg1
   : (!transform.any_op) -> !transform.any_op
%sz = transform.structured.match ...
%3:2 = transform.structured.tile_using_forall %0 tile_sizes [0, %sz, 20]
   : (!transform.any_op, !transform.any_op) -> (!transform.any_op, !transform.any_op)

```

--------------------------------

### getStartOperation

Source: https://mlir.llvm.org/doxygen/classmlir_1_1LivenessBlockInfo.html

Gets the start operation for the given value.

```APIDOC
## getStartOperation

### Description
Gets the start operation for the given value.

### Method
Member function

### Signature
`Operation *getStartOperation(Value value) const`

### Parameters
- **value** (Value) - The value to get the start operation for.

### Returns
- **Operation*** - The start operation for the given value.
```

--------------------------------

### mbarrier_init()

Source: https://mlir.llvm.org/python-bindings/autoapi/mlir/dialects/nvgpu/index.html

Creates an MBarrierInitOp.

```APIDOC
## mbarrier_init()

### Description
Creates an MBarrierInitOp.

### Method
Function call

### Signature
`mbarrier_init()`

### Parameters
No parameters documented.

### Returns
An MBarrierInitOp instance.
```

--------------------------------

### affine.dma_start example 2: with stride arguments

Source: https://mlir.llvm.org/docs/Dialects/Affine

Example of an affine.dma_start operation with optional stride and number_of_elements_per_stride arguments, transferring chunks of elements every stride.

```mlir
affine.dma_start %src[%i, %j], %dst[%k, %l], %tag[%idx], %num_elements,
  %stride, %num_elt_per_stride : ...

```

--------------------------------

### tensor.dim syntax and examples

Source: https://mlir.llvm.org/docs/Dialects/TensorOps

Shows the syntax and three examples of the tensor.dim operation. The first two examples get static and dynamic dimensions, and the third shows the equivalent generic form.

```mlir
operation ::= `tensor.dim` attr-dict $source `,` $index `:` type($source)
```

```mlir
// Always returns 4, can be constant folded:
%c0 = arith.constant 0 : index
%x = tensor.dim %A, %c0 : tensor<4x?xf32>

// Return the dynamic dimension of %A.
%c1 = arith.constant 1 : index
%y = tensor.dim %A, %c1 : tensor<4x?xf32>

// Equivalent generic form:
%x = "tensor.dim"(%A, %c0) : (tensor<4x?xf32>, index) -> index
%y = "tensor.dim"(%A, %c1) : (tensor<4x?xf32>, index) -> index
```

--------------------------------

### nvgpu.mbarrier.get example

Source: https://mlir.llvm.org/docs/Dialects/NVGPU

Example of the nvgpu.mbarrier.get operation.

```mlir
  %mbars = nvgpu.mbarrier.create -> !nvgpu.mbarrier.group<memorySpace = #gpu.address_space<workgroup>, num_barriers = 10>
  %mbar_pointer = nvgpu.mbarrier.get %mbars[%c2] : !nvgpu.mbarrier.group<memorySpace = #gpu.address_space<workgroup>>

```

--------------------------------

### pdl_interp.get_value_type syntax and examples

Source: https://mlir.llvm.org/docs/Dialects/PDLInterpOps

Shows the syntax and two usage examples for the pdl_interp.get_value_type operation, which gets the type of a value or range of values.

```mlir
operation ::= `pdl_interp.get_value_type` `of` $value `:` type($result) attr-dict
```

```mlir
// Get the type of a single value.
%type = pdl_interp.get_value_type of %value : !pdl.type

// Get the types of a value range.
%type = pdl_interp.get_value_type of %values : !pdl.range<type>
```

--------------------------------

### nvgpu.mbarrier.arrive.nocomplete example

Source: https://mlir.llvm.org/docs/Dialects/NVGPU

Example of the nvgpu.mbarrier.arrive.nocomplete operation.

```mlir
  %token = nvgpu.mbarrier.arrive.noComplete %barrier, %count : !nvgpu.mbarrier.barrier<memorySpace = #gpu.address_space<workgroup>> -> !nvgpu.mbarrier.token

```

--------------------------------

### spirv.Tosa.Slice examples

Source: https://mlir.llvm.org/docs/Dialects/SPIR-V

Two examples of the spirv.Tosa.Slice operation. The first example uses %start = [3, 10, 20] and %size = [21, 5, 2] to extract %input1[3:24, 10:15, 20:22].

```mlir
%2 = spirv.Tosa.Slice %input1, %start, %size : !spirv.arm.tensor<32x19x41xi8>, !spirv.arm.tensor<3xi32>, !spirv.arm.tensor<3xi32> -> !spirv.arm.tensor<21x5x2xi8>
%2 = spirv.Tosa.Slice %input1, %start, %size : !spirv.arm.tensor<30x45x29xf32>, !spirv.arm.tensor<3xi32>, !spirv.arm.tensor<3xi32> -> !spirv.arm.tensor<5x12x11xf32>

```

--------------------------------

### FileLineColLoc.start_col (property)

Source: https://mlir.llvm.org/python-bindings/_sources/autoapi/mlir/_mlir_libs/_mlir/ir/index.rst.txt

Gets the start column number from a FileLineColLoc.

```APIDOC
## FileLineColLoc.start_col (property)

### Description
Gets the start column number from a ``FileLineColLoc``.

### Property
- **start_col** (int) - The start column number.
```

--------------------------------

### MBarrierInitOp

Source: https://mlir.llvm.org/python-bindings/autoapi/mlir/dialects/_nvgpu_ops_gen/index.html

Class representing the `nvgpu.mbarrier.init` operation.

```APIDOC
## MBarrierInitOp

### Description
Represents the `nvgpu.mbarrier.init` operation. Initializes the `mbarrier` object with the given number of threads.

### Operation Name
`nvgpu.mbarrier.init`

### Parameters
- `_barriers` (ODS IR) - Required - The mbarrier barrier operand.
- `_count` (ODS IR) - Required - The number of threads.
- `_mbarId` (ODS IR) - Required - The mbarrier ID operand.
- `_predicate` (ODS IR | None) - Optional - Predicate operand.
- `_loc` (ODS IR | None) - Optional - Source location.
- `_ip` (ODS IR | None) - Optional - Insertion point.
```

--------------------------------

### FileLineColLoc.start_line (property)

Source: https://mlir.llvm.org/python-bindings/_sources/autoapi/mlir/_mlir_libs/_mlir/ir/index.rst.txt

Gets the start line number from a FileLineColLoc.

```APIDOC
## FileLineColLoc.start_line (property)

### Description
Gets the start line number from a ``FileLineColLoc``.

### Property
- **start_line** (int) - The start line number.
```

--------------------------------

### nvvm.cp.async.bulk.prefetch example

Source: https://mlir.llvm.org/docs/Dialects/NVVMDialect

Example usage of nvvm.cp.async.bulk.prefetch with and without the optional l2_cache_hint operand.

```mlir
  nvvm.cp.async.bulk.prefetch %src, %size : !llvm.ptr<1>

  // with l2_cache_hint
  nvvm.cp.async.bulk.prefetch %src, %size l2_cache_hint = %ch : !llvm.ptr<1>
```

--------------------------------

### async.await_all example

Source: https://mlir.llvm.org/docs/Dialects/AsyncDialect

Example of waiting for all tokens in a group to become ready.

```mlir
%0 = async.create_group %size : !async.group

%1 = ... : !async.token
%2 = async.add_to_group %1, %0 : !async.token

%3 = ... : !async.token
%4 = async.add_to_group %2, %0 : !async.token

async.await_all %0
```

--------------------------------

### spirv.SMulExtended example

Source: https://mlir.llvm.org/docs/Dialects/SPIR-V

Shows an example for the spirv.SMulExtended operation. The result type must be a struct with two members of the same integer scalar or vector type; operands have that same type. Member 0 gets the low-order bits and member 1 gets the high-order bits of the signed multiplication.

```mlir
%2 = spirv.SMulExtended %0, %1 : !spirv.struct<(i32, i32)>
%2 = spirv.SMulExtended %0, %1 : !spirv.struct<(vector<2xi32>, vector<2xi32>)>

```

--------------------------------

### async.execute example

Source: https://mlir.llvm.org/docs/Dialects/AsyncDialect

Example of async.execute: starts asynchronous execution of a body region after dependencies and value operands are ready. The body computes a value and yields it.

```mlir
%dependency = ... : !async.token
%value = ... : !async.value<f32>

%token, %results =
  async.execute [%dependency](%value as %unwrapped: !async.value<f32>)
             -> !async.value<!some.type>
  {
    %0 = "compute0"(%unwrapped): (f32) -> !some.type
    async.yield %0 : !some.type
  }

%1 = "compute1"(...) : !some.type
```

--------------------------------

### Example of nvgpu.device_async_create_group

Source: https://mlir.llvm.org/docs/Dialects/NVGPU

An example of creating an empty async group with no input tokens.

```mlir
%0 = nvgpu.device_async_create_group
```

--------------------------------

### mlirExecutionEngineInitialize

Source: https://mlir.llvm.org/doxygen/CAPI_2ExecutionEngine_2ExecutionEngine_8cpp.html

Initialize the ExecutionEngine.

```APIDOC
## mlirExecutionEngineInitialize

### Description
Initialize the ExecutionEngine.

### Signature
```c
void mlirExecutionEngineInitialize(MlirExecutionEngine jit)
```

### Parameters
- **jit** (MlirExecutionEngine) - Required - The execution engine to initialize.
```

--------------------------------

### pdl_interp.get_attribute syntax and example

Source: https://mlir.llvm.org/docs/Dialects/PDLInterpOps

Shows the syntax and an example of the pdl_interp.get_attribute operation, which tries to get a specific attribute from an operation, returning null if the attribute is absent.

```mlir
operation ::= `pdl_interp.get_attribute` $name `of` $inputOp attr-dict
```

```mlir
%attr = pdl_interp.get_attribute "attr" of %op
```

--------------------------------

### pdl_interp.get_attribute_type syntax and example

Source: https://mlir.llvm.org/docs/Dialects/PDLInterpOps

Shows the syntax and usage of the pdl_interp.get_attribute_type operation, which gets the result type of a specified attribute. The example retrieves the type of %attr into %type.

```mlir
operation ::= `pdl_interp.get_attribute_type` `of` $value attr-dict
```

```mlir
%type = pdl_interp.get_attribute_type of %attr
```

--------------------------------

### Example of nvgpu.device_async_wait

Source: https://mlir.llvm.org/docs/Dialects/NVGPU

An example of waiting on an async token.

```mlir
nvgpu.device_async_wait %0
```

--------------------------------

### nvgpu.mbarrier.init syntax

Source: https://mlir.llvm.org/docs/Dialects/NVGPU

Syntax for the nvgpu.mbarrier.init operation.

```mlir
operation ::= `nvgpu.mbarrier.init` $barriers `[` $mbarId `]` `,` $count (`,` `predicate` `=` $predicate^)? attr-dict `:` type($barriers)

```

--------------------------------

### x86.avx.intr.dot MLIR example

Source: https://mlir.llvm.org/python-bindings/_sources/autoapi/mlir/dialects/x86/index.rst.txt

Example of the x86.avx.intr.dot operation, which computes 4-way dot products and broadcasts results, followed by extraction and addition to get the full dot product.

```mlir
%0 = x86.avx.intr.dot %a, %b : vector<8xf32>
%1 = vector.extract %0[%i0] : f32 from vector<8xf32>
%2 = vector.extract %0[%i4] : f32 from vector<8xf32>
%d = arith.addf %1, %2 : f32
```

--------------------------------

### makeVscaleConstantBuilder example

Source: https://mlir.llvm.org/doxygen/VectorUtils_8h_source.html

Shows how to use makeVscaleConstantBuilder to create a functor that returns a constant vscale multiple. The example creates a builder and uses it to get a value representing 4 * vector.vscale.

```c++
auto createVscaleMultiple = makeVscaleConstantBuilder(rewriter, loc);
auto c4Vscale = createVscaleMultiple(4); // 4 * vector.vscale
```

--------------------------------

### nvgpu.mbarrier.arrive MLIR example

Source: https://mlir.llvm.org/python-bindings/_sources/autoapi/mlir/dialects/_nvgpu_ops_gen/index.rst.txt

Performs an arrive-on operation on an mbarrier and returns a token. The example shows the MLIR syntax for the basic arrive operation.

```mlir
%token = nvgpu.mbarrier.arrive %barrier : !nvgpu.mbarrier.barrier<memorySpace = #gpu.address_space<workgroup>> -> !nvgpu.mbarrier.token
```

--------------------------------

### mlir::getPerfectlyNestedLoops

Source: https://mlir.llvm.org/doxygen/Dialect_2SCF_2Utils_2Utils_8cpp_source.html

Gets a perfectly nested sequence of loops starting at the root of a loop nest.

```APIDOC
## mlir::getPerfectlyNestedLoops

### Description
Get perfectly nested sequence of loops starting at root of loop nest.

### Method
void getPerfectlyNestedLoops(SmallVectorImpl<scf::ForOp> &nestedLoops, scf::ForOp root)

### Parameters
- **nestedLoops** (SmallVectorImpl<scf::ForOp>&) - Required - Output vector to store the nested loops.
- **root** (scf::ForOp) - Required - The root loop.
```

--------------------------------

### affine.dma_start example with and without stride arguments

Source: https://mlir.llvm.org/docs/Dialects/Affine

Shows two forms of affine.dma_start: one without stride arguments and one with stride and num_elt_per_stride. The first transfers 256 elements from %src to %dst using a tag. The second demonstrates the optional stride arguments for slower memory space transfers.

```mlir
%num_elements = arith.constant 256
%idx = arith.constant 0 : index
%tag = memref.alloc() : memref<1xi32, 4>
affine.dma_start %src[%i + 3, %j], %dst[%k + 7, %l], %tag[%idx],
  %num_elements :
    memref<40x128xf32, 0>, memref<2x1024xf32, 1>, memref<1xi32, 2>

// If %stride and %num_elt_per_stride are specified, the DMA is expected to
// transfer %num_elt_per_stride elements every %stride elements apart from
// memory space 0 until %num_elements are transferred.
affine.dma_start %src[%i, %j], %dst[%k, %l], %tag[%idx], %num_elements,
  %stride, %num_elt_per_stride : ...
```

--------------------------------

### acc.on_device syntax and example

Source: https://mlir.llvm.org/docs/Dialects/OpenACCDialect

Shows the assembly syntax for the acc.on_device operation and an example that folds a constant device type to a compile-time result when device specialization runs. The device type must be one of the defined accelerator types.

```mlir
operation ::= `acc.on_device` $deviceType attr-dict `:` type($deviceType) `->` type($result)

```

```mlir
%host = arith.constant 1 : i32
%on_host = acc.on_device %host : i32 -> i1

```

--------------------------------

### Create a TypeConverter converting FloatType to IntegerType

Source: https://mlir.llvm.org/docs/DialectConversion

Shows how to define a type converter that turns FloatTypes into IntegerTypes, used as the setup for the MyConversionPattern example below. This is one part of the complete conversion pattern example.

```cpp
TypeConverter converter;
converter.addConversion([](FloatType t) {
    return IntegerType::get(t.getContext(), t.getWidth());
});

```

--------------------------------

### affine.load example 1

Source: https://mlir.llvm.org/docs/Dialects/Affine

Basic affine.load example with affine expressions using loop induction variables.

```mlir
%1 = affine.load %0[%i0 + 3, %i1 + 7] : memref<100x100xf32>

```

--------------------------------

### mlir::affine::getPerfectlyNestedLoops

Source: https://mlir.llvm.org/doxygen/LoopUtils_8h.html

Gets the perfectly nested sequence of loops starting at a given root loop.

```APIDOC
## mlir::affine::getPerfectlyNestedLoops

### Description
Get perfectly nested sequence of loops starting at root of loop nest (the first op being another AffineFor, and the second op - a terminator).

### Signature
```cpp
void mlir::affine::getPerfectlyNestedLoops(SmallVectorImpl<AffineForOp> &nestedLoops, AffineForOp root)
```

### Parameters
- **nestedLoops** (SmallVectorImpl<AffineForOp>&) - Output parameter to store the nested loops.
- **root** (AffineForOp) - The root loop to start from.

### Returns
- **void** - This function does not return a value.
```

--------------------------------

### mlir::FlatLinearValueConstraints::mergeAndAlignVarsWithOther

Source: https://mlir.llvm.org/doxygen/Dialect_2Affine_2Analysis_2Utils_8cpp_source.html

Merge and align the variables of this and other starting at offset, so that both constraint systems get the same variables.

```APIDOC
## mlir::FlatLinearValueConstraints::mergeAndAlignVarsWithOther

### Description
Merge and align the variables of this and other starting at offset, so that both constraint systems get the same variables.

### Method
FlatLinearValueConstraints::mergeAndAlignVarsWithOther

### Signature
void mergeAndAlignVarsWithOther(unsigned offset, FlatLinearValueConstraints *other)

### Parameters
- **offset** (unsigned) - The offset at which to start merging.
- **other** (FlatLinearValueConstraints*) - The other constraint system.

### Source Definition
FlatLinearValueConstraints.cpp:1371
```

--------------------------------

### llvm.mlir.ifunc example

Source: https://mlir.llvm.org/docs/Dialects/LLVM

Example showing how to define an ifunc, a resolver function, and how to call the ifunc as a regular function.

```mlir
// IFuncs resolve a symbol at runtime using a resovler function.
llvm.mlir.ifunc external @foo: !llvm.func<f32 (i64)>, !llvm.ptr @resolver

llvm.func @foo_1(i64) -> f32
llvm.func @foo_2(i64) -> f32

llvm.func @resolve_foo() -> !llvm.ptr attributes {
  %0 = llvm.mlir.addressof @foo_2 : !llvm.ptr
  %1 = llvm.mlir.addressof @foo_1 : !llvm.ptr

  // ... Logic selecting from foo_{1, 2}

  // Return function pointer to the selected function
  llvm.return %7 : !llvm.ptr
}

llvm.func @use_foo() {
  // IFuncs are called as regular functions
  %res = llvm.call @foo(%value) : i64 -> f32
}
```

--------------------------------

### mlir.dialects._nvgpu_ops_gen.mbarrier_arrive

Source: https://mlir.llvm.org/python-bindings/_sources/autoapi/mlir/dialects/_nvgpu_ops_gen/index.rst.txt

Creates an MBarrierArriveOp for arriving at an mbarrier.

```APIDOC
## Function: mbarrier_arrive

### Description
Creates an `MBarrierArriveOp` for arriving at an mbarrier.

### Signature
`mbarrier_arrive(...)`

### Parameters
- (Parameters not explicitly documented in source.)

### Returns
- **MBarrierArriveOp** - The created operation.
```
Tokens:
Raw