Return op name nvvm.addf as a bitstring.

nvvm.addf -

Performs floating point addition of the given arguments `lhs` and `rhs`

This op has support for result type inference.

Attributes

• rnd - Single, FPArithRoundingMode, NVVM FPRoundingMode kind whose value is one of {none, rm, rn, rp, rz}
• sat - Single, SaturationModeSatOrNone, Describes the saturation mode whose value is one of {none, sat}
• ftz - Single, BoolAttr, bool attribute

Operands

• lhs - Single, SIMTFloatType, 16-bit float or bfloat16 type or 32-bit float or 64-bit float or vector of 16-bit float or bfloat16 type or 32-bit float or 64-bit float values of length 2
• rhs - Single, SIMTFloatType, 16-bit float or bfloat16 type or 32-bit float or 64-bit float or vector of 16-bit float or bfloat16 type or 32-bit float or 64-bit float values of length 2

Results

• res - Single, SIMTFloatType, 16-bit float or bfloat16 type or 32-bit float or 64-bit float or vector of 16-bit float or bfloat16 type or 32-bit float or 64-bit float values of length 2

Description

The nvvm.addf operation performs floating point addition of two floating point operands of the same type.

The rounding mode is specified by the rnd attribute, saturation mode by the sat attribute, and flush-to-zero by the ftz attribute.

For more information, see PTX ISA:

• floating point addition
• half-precision floating point addition

Return op name nvvm.bar.warp.sync as a bitstring.

nvvm.bar.warp.sync - Warp Barrier Synchronization Op

Operands

• mask - Single, I32, 32-bit signless integer

Description

The nvvm.bar.warp.sync operation performs barrier synchronization for threads within a warp.

This operation causes the executing thread to wait until all threads corresponding to the mask operand have executed a bar.warp.sync with the same mask value before resuming execution.

The mask operand specifies the threads participating in the barrier, where each bit position corresponds to the thread's lane ID within the warp. Only threads with their corresponding bit set in the mask participate in the barrier synchronization.

Important constraints:

• The behavior is undefined if the executing thread is not included in the mask (i.e., the bit corresponding to the thread's lane ID is not set)
• For compute capability sm_6x or below, all threads in the mask must execute the same bar.warp.sync instruction in convergence

This operation also guarantees memory ordering among participating threads. Threads within the warp that wish to communicate via memory can store to memory, execute bar.warp.sync, and then safely read values stored by other threads in the warp.

For more information, see PTX ISA

Return op name nvvm.barrier as a bitstring.

nvvm.barrier - CTA Barrier Synchronization Op

Attributes

• aligned - Single, BoolAttr, bool attribute

Operands

• barrierId - Optional, I32, 32-bit signless integer
• numberOfThreads - Optional, I32, 32-bit signless integer

Description

The nvvm.barrier operation performs barrier synchronization and communication within a CTA (Cooperative Thread Array). It causes executing threads to wait for all non-exited threads participating in the barrier to arrive.

The operation takes the following optional operands and attributes:

• barrierId: Specifies a logical barrier resource with value 0 through 15. Each CTA instance has sixteen barriers numbered 0..15. Defaults to 0 if not specified.
• numberOfThreads: Specifies the number of threads participating in the barrier. When specified, the value must be a multiple of the warp size. If not specified, all threads in the CTA participate in the barrier.
• aligned: Selects between the .aligned and non-.aligned forms of the underlying @llvm.nvvm.barrier.cta.* intrinsic family. Defaults to true, which requires every thread in the CTA to reach this same barrier instruction, otherwise the behavior is undefined. Set it to false to emit the non-.aligned form.

Reduction variants of the barrier instruction are modeled by the nvvm.barrier.reduction op.

The barrier operation guarantees that when the barrier completes, prior memory accesses requested by participating threads are performed relative to all threads participating in the barrier. It also ensures that no new memory access is requested by participating threads before the barrier completes.

When a barrier completes, the waiting threads are restarted without delay, and the barrier is reinitialized so that it can be immediately reused.

For more information, see PTX ISA

Return op name nvvm.barrier.arrive as a bitstring.

nvvm.barrier.arrive

Attributes

• aligned - Single, BoolAttr, bool attribute

Operands

• barrierId - Optional, I32, 32-bit signless integer
• numberOfThreads - Single, I32, 32-bit signless integer

Description

Thread that executes this op announces their arrival at the barrier with given id and continue their execution.

The default barrier id is 0 that is similar to nvvm.barrier Op. When barrierId is not present, the default barrier id is used.

The aligned attribute, which defaults to true, generates the aligned form of the barrier (all threads in the CTA execute the same barrier instruction). When set to false, the unaligned form is generated.

For more information, see PTX ISA

Return op name nvvm.barrier.reduction as a bitstring.

nvvm.barrier.reduction - CTA Barrier Reduction Op

This op has support for result type inference.

Attributes

• reductionOp - Single, BarrierReductionAttr, NVVM barrier reduction operation
• aligned - Single, BoolAttr, bool attribute

Operands

• barrierId - Optional, I32, 32-bit signless integer
• reductionPredicate - Single, I32, 32-bit signless integer

Results

• res - Single, I32, 32-bit signless integer

Description

The nvvm.barrier.reduction operation performs barrier synchronization with a reduction across the per-thread predicates contributed by participating threads in a CTA.

• barrierId: Specifies a logical barrier resource with value 0 through 15. Optional; defaults to barrier id 0 when not specified.
• reductionOp: The reduction kind (popc, and, or) applied across the per-thread predicates.
• reductionPredicate: The per-thread i32 predicate. It is compared against zero to form the i1 value fed into the reduction.
• aligned: Selects between the .aligned and non-.aligned forms of the underlying @llvm.nvvm.barrier.cta.red.* intrinsic family. Defaults to true, which requires every thread in the CTA to reach this same barrier instruction, otherwise the behavior is undefined. Set it to false to emit the non-.aligned form.

The result is the i32 reduction value computed across all threads participating in the barrier.

For more information, see PTX ISA

Return op name nvvm.breakpoint as a bitstring.

nvvm.breakpoint - Breakpoint Op

Description

Breakpoint suspends execution of the program for debugging. For more information, see PTX ISA

Return op name nvvm.cluster.arrive as a bitstring.

nvvm.cluster.arrive - Cluster Barrier Arrive Op

Attributes

• aligned - Optional, UnitAttr, unit attribute

Description

The cluster.arrive can be used by the threads within the cluster for synchronization and communication. The cluster.arrive instruction marks the warps' arrival at the barrier without causing the executing thread to wait for other participating threads.

The aligned attribute, when provided, generates the .aligned version of the PTX instruction.

For more information, see PTX ISA

Return op name nvvm.cluster.arrive.relaxed as a bitstring.

nvvm.cluster.arrive.relaxed - Cluster Barrier Relaxed Arrive Op

Attributes

• aligned - Optional, UnitAttr, unit attribute

Description

The cluster.arrive can be used by the threads within the cluster for synchronization and communication. The cluster.arrive instruction marks the warps' arrival at the barrier without causing the executing thread to wait for other participating threads.

The aligned attribute, when provided, generates the .aligned version of the PTX instruction. The .relaxed qualifier on cluster.arrive specifies that there are no memory ordering and visibility guarantees provided for the memory accesses performed prior to cluster.arrive.

For more information, see PTX ISA

Return op name nvvm.cluster.wait as a bitstring.

nvvm.cluster.wait - Cluster Barrier Wait Op

Attributes

• aligned - Optional, UnitAttr, unit attribute

Description

The cluster.wait causes the executing thread to wait for all non-exited threads of the cluster to perform cluster.arrive. The aligned attribute, when provided, generates the .aligned version of the PTX instruction.

For more information, see PTX ISA

Return op name nvvm.clusterlaunchcontrol.query.cancel as a bitstring.

nvvm.clusterlaunchcontrol.query.cancel - Query the response of a clusterlaunchcontrol.try.cancel operation

This op has support for result type inference.

Attributes

• query_type - Single, ClusterLaunchControlQueryTypeAttr, NVVM ClusterLaunchControlQueryType

Operands

• try_cancel_response - Single, I128, 128-bit signless integer

Results

• res - Single, anonymous/composite constraint, 1-bit signless integer or 32-bit signless integer

Description

clusterlaunchcontrol.query.cancel queries the response of a clusterlaunchcontrol.try.cancel operation specified by operand try_cancel_response.

Operand query_type specifies the type of query to perform and can be one of the following:

• is_canceled : Returns true if the try cancel request succeeded, and false otherwise.
• get_first_cta_id_{x/y/z} : Returns the x, y, or z coordinate of the first CTA in the canceled cluster. Behaviour is defined only if the try cancel request succeeded.

For more information, see PTX ISA

Return op name nvvm.clusterlaunchcontrol.try.cancel as a bitstring.

nvvm.clusterlaunchcontrol.try.cancel - Request atomically canceling the launch of a cluster that has not started running yet

Attributes

• multicast - Optional, UnitAttr, unit attribute

Operands

• smemAddress - Single, LLVM_PointerShared, LLVM pointer in address space 3
• mbarrier - Single, LLVM_PointerShared, LLVM pointer in address space 3

Description

clusterlaunchcontrol.try.cancel requests atomically canceling the launch of a cluster that has not started running yet. It asynchronously writes an opaque response to shared memory indicating whether the operation succeeded or failed.

Operand smemAddress specifies the naturally aligned address of the 16-byte wide shared memory location where the request's response is written.

Operand mbarrier specifies the mbarrier object used to track the completion of the asynchronous operation.

If multicast is specified, the response is asynchronously written to the corresponding local shared memory location (specifed by addr) of each CTA in the requesting cluster.

For more information, see PTX ISA

Return op name nvvm.convert.bf16x2.to.f4x2 as a bitstring.

nvvm.convert.bf16x2.to.f4x2 - Convert an bf16x2 input to f4x2

This op has support for result type inference.

Attributes

• relu - Single, BoolAttr, bool attribute
• dstTy - Single, anonymous/composite constraint, type attribute of f4E2M1FN type

Operands

• src - Single, anonymous/composite constraint, vector of bfloat16 type values of length 2

Results

• dst - Single, I8, 8-bit signless integer

Description

This Op converts each of the given BF16 inputs in an bf16x2 vector to the specified fp4 type. The result dst is returned as an i8 type where the converted values are packed such that the value converted from the first element of a is stored in the lower 4 bits of dst and the value converted from the second element of a is stored in the upper 4 bits of dst. The relu attribute, when set, lowers to the '.relu' variant of the cvt instruction.

Return op name nvvm.convert.bf16x2.to.f6x2 as a bitstring.

nvvm.convert.bf16x2.to.f6x2 - Convert an bf16x2 input to f6x2

Attributes

• relu - Single, BoolAttr, bool attribute
• dstTy - Single, anonymous/composite constraint, type attribute of f6E2M3FN type or f6E3M2FN type

Operands

• src - Single, anonymous/composite constraint, vector of bfloat16 type values of length 2

Results

• dst - Single, anonymous/composite constraint, 16-bit signless integer or vector of 8-bit signless integer values of length 2

Description

This Op converts each of the given BF16 inputs in an bf16x2 vector to the specified fp6 type. The result dst is represented either as an i16 type or as a vector of two i8 types. If dst is returned as an i16 type, the converted values are packed such that the value converted from the first element of a is stored in the lower 8 bits of dst with 2 MSB bits padded with zeros and the value converted from the second element of a is stored in the upper 8 bits of dst with 2 MSB bits padded with zeros. If dst is returned as a vector type, each converted value is stored as an i8 element in the vector with 2 MSB bits padded with zeros. The relu attribute, when set, lowers to the '.relu' variant of the cvt instruction.

Return op name nvvm.convert.bf16x2.to.f8x2 as a bitstring.

nvvm.convert.bf16x2.to.f8x2 - Convert a pair of bf16 inputs to f8x2

Attributes

• rnd - Single, FPRoundingModeAttr, NVVM FPRoundingMode kind
• sat - Single, SaturationModeAttr, Describes the saturation mode
• relu - Single, BoolAttr, bool attribute
• dstTy - Single, anonymous/composite constraint, type attribute of f8E8M0FNU type or f8E4M3FN type or f8E5M2 type

Operands

• src - Single, anonymous/composite constraint, vector of bfloat16 type values of length 2

Results

• dst - Single, anonymous/composite constraint, 16-bit signless integer or vector of 8-bit signless integer values of length 2

Description

This Op converts the given bf16 inputs in a bf16x2 vector to the specified f8 type. The result dst is represented either as a packed i16 type or as a vector of two i8 types. If dst is returned as an i16 type, the converted values are packed such that the value converted from the first element of a is stored in the lower 8 bits of dst and the value converted from the second element of a is stored in the upper 8 bits of dst. If dst is returned as a vector type, each converted value is stored as an i8 element in the vector. The rnd and sat attributes specify the rounding and saturation modes respectively.

For more information, see PTX ISA

Return op name nvvm.convert.bf16x2.to.s2f6x2 as a bitstring.

nvvm.convert.bf16x2.to.s2f6x2 - Convert a pair of BF16 inputs to S2F6x2

Attributes

• relu - Single, BoolAttr, bool attribute

Operands

• src - Single, anonymous/composite constraint, vector of bfloat16 type values of length 2
• scaleFactor - Optional, I16, 16-bit signless integer

Results

• dst - Single, anonymous/composite constraint, 16-bit signless integer or vector of 8-bit signless integer values of length 2

Description

This Op converts each of the given BF16 inputs in a bf16x2 vector to the S2F6x2 type. The result dst can be either a packed i16 type or a vector of two i8 types. If dst is returned as an i16 type, the converted values are packed such that the value converted from the first element of a is stored in the lower 8 bits of dst and the value converted from the second element of a is stored in the upper 8 bits of dst. If dst is returned as a vector type, each converted value is stored as an i8 element in the vector. The relu attribute, when set, lowers to the '.relu' variant of the cvt instruction. The optional scaling-factors for each of the inputs are provided through the operand scaleFactor as a packed i16 type. Only ue8m0 is supported as the type of the scale-factor currently.

For more information, see PTX ISA

Return op name nvvm.convert.f4x2.to.bf16x2 as a bitstring.

nvvm.convert.f4x2.to.bf16x2 - Convert a pair of f4 inputs to bf16x2

Attributes

• srcType - Single, anonymous/composite constraint, type attribute of f4E2M1FN type
• sat - Single, SaturationModeSatfiniteOrNone, Describes the saturation mode whose value is one of {none, satfinite}
• relu - Single, BoolAttr, bool attribute

Operands

• src - Single, I8, 8-bit signless integer
• scaleFactor - Optional, I16, 16-bit signless integer

Results

• dst - Single, anonymous/composite constraint, vector of bfloat16 type values of length 2

Description

This Op converts the given f4 inputs in a packed i8 to bf16.

The result dst is represented as a vector of bf16 elements.

The relu attribute, when set, lowers to the '.relu' variant of the cvt instruction.

The sat attribute specifies the saturation mode.

The optional scaling-factors for each of the inputs are provided through the operand scaleFactor as a packed i16 type. Only ue8m0 is supported as the type of the scale-factor currently.

For more information, see PTX ISA

Example:

// Basic conversion; the f4x2 source is packed in a single i8.
%res1 = nvvm.convert.f4x2.to.bf16x2 %src
    : i8 (f4E2M1FN) -> vector<2xbf16>

// Conversion with relu and saturation.
%res2 = nvvm.convert.f4x2.to.bf16x2 %src
    {relu = true, sat = #nvvm.sat_mode<satfinite>}
    : i8 (f4E2M1FN) -> vector<2xbf16>

// Conversion with a packed ue8m0 scale-factor.
%res3 = nvvm.convert.f4x2.to.bf16x2 %src, %scaleFactor
    : i8 (f4E2M1FN) -> vector<2xbf16>

Return op name nvvm.convert.f4x2.to.f16x2 as a bitstring.

nvvm.convert.f4x2.to.f16x2 - Convert a pair of f4 inputs to f16x2

Attributes

• srcType - Single, anonymous/composite constraint, type attribute of f4E2M1FN type
• relu - Single, BoolAttr, bool attribute

Operands

• src - Single, I8, 8-bit signless integer

Results

• dst - Single, anonymous/composite constraint, vector of 16-bit float values of length 2

Description

This Op converts the given f4 inputs in a packed i8 to f16.

The result dst is represented as a vector of f16 elements.

The relu attribute, when set, lowers to the '.relu' variant of the cvt instruction.

For more information, see PTX ISA

Return op name nvvm.convert.f6x2.to.bf16x2 as a bitstring.

nvvm.convert.f6x2.to.bf16x2 - Convert a pair of f6 inputs to bf16x2

Attributes

• srcType - Single, anonymous/composite constraint, type attribute of f6E2M3FN type or f6E3M2FN type
• sat - Single, SaturationModeSatfiniteOrNone, Describes the saturation mode whose value is one of {none, satfinite}
• relu - Single, BoolAttr, bool attribute

Operands

• src - Single, anonymous/composite constraint, vector of 8-bit signless integer values of length 2
• scaleFactor - Optional, I16, 16-bit signless integer

Results

• dst - Single, anonymous/composite constraint, vector of bfloat16 type values of length 2

Description

This Op converts the given f6 inputs in a i8x2 vector to bf16.

The result dst is represented as a vector of bf16 elements.

The relu attribute, when set, lowers to the '.relu' variant of the cvt instruction.

The sat attribute specifies the saturation mode.

The optional scaling-factors for each of the inputs are provided through the operand scaleFactor as a packed i16 type. Only ue8m0 is supported as the type of the scale-factor currently.

For more information, see PTX ISA

Example:

// Basic conversion from f6E2M3FN.
%res1 = nvvm.convert.f6x2.to.bf16x2 %src
    : vector<2xi8> (f6E2M3FN) -> vector<2xbf16>

// Conversion from f6E3M2FN with relu and saturation.
%res2 = nvvm.convert.f6x2.to.bf16x2 %src
    {relu = true, sat = #nvvm.sat_mode<satfinite>}
    : vector<2xi8> (f6E3M2FN) -> vector<2xbf16>

// Conversion with a packed ue8m0 scale-factor.
%res3 = nvvm.convert.f6x2.to.bf16x2 %src, %scaleFactor
    : vector<2xi8> (f6E2M3FN) -> vector<2xbf16>

Return op name nvvm.convert.f6x2.to.f16x2 as a bitstring.

nvvm.convert.f6x2.to.f16x2 - Convert a pair of f6 inputs to f16x2

Attributes

• srcType - Single, anonymous/composite constraint, type attribute of f6E2M3FN type or f6E3M2FN type
• relu - Single, BoolAttr, bool attribute

Operands

• src - Single, anonymous/composite constraint, vector of 8-bit signless integer values of length 2

Results

• dst - Single, anonymous/composite constraint, vector of 16-bit float values of length 2

Description

This Op converts the given f6 inputs in a i8x2 vector to f16.

The result dst is represented as a vector of f16 elements.

The relu attribute, when set, lowers to the '.relu' variant of the cvt instruction.

For more information, see PTX ISA

Return op name nvvm.convert.f8x2.to.bf16x2 as a bitstring.

nvvm.convert.f8x2.to.bf16x2 - Convert a pair of f8 inputs to bf16x2

Attributes

• srcType - Single, anonymous/composite constraint, type attribute of f8E8M0FNU type or f8E4M3FN type or f8E5M2 type
• sat - Single, SaturationModeSatfiniteOrNone, Describes the saturation mode whose value is one of {none, satfinite}
• relu - Single, BoolAttr, bool attribute

Operands

• src - Single, anonymous/composite constraint, vector of 8-bit signless integer values of length 2
• scaleFactor - Optional, I16, 16-bit signless integer

Results

• dst - Single, anonymous/composite constraint, vector of bfloat16 type values of length 2

Description

This Op converts the given f8 inputs in a i8x2 vector to bf16.

The result dst is represented as a vector of bf16 elements.

The relu attribute, when set, lowers to the '.relu' variant of the cvt instruction.

The sat attribute specifies the saturation mode.

The optional scaling-factors for each of the inputs are provided through the operand scaleFactor as a packed i16 type. Only ue8m0 is supported as the type of the scale-factor currently.

For more information, see PTX ISA

Example:

// Basic conversion from f8E4M3FN.
%res1 = nvvm.convert.f8x2.to.bf16x2 %src
    : vector<2xi8> (f8E4M3FN) -> vector<2xbf16>

// Conversion from f8E5M2 with relu and saturation.
%res2 = nvvm.convert.f8x2.to.bf16x2 %src
    {relu = true, sat = #nvvm.sat_mode<satfinite>}
    : vector<2xi8> (f8E5M2) -> vector<2xbf16>

// Conversion with a packed ue8m0 scale-factor.
%res3 = nvvm.convert.f8x2.to.bf16x2 %src, %scaleFactor
    : vector<2xi8> (f8E4M3FN) -> vector<2xbf16>

Return op name nvvm.convert.f8x2.to.f16x2 as a bitstring.

nvvm.convert.f8x2.to.f16x2 - Convert a pair of f8 inputs to f16x2

Attributes

• srcType - Single, anonymous/composite constraint, type attribute of f8E4M3FN type or f8E5M2 type
• relu - Single, BoolAttr, bool attribute

Operands

• src - Single, anonymous/composite constraint, vector of 8-bit signless integer values of length 2

Results

• dst - Single, anonymous/composite constraint, vector of 16-bit float values of length 2

Description

This Op converts the given f8 inputs in a i8x2 vector to f16.

The result dst is represented as a vector of f16 elements.

The relu attribute, when set, lowers to the '.relu' variant of the cvt instruction.

For more information, see PTX ISA

Return op name nvvm.convert.f16x2.to.f4x2 as a bitstring.

nvvm.convert.f16x2.to.f4x2 - Convert an f16x2 input to f4x2

This op has support for result type inference.

Attributes

• relu - Single, BoolAttr, bool attribute
• dstTy - Single, anonymous/composite constraint, type attribute of f4E2M1FN type

Operands

• src - Single, anonymous/composite constraint, vector of 16-bit float values of length 2

Results

• dst - Single, I8, 8-bit signless integer

Description

This Op converts each of the given F16 inputs in an f16x2 vector to the specified fp4 type. The result dst is returned as an i8 type where the converted values are packed such that the value converted from the first element of a is stored in the lower 4 bits of dst and the value converted from the second element of a is stored in the upper 4 bits of dst. The relu attribute, when set, lowers to the '.relu' variant of the cvt instruction.

Return op name nvvm.convert.f16x2.to.f6x2 as a bitstring.

nvvm.convert.f16x2.to.f6x2 - Convert an f16x2 input to f6x2

Attributes

• relu - Single, BoolAttr, bool attribute
• dstTy - Single, anonymous/composite constraint, type attribute of f6E2M3FN type or f6E3M2FN type

Operands

• src - Single, anonymous/composite constraint, vector of 16-bit float values of length 2

Results

• dst - Single, anonymous/composite constraint, 16-bit signless integer or vector of 8-bit signless integer values of length 2

Description

This Op converts each of the given F16 inputs in an f16x2 vector to the specified fp6 type. The result dst is represented either as an i16 type or as a vector of two i8 types. If dst is returned as an i16 type, the converted values are packed such that the value converted from the first element of a is stored in the lower 8 bits of dst with 2 MSB bits padded with zeros and the value converted from the second element of a is stored in the upper 8 bits of dst with 2 MSB bits padded with zeros. If dst is returned as a vector type, each converted value is stored as an i8 element in the vector with 2 MSB bits padded with zeros. The relu attribute, when set, lowers to the '.relu' variant of the cvt instruction.

Return op name nvvm.convert.f16x2.to.f8x2 as a bitstring.

nvvm.convert.f16x2.to.f8x2 - Convert an f16x2 input to f8x2

Attributes

• relu - Single, BoolAttr, bool attribute
• dstTy - Single, TypeAttr, any type attribute

Operands

• a - Single, anonymous/composite constraint, vector of 16-bit float values of length 2

Results

• dst - Single, anonymous/composite constraint, 16-bit signless integer or vector of 8-bit signless integer values of length 2

Description

This Op converts the given f16 inputs in an f16x2 vector to the specified f8 type. The result dst is represented as an i16 type or as a vector of two i8 types. If dst is returned as an i16 type, the converted values from a are packed such that the value converted from the first element of a is stored in the upper 8 bits of dst and the value converted from the second element of a is stored in the lower 8 bits of dst. If dst is returned as a vector type, each converted value is stored as an i8 element in the vector. The relu attribute, when set, lowers to the '.relu' variant of the cvt instruction.

For more information, see PTX ISA

Return op name nvvm.convert.f32x2.to.bf16x2 as a bitstring.

nvvm.convert.f32x2.to.bf16x2 - Convert two F32 values to packed bf16x2.

Attributes

• rnd - Single, FPRoundingModeAttr, NVVM FPRoundingMode kind
• sat - Single, SaturationModeAttr, Describes the saturation mode
• relu - Single, BoolAttr, bool attribute

Operands

• src_hi - Single, F32, 32-bit float
• src_lo - Single, F32, 32-bit float
• random_bits - Optional, I32, 32-bit signless integer

Results

• dst - Single, anonymous/composite constraint, vector of bfloat16 type values of length 2

Description

Converts two F32 values to packed bf16x2 format with the specified rounding mode. The src_hi and src_lo parameters correspond to operands a and b in the PTX ISA, respectively.

The random_bits parameter is required for stochastic rounding and provides the random bits to be used for the conversion.

The relu attribute clamps negative results to 0.

The sat attribute determines saturation behavior.

For more information, see PTX ISA

Return op name nvvm.convert.f32x2.to.f4x2 as a bitstring.

nvvm.convert.f32x2.to.f4x2 - Convert a pair of float inputs to f4x2

This op has support for result type inference.

Attributes

• relu - Single, BoolAttr, bool attribute
• dstTy - Single, TypeAttr, any type attribute

Operands

• a - Single, F32, 32-bit float
• b - Single, F32, 32-bit float

Results

• dst - Single, I8, 8-bit signless integer

Description

This Op converts each of the given float inputs to the specified fp4 type. The result dst is returned as an i8 type where the converted values are packed such that the value converted from a is stored in the upper 4 bits of dst and the value converted from b is stored in the lower 4 bits of dst. The relu attribute, when set, lowers to the '.relu' variant of the cvt instruction.

For more information, see PTX ISA

Return op name nvvm.convert.f32x2.to.f6x2 as a bitstring.

nvvm.convert.f32x2.to.f6x2 - Convert a pair of float inputs to f6x2

Attributes

• relu - Single, BoolAttr, bool attribute
• dstTy - Single, TypeAttr, any type attribute

Operands

• a - Single, F32, 32-bit float
• b - Single, F32, 32-bit float

Results

• dst - Single, anonymous/composite constraint, 16-bit signless integer or vector of 8-bit signless integer values of length 2

Description

This Op converts each of the given float inputs to the specified fp6 type. The result dst is represented either as an i16 type or as a vector of two i8 types. If dst is returned as an i16 type, the converted values are packed such that the value converted from a is stored in the upper 8 bits of dst with 2 MSB bits padded with zeros and the value converted from b is stored in the lower 8 bits of dst with 2 MSB bits padded with zeros. If dst is returned as a vector type, each converted value is stored as an i8 element in the vector. The relu attribute, when set, lowers to the '.relu' variant of the cvt instruction.

For more information, see PTX ISA

Return op name nvvm.convert.f32x2.to.f8x2 as a bitstring.

nvvm.convert.f32x2.to.f8x2 - Convert a pair of float inputs to f8x2

Attributes

• rnd - Single, FPRoundingModeAttr, NVVM FPRoundingMode kind
• sat - Single, SaturationModeAttr, Describes the saturation mode
• relu - Single, BoolAttr, bool attribute
• dstTy - Single, TypeAttr, any type attribute

Operands

• a - Single, F32, 32-bit float
• b - Single, F32, 32-bit float

Results

• dst - Single, anonymous/composite constraint, 16-bit signless integer or vector of 8-bit signless integer values of length 2

Description

This Op converts each of the given float inputs to the specified fp8 type. The result dst is represented as an i16 type or as a vector of two i8 types. If dst is returned as an i16 type, the converted values are packed such that the value converted from a is stored in the upper 8 bits of dst and the value converted from b is stored in the lower 8 bits of dst. If dst is returned as a vector type, each converted value is stored as an i8 element in the vector. The rnd and sat attributes specify the rounding and saturation modes respectively. The relu attribute, when set, lowers to the '.relu' variant of the cvt instruction.

For more information, see PTX ISA

Return op name nvvm.convert.f32x2.to.f16x2 as a bitstring.

nvvm.convert.f32x2.to.f16x2 - Convert two F32 values to packed f16x2.

Attributes

• rnd - Single, FPRoundingModeAttr, NVVM FPRoundingMode kind
• sat - Single, SaturationModeAttr, Describes the saturation mode
• relu - Single, BoolAttr, bool attribute

Operands

• src_hi - Single, F32, 32-bit float
• src_lo - Single, F32, 32-bit float
• random_bits - Optional, I32, 32-bit signless integer

Results

• dst - Single, anonymous/composite constraint, vector of 16-bit float values of length 2

Description

Converts two F32 values to packed f16x2 format with the specified rounding mode. The src_hi and src_lo parameters correspond to operands a and b in the PTX ISA, respectively.

The random_bits parameter is required for stochastic rounding and provides the random bits to be used for the conversion.

The relu attribute clamps negative results to 0.

The sat attribute determines saturation behavior.

For more information, see PTX ISA

Return op name nvvm.convert.f32x2.to.s2f6x2 as a bitstring.

nvvm.convert.f32x2.to.s2f6x2 - Convert a pair of f32 inputs to S2F6x2

Attributes

• relu - Single, BoolAttr, bool attribute

Operands

• a - Single, F32, 32-bit float
• b - Single, F32, 32-bit float
• scaleFactor - Optional, I16, 16-bit signless integer

Results

• dst - Single, anonymous/composite constraint, 16-bit signless integer or vector of 8-bit signless integer values of length 2

Description

This Op converts each of the given f32 inputs to the S2F6x2 type. The result dst can be either a packed i16 type or a vector of two i8 types. If dst is returned as an i16 type, the converted values are packed such that the value converted from a is stored in the upper 8 bits of dst and the value converted from b is stored in the lower 8 bits of dst. If dst is returned as a vector type, each converted value is stored as an i8 element in the vector. The relu attribute, when set, lowers to the '.relu' variant of the cvt instruction. The optional scaling-factors for each of the inputs are provided through the operand scaleFactor as a packed i16 type. Only ue8m0 is supported as the type of the scale-factor currently.

For more information, see PTX ISA

Return op name nvvm.convert.f32x4.to.f4x4 as a bitstring.

nvvm.convert.f32x4.to.f4x4 - Convert vector<4xf32> to packed f4x4 with stochastic rounding (.rs) and satfinite

This op has support for result type inference.

Attributes

• relu - Single, BoolAttr, bool attribute
• dstTy - Single, TypeAttr, any type attribute

Operands

• src - Single, anonymous/composite constraint, vector of 32-bit float values of length 4
• rbits - Single, I32, 32-bit signless integer

Results

• dst - Single, I16, 16-bit signless integer

Description

Converts a vector<4xf32> to packed f4x4 format using stochastic rounding (.rs) mode with SATFINITE saturation. Randomness is provided by the rbits parameter. The dstTy attribute specifies the target floating-point format. The relu attribute clamps negative results to 0.

Note: These operations always use RS rounding mode and SATFINITE saturation mode.

For more information, see PTX ISA

Return op name nvvm.convert.f32x4.to.f6x4 as a bitstring.

nvvm.convert.f32x4.to.f6x4 - Convert vector<4xf32> to packed f6x4 with stochastic rounding (.rs) and satfinite

Attributes

• relu - Single, BoolAttr, bool attribute
• dstTy - Single, TypeAttr, any type attribute

Operands

• src - Single, anonymous/composite constraint, vector of 32-bit float values of length 4
• rbits - Single, I32, 32-bit signless integer

Results

• dst - Single, anonymous/composite constraint, vector of 8-bit signless integer values of length 4

Description

Converts a vector<4xf32> to packed f6x4 format using stochastic rounding (.rs) mode with SATFINITE saturation. Randomness is provided by the rbits parameter. The dstTy attribute specifies the target floating-point format. The relu attribute clamps negative results to 0.

Note: These operations always use RS rounding mode and SATFINITE saturation mode.

For more information, see PTX ISA

Return op name nvvm.convert.f32x4.to.f8x4 as a bitstring.

nvvm.convert.f32x4.to.f8x4 - Convert vector<4xf32> to packed f8x4 with stochastic rounding (.rs) and satfinite

Attributes

• relu - Single, BoolAttr, bool attribute
• dstTy - Single, TypeAttr, any type attribute

Operands

• src - Single, anonymous/composite constraint, vector of 32-bit float values of length 4
• rbits - Single, I32, 32-bit signless integer

Results

• dst - Single, anonymous/composite constraint, vector of 8-bit signless integer values of length 4

Description

Converts a vector<4xf32> to packed f8x4 format using stochastic rounding (.rs) mode with SATFINITE saturation. Randomness is provided by the rbits parameter. The dstTy attribute specifies the target floating-point format. The relu attribute clamps negative results to 0.

Note: These operations always use RS rounding mode and SATFINITE saturation mode.

For more information, see PTX ISA

Return op name nvvm.convert.float.to.tf32 as a bitstring.

nvvm.convert.float.to.tf32 - Convert the given float input to TF32

This op has support for result type inference.

Attributes

• rnd - Single, FPRoundingModeAttr, NVVM FPRoundingMode kind
• sat - Single, SaturationModeAttr, Describes the saturation mode
• relu - Single, BoolAttr, bool attribute

Operands

• src - Single, F32, 32-bit float

Results

• res - Single, I32, 32-bit signless integer

Description

This Op converts the given f32 input to tf32. The result res is represented as an i32 type. The relu attribute, when set, lowers to the '.relu' variant of the cvt instruction. The rnd and sat attributes specify the the rounding and saturation modes respectively.

For more information, see PTX ISA

Return op name nvvm.convert.s2f6x2.to.bf16x2 as a bitstring.

nvvm.convert.s2f6x2.to.bf16x2 - Convert s2f6x2 to bf16x2

Attributes

• sat - Single, SaturationModeAttr, Describes the saturation mode
• relu - Single, BoolAttr, bool attribute

Operands

• src - Single, anonymous/composite constraint, vector of 8-bit signless integer values of length 2
• scaleFactor - Optional, I16, 16-bit signless integer

Results

• dst - Single, anonymous/composite constraint, vector of bfloat16 type values of length 2

Description

This Op converts a pair of s2f6x2 inputs to bf16x2 type. The result dst is represented as a vector of two bf16 elements.

The relu attribute, when set, lowers to the '.relu' variant of the cvt instruction.

The optional scaling-factors for each of the inputs are provided through the operand scaleFactor as a packed i16 type. Only ue8m0 is supported as the type of the scale-factor currently.

For more information, see PTX ISA

Return op name nvvm.cos as a bitstring.

nvvm.cos - Cosine (fast approximation)

This op has support for result type inference.

Attributes

• ftz - Single, BoolAttr, bool attribute

Operands

• src - Single, F32, 32-bit float

Results

• res - Single, F32, 32-bit float

Description

Computes a fast approximation of the cosine of the input value (in radians). The ftz attribute, when set, flushes subnormal inputs and results to sign-preserving zero.

Return op name nvvm.cp.async.bulk.commit.group as a bitstring.

nvvm.cp.async.bulk.commit.group

Description

This Op commits all prior initiated but uncommitted cp.async.bulk instructions into a cp.async.bulk-group.

For more information, see PTX ISA

Return op name nvvm.cp.async.bulk.global.shared.cta as a bitstring.

nvvm.cp.async.bulk.global.shared.cta - Async bulk copy from Shared CTA memory to Global memory

Operands

• dstMem - Single, LLVM_PointerGlobal, LLVM pointer in address space 1
• srcMem - Single, LLVM_PointerShared, LLVM pointer in address space 3
• size - Single, I32, 32-bit signless integer
• l2CacheHint - Optional, I64, 64-bit signless integer
• byteMask - Optional, I16, 16-bit signless integer

Description

Initiates an asynchronous copy operation from Shared CTA memory to global memory. The 32-bit operand size specifies the amount of memory to be copied, in terms of number of bytes. size must be a multiple of 16. The l2CacheHint operand is optional, and it is used to specify cache eviction policy that may be used during the memory access. The byteMask operand is optional. The i-th bit in the 16-bit wide byteMask specifies whether the i-th byte of each 16-byte wide chunk of source data is copied to the destination. If the bit is set, the byte is copied.

Example:

  nvvm.cp.async.bulk.global.shared.cta %dst, %src, %size
      : !llvm.ptr<1>, !llvm.ptr<3>

  // with l2_cache_hint
  nvvm.cp.async.bulk.global.shared.cta %dst, %src, %size l2_cache_hint = %ch
      : !llvm.ptr<1>, !llvm.ptr<3>

  // with byte_mask
  nvvm.cp.async.bulk.global.shared.cta %dst, %src, %size byte_mask = %mask
      : !llvm.ptr<1>, !llvm.ptr<3>

  // with both l2_cache_hint and byte_mask
  nvvm.cp.async.bulk.global.shared.cta %dst, %src, %size l2_cache_hint = %ch byte_mask = %mask
      : !llvm.ptr<1>, !llvm.ptr<3>

For more information, see PTX ISA

Return op name nvvm.cp.async.bulk.prefetch as a bitstring.

nvvm.cp.async.bulk.prefetch - Async bulk prefetch from global memory to L2 cache

Operands

• srcMem - Single, LLVM_PointerGlobal, LLVM pointer in address space 1
• size - Single, I32, 32-bit signless integer
• l2CacheHint - Optional, I64, 64-bit signless integer

Description

Initiates an asynchronous prefetch of data from the location specified by srcMem to the L2 cache.

The l2CacheHint operand is optional, and it is used to specify cache eviction policy that may be used during the memory access.

Example:

  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>

For more information, see PTX ISA

Return op name nvvm.cp.async.bulk.shared.cluster.global as a bitstring.

nvvm.cp.async.bulk.shared.cluster.global - Async bulk copy from global to Shared {cta or cluster} memory

Operands

• dstMem - Single, anonymous/composite constraint, LLVM pointer in address space 3 or LLVM pointer in address space 7
• srcMem - Single, LLVM_PointerGlobal, LLVM pointer in address space 1
• mbar - Single, LLVM_PointerShared, LLVM pointer in address space 3
• size - Single, I32, 32-bit signless integer
• multicastMask - Optional, I16, 16-bit signless integer
• l2CacheHint - Optional, I64, 64-bit signless integer

Description

Initiates an asynchronous copy operation from global memory to shared memory or shared_cluster memory.

The multicastMask operand is optional and can be used only when the destination is shared::cluster memory. When it is present, this Op copies data from global memory to shared memory of multiple CTAs in the cluster. Operand multicastMask specifies the destination CTAs in the cluster such that each bit position in the 16-bit multicastMask operand corresponds to the nvvm.read.ptx.sreg.ctaid of the destination CTA.

The l2CacheHint operand is optional, and it is used to specify cache eviction policy that may be used during the memory access.

For more information, see PTX ISA

Return op name nvvm.cp.async.bulk.shared.cluster.shared.cta as a bitstring.

nvvm.cp.async.bulk.shared.cluster.shared.cta - Async bulk copy from Shared CTA memory to Shared cluster memory

Operands

• dstMem - Single, LLVM_PointerSharedCluster, LLVM pointer in address space 7
• srcMem - Single, LLVM_PointerShared, LLVM pointer in address space 3
• mbar - Single, LLVM_PointerShared, LLVM pointer in address space 3
• size - Single, I32, 32-bit signless integer

Description

Initiates an asynchronous copy operation from Shared CTA memory to Shared cluster memory.

For more information, see PTX ISA

Return op name nvvm.cp.async.bulk.tensor.global.shared.cta as a bitstring.

nvvm.cp.async.bulk.tensor.global.shared.cta

Attributes

• mode - Single, TMAStoreModeAttr, NVVM TMA Store Mode

Operands

• tmaDescriptor - Single, LLVM_PointerGeneric, LLVM pointer in address space 0
• srcMem - Single, LLVM_PointerShared, LLVM pointer in address space 3
• coordinates - Variadic, I32, variadic of 32-bit signless integer
• l2CacheHint - Optional, I64, 64-bit signless integer
• predicate - Optional, PtxPredicate, 1-bit signless integer

Description

Initiates an asynchronous copy of the tensor data from shared::cta memory to global memory. This Op supports all the store modes specified in TMAStoreMode.

The l2CacheHint operand is optional, and it is used to specify cache eviction policy that may be used during the memory access.

For more information, see PTX ISA

Return op name nvvm.cp.async.bulk.tensor.prefetch as a bitstring.

nvvm.cp.async.bulk.tensor.prefetch

Attributes

• mode - Single, TMALoadModeAttr, List of Load-Modes supported for TMA Tensor Ops

Operands

• tmaDescriptor - Single, LLVM_PointerGeneric, LLVM pointer in address space 0
• coordinates - Variadic, I32, variadic of 32-bit signless integer
• im2colOffsets - Variadic, I16, variadic of 16-bit signless integer
• l2CacheHint - Optional, I64, 64-bit signless integer

Description

Initiates an asynchronous prefetch operation on the tensor data from global memory to L2 cache. This Op supports all the load modes specified in TMALoadMode.

The l2CacheHint operand is optional, and it is used to specify cache eviction policy that may be used during the memory access.

For more information, see PTX ISA

Return op name nvvm.cp.async.bulk.tensor.reduce as a bitstring.

nvvm.cp.async.bulk.tensor.reduce

Attributes

• redKind - Single, TMAReduxKindAttr, NVVM TMA redux kind
• mode - Single, TMAStoreModeAttr, NVVM TMA Store Mode

Operands

• tmaDescriptor - Single, LLVM_AnyPointer, LLVM pointer type
• srcMem - Single, LLVM_PointerShared, LLVM pointer in address space 3
• coordinates - Variadic, I32, variadic of 32-bit signless integer
• l2CacheHint - Optional, I64, 64-bit signless integer

Description

Initiates an asynchronous reduction operation of tensor data in global memory with tensor data in shared memory.

The mode attribute indicates whether the copy mode is tile or im2col. The redOp attribute specifies the reduction operations applied. The supported reduction operations are: {add, min, max, inc, dec, and, or, xor}

The l2CacheHint operand is optional, and it is used to specify cache eviction policy that may be used during the memory access.

For more information, see PTX ISA

Return op name nvvm.cp.async.bulk.tensor.shared.cluster.global as a bitstring.

nvvm.cp.async.bulk.tensor.shared.cluster.global

Attributes

• mode - Single, TMALoadModeAttr, List of Load-Modes supported for TMA Tensor Ops
• isCTAOnly - Single, BoolAttr, bool attribute
• group - Optional, CTAGroupKindAttr, NVVM CTA group kind

Operands

• dstMem - Single, anonymous/composite constraint, LLVM pointer in address space 3 or LLVM pointer in address space 7
• tmaDescriptor - Single, LLVM_PointerGeneric, LLVM pointer in address space 0
• coordinates - Variadic, I32, variadic of 32-bit signless integer
• mbar - Single, LLVM_PointerShared, LLVM pointer in address space 3
• im2colOffsets - Variadic, I16, variadic of 16-bit signless integer
• multicastMask - Optional, I16, 16-bit signless integer
• l2CacheHint - Optional, I64, 64-bit signless integer
• predicate - Optional, PtxPredicate, 1-bit signless integer

Description

Initiates an asynchronous copy operation on the tensor data from global memory to shared::cluster (or) shared::cta memory. This Op supports all the load modes specified in TMALoadMode.

The multicastMask operand is optional. When it is present, the Op copies data from global memory to shared memory of multiple CTAs in the cluster. Operand multicastMask specifies the destination CTAs in the cluster such that each bit position in the 16-bit multicastMask operand corresponds to the nvvm.read.ptx.sreg.ctaid of the destination CTA.

The l2CacheHint operand is optional, and it is used to specify cache eviction policy that may be used during the memory access.

When the isCTAOnly attribute is set to true, the destination is shared::cta only. Hence, multicastMask and CTAGroup are not applicable when isCTAOnly is true.

For more information, see PTX ISA

Return op name nvvm.cp.async.bulk.wait_group as a bitstring.

nvvm.cp.async.bulk.wait_group

Attributes

• group - Single, I32Attr, 32-bit signless integer attribute whose minimum value is 0
• read - Optional, UnitAttr, unit attribute

Description

Op waits for completion of the most recent bulk async-groups.

The $group operand tells waiting has to be done until for $group or fewer of the most recent bulk async-groups. If `$group` is 0, the op wait until all the most recent bulk async-groups have completed.

The $read indicates that the waiting has to be done until all the bulk async operations in the specified bulk async-group have completed reading from their source locations.

For more information, see PTX ISA

Return op name nvvm.cp.async.commit.group as a bitstring.

nvvm.cp.async.commit.group

Return op name nvvm.cp.async.mbarrier.arrive as a bitstring.

nvvm.cp.async.mbarrier.arrive - NVVM Dialect Op for cp.async.mbarrier.arrive

Attributes

• noinc - Single, I1Attr, 1-bit signless integer attribute

Operands

• addr - Single, anonymous/composite constraint, LLVM pointer in address space 0 or LLVM pointer in address space 3

Description

The cp.async.mbarrier.arrive Op makes the mbarrier object track all prior cp.async operations initiated by the executing thread. The addr operand specifies the address of the mbarrier object in generic or shared::cta address space. When it is generic, the underlying memory should fall within the shared::cta space; otherwise the behavior is undefined. The noinc attr impacts how the mbarrier's state is updated.

For more information, see PTX ISA

Return op name nvvm.cp.async.shared.global as a bitstring.

nvvm.cp.async.shared.global

Attributes

• size - Single, I32Attr, 32-bit signless integer attribute
• modifier - Single, LoadCacheModifierAttr, NVVM load cache modifier kind

Operands

• dst - Single, LLVM_PointerShared, LLVM pointer in address space 3
• src - Single, LLVM_PointerGlobal, LLVM pointer in address space 1
• cpSize - Optional, I32, 32-bit signless integer

Return op name nvvm.cp.async.wait.group as a bitstring.

nvvm.cp.async.wait.group

Attributes

• n - Single, I32Attr, 32-bit signless integer attribute

Return op name nvvm.divf as a bitstring.

nvvm.divf - Divide one value by another

This op has support for result type inference.

Attributes

• rnd - Single, FPArithRoundingMode, NVVM FPRoundingMode kind whose value is one of {none, rm, rn, rp, rz}
• ftz - Single, BoolAttr, bool attribute
• approx - Single, BoolAttr, bool attribute
• full - Single, BoolAttr, bool attribute

Operands

• lhs - Single, anonymous/composite constraint, 32-bit float or 64-bit float
• rhs - Single, anonymous/composite constraint, 32-bit float or 64-bit float

Results

• res - Single, anonymous/composite constraint, 32-bit float or 64-bit float

Description

Divides lhs by rhs, stores result in res (res = lhs / rhs).

For more information, see PTX ISA

Return op name nvvm.dot.accumulate.2way as a bitstring.

nvvm.dot.accumulate.2way - Two-way 16-bit to 8-bit dot product-accumulate instruction

This op has support for result type inference.

Attributes

• a_type - Single, DotAccumulateTypeAttr, NVVM DotAccumulateType
• b_type - Single, DotAccumulateTypeAttr, NVVM DotAccumulateType
• b_hi - Single, BoolAttr, bool attribute

Operands

• a - Single, anonymous/composite constraint, vector of 16-bit signless integer values of length 2
• b - Single, anonymous/composite constraint, vector of 8-bit signless integer values of length 4
• c - Single, I32, 32-bit signless integer

Results

• res - Single, I32, 32-bit signless integer

Description

Performs a two-way 16-bit to 8-bit dot-product which is accumulated in a 32-bit result. Operand a is a vector of two 16-bit elements and operand b a vector of four 8-bit elements between which the dot product is computed.

The a_type and b_type attributes specify the type of the elements in a and b respectively. If a_type or b_type is s, then the elements in the corresponding vector are sign-extended to 32-bit before the dot product is computed. If a_type or b_type is u, then the elements in the corresponding vector are zero-extended to 32-bit instead.

The b_hi boolean attribute specifies which two bytes of b are used for the dot product. If b_hi is true, then the dot product is computed between a and elements at indices 2 and 3 of b. If b_hi is false, then the dot product is computed between a and elements at indices 0 and 1 of b.

Operand c is a 32-bit integer to which the result is accumulated. It is treated as holding a signed integer if any of a_type or b_type is signed.

For more information, see PTX ISA

Return op name nvvm.dot.accumulate.4way as a bitstring.

nvvm.dot.accumulate.4way - Four-way byte dot product-accumulate instruction

This op has support for result type inference.

Attributes

• a_type - Single, DotAccumulateTypeAttr, NVVM DotAccumulateType
• b_type - Single, DotAccumulateTypeAttr, NVVM DotAccumulateType

Operands

• a - Single, anonymous/composite constraint, vector of 8-bit signless integer values of length 4
• b - Single, anonymous/composite constraint, vector of 8-bit signless integer values of length 4
• c - Single, I32, 32-bit signless integer

Results

• res - Single, I32, 32-bit signless integer

Description

Performs a four-way byte dot-product which is accumulated in a 32-bit result. Operand a and b are vectors of 4 bytes between which the dot product is computed.

The a_type and b_type attributes specify the type of the elements in a and b respectively. If a_type or b_type is signed, then the elements in the corresponding vector are sign-extended to 32-bit before the dot product is computed. If a_type or b_type is unsigned, then the elements in the corresponding vector are zero-extended to 32-bit instead.

Operand c is a 32-bit integer to which the result is accumulated. It is treated as holding a signed integer if any of a_type or b_type is s8.

For more information, see PTX ISA

Return op name nvvm.elect.sync as a bitstring.

nvvm.elect.sync - Elect one leader thread

This op has support for result type inference.

Operands

• membermask - Optional, I32, 32-bit signless integer

Results

• pred - Single, I1, 1-bit signless integer

Description

The elect.sync instruction elects one predicated active leader thread from among a set of threads specified in the membermask. When the membermask is not provided explicitly, a default value of 0xFFFFFFFF is used. The predicate result is set to True for the leader thread, and False for all other threads.

For more information, see PTX ISA

Return op name nvvm.ex2 as a bitstring.

nvvm.ex2 - Base-2 exponential (fast approximation)

This op has support for result type inference.

Attributes

• ftz - Single, BoolAttr, bool attribute

Operands

• src - Single, F32, 32-bit float

Results

• res - Single, F32, 32-bit float

Description

Computes a fast approximation of 2 raised to the power of the input value. The ftz attribute, when set, flushes subnormal inputs and results to sign-preserving zero.

Return op name nvvm.exit as a bitstring.

nvvm.exit - Exit Op

Description

Ends execution of a thread. For more information, see PTX ISA

Return op name nvvm.fence.mbarrier.init as a bitstring.

nvvm.fence.mbarrier.init

Description

Fence operation that applies on the prior nvvm.mbarrier.init

For more information, see PTX ISA

Return op name nvvm.fence.proxy as a bitstring.

nvvm.fence.proxy

Attributes

• kind - Single, ProxyKindAttr, Proxy kind whose value is none of {tensormap, generic}
• space - Optional, SharedSpaceAttr, Shared memory space

Description

Fence operation with proxy to establish an ordering between memory accesses that may happen through different proxies.

For more information, see PTX ISA

Return op name nvvm.fence.proxy.acquire as a bitstring.

nvvm.fence.proxy.acquire - Uni-directional proxy fence operation with acquire semantics

Attributes

• scope - Single, MemScopeKindAttr, NVVM Memory Scope kind
• fromProxy - Single, ProxyKindAttr, Proxy kind
• toProxy - Single, ProxyKindAttr, Proxy kind

Operands

• addr - Single, LLVM_PointerGeneric, LLVM pointer in address space 0
• size - Single, I32, 32-bit signless integer

Description

fence.proxy.acquire is a uni-directional fence used to establish ordering between a prior memory access performed via the generic proxy and a subsequent memory access performed via the tensormap proxy

The address operand addr and the operand size together specify the memory range [addr, addr+size) on which the ordering guarantees on the memory accesses across the proxies is to be provided. The only supported value for the size operand is 128 and must be an immediate. Generic Addressing is used unconditionally, and the address specified by the operand addr must fall within the .global state space. Otherwise, the behavior is undefined

For more information, see PTX ISA

Return op name nvvm.fence.proxy.release as a bitstring.

nvvm.fence.proxy.release - Uni-directional proxy fence operation with release semantics

Attributes

• scope - Single, MemScopeKindAttr, NVVM Memory Scope kind
• fromProxy - Single, ProxyKindAttr, Proxy kind
• toProxy - Single, ProxyKindAttr, Proxy kind

Description

fence.proxy.release is a uni-directional fence used to establish ordering between a prior memory access performed via the generic proxy and a subsequent memory access performed via the tensormap proxy. fence.proxy.release operation can form a release sequence that synchronizes with an acquire sequence that contains the fence.proxy.acquire proxy fence operation

For more information, see PTX ISA

Return op name nvvm.fence.proxy.sync_restrict as a bitstring.

nvvm.fence.proxy.sync_restrict - Uni-directional proxy fence operation with sync_restrict

Attributes

• order - Single, MemOrderKindAttr, NVVM Memory Ordering kind whose value is one of {acquire, release}
• fromProxy - Single, ProxyKindAttr, Proxy kind
• toProxy - Single, ProxyKindAttr, Proxy kind

Description

The nvvm.fence.proxy.sync_restrict Op used to establish ordering between a prior memory access performed between proxies. Currently, the ordering is only supported between async and generic proxies. sync_restrict restricts acquire memory semantics to shared_cluster and release memory semantics to shared_cta with cluster scope. For more information, see PTX ISA

Return op name nvvm.fence.sc.cluster as a bitstring.

nvvm.fence.sc.cluster

Return op name nvvm.fence.sync_restrict as a bitstring.

nvvm.fence.sync_restrict - Uni-directional thread fence operation

Attributes

• order - Single, MemOrderKindAttr, NVVM Memory Ordering kind whose value is one of {acquire, release}

Description

The nvvm.fence.sync_restrict Op restricts the class of memory operations for which the fence instruction provides the memory ordering guarantees. sync_restrict restricts acquire memory semantics to shared_cluster and release memory semantics to shared_cta with cluster scope. For more information, see PTX ISA

Return op name nvvm.fma as a bitstring.

nvvm.fma -

Performs floating point fused multiply-add operation with support for mixed 
precision operands

This op has support for result type inference.

Attributes

• rnd - Single, FPArithRoundingMode, NVVM FPRoundingMode kind whose value is one of {none, rm, rn, rp, rz}
• sat - Single, SaturationModeSatOrNone, Describes the saturation mode whose value is one of {none, sat}
• ftz - Single, BoolAttr, bool attribute
• relu - Single, BoolAttr, bool attribute
• oob - Single, BoolAttr, bool attribute

Operands

• a - Single, SIMTFloatType, 16-bit float or bfloat16 type or 32-bit float or 64-bit float or vector of 16-bit float or bfloat16 type or 32-bit float or 64-bit float values of length 2
• b - Single, SIMTFloatType, 16-bit float or bfloat16 type or 32-bit float or 64-bit float or vector of 16-bit float or bfloat16 type or 32-bit float or 64-bit float values of length 2
• c - Single, SIMTFloatType, 16-bit float or bfloat16 type or 32-bit float or 64-bit float or vector of 16-bit float or bfloat16 type or 32-bit float or 64-bit float values of length 2

Results

• res - Single, SIMTFloatType, 16-bit float or bfloat16 type or 32-bit float or 64-bit float or vector of 16-bit float or bfloat16 type or 32-bit float or 64-bit float values of length 2

Description

The nvvm.fma operation performs floating point fused multiply-add of three operands of the same type.

The rounding mode is specified by the rnd attribute, saturation mode by the sat attribute, flush-to-zero by the ftz attribute, and ReLU by the relu attribute.

Out-of-bounds (OOB) behavior is controlled by the oob attribute. oob clamps the result to 0 if either of the operands is OOB NaN (see Tensors).

For more information, see PTX ISA:

• floating point fused multiply-add
• half-precision floating point fused multiply-add

Return op name nvvm.griddepcontrol as a bitstring.

nvvm.griddepcontrol

Attributes

• kind - Single, GridDepActionAttr, Action kind for grid dependency control

Description

If the $kind attribute is set to wait, it causes the executing thread to wait until all prerequisite grids in flight have completed and all the memory operations from the prerequisite grids are performed and made visible to the current grid.

When the $kind is launch_dependents, it signals that specific dependents the runtime system designated to react to this instruction can be scheduled as soon as all other CTAs in the grid issue the same instruction or have completed.

For more information, see PTX ISA

Return op name nvvm.inline_ptx as a bitstring.

nvvm.inline_ptx - Inline PTX Op

Attributes

• ptxCode - Single, StrAttr, string attribute
• memoryClobber - Single, BoolAttr, bool attribute

Operands

• readOnlyArgs - Variadic, AnyType, variadic of any non-token type
• readWriteArgs - Variadic, AnyType, variadic of any non-token type
• predicate - Optional, PtxPredicate, 1-bit signless integer

Results

• writeOnlyArgs - Variadic, AnyType, variadic of any non-token type

Description

This op allows using PTX directly within the NVVM

dialect, while greatly simplifying llvm.inline_asm generation. It 
automatically handles register size selection and sets the correct 
read/write access for each operand. The operation leverages the 
`BasicPtxBuilderInterface` to abstract away low-level details of 
PTX assembly formatting.

The `predicate` attribute is used to specify a predicate for the
PTX instruction.

The `memory_clobber` attribute appends a "~{memory}" clobber to the
constraints of the generated inline assembly. Set it when the PTX reads
or writes memory beyond its listed operands (e.g. stores, atomics, or
instructions with acquire/release semantics such as mbarrier), so that
LLVM does not reorder memory accesses across the inline assembly.

Example 1: Read-only Parameters
```mlir
nvvm.inline_ptx "mbarrier.init.b64 [$0], $1;" ro(%barrier_gen, %count : !llvm.ptr, i32) memory_clobber = true

// Lowers to:
llvm.inline_asm has_side_effects asm_dialect = att
  "mbarrier.init.b64 [$0], $1;", "l,r,~{memory}" %arg0, %arg2 : (!llvm.ptr, i32) -> ()
```

Example 2: Read-only and Write-only Parameters
```mlir
%0 = nvvm.inline_ptx "ex2.approx.ftz.f32 $0, $1;" (%input) : f32 -> f32

// Lowers to:
%0 = llvm.inline_asm has_side_effects asm_dialect = att 
  "ex2.approx.ftz.f32 $0, $1;", "=f,f" %arg0 : (f32) -> f32
```

Example 3: Predicate Usage
```mlir
nvvm.inline_ptx "mbarrier.init.b64 [$0], $1;" (%barrier_gen, %count), 
  predicate = %pred : !llvm.ptr, i32, i1

// Lowers to:
llvm.inline_asm has_side_effects asm_dialect = att 
  "@$2 mbarrier.init.b64 [$0], $1;", "l,r,b" %arg0, %arg2, %arg3 
  : (!llvm.ptr, i32, i1) -> ()
```

Return op name nvvm.ldmatrix as a bitstring.

nvvm.ldmatrix - cooperative matrix load

This op has support for result type inference.

Attributes

• num - Single, I32Attr, 32-bit signless integer attribute
• layout - Single, MMALayoutAttr, NVVM MMA layout
• shape - Single, LdStMatrixShapeAttr, Matrix shape for ldmatrix and stmatrix
• eltType - Single, LdStMatrixEltTypeAttr, Element type for ldmatrix and stmatrix

Operands

• ptr - Single, LLVM_PointerShared, LLVM pointer in address space 3

Results

• res - Single, AnyType, any non-token type

Return op name nvvm.log2 as a bitstring.

nvvm.log2 - Base-2 logarithm (fast approximation)

This op has support for result type inference.

Attributes

• ftz - Single, BoolAttr, bool attribute

Operands

• src - Single, F32, 32-bit float

Results

• res - Single, F32, 32-bit float

Description

Computes a fast approximation of the base-2 logarithm of the input value. The ftz attribute, when set, flushes subnormal inputs and results to sign-preserving zero.

Return op name nvvm.mapa as a bitstring.

nvvm.mapa

Operands

• a - Single, anonymous/composite constraint, LLVM pointer in address space 0 or LLVM pointer in address space 3
• b - Single, I32, 32-bit signless integer

Results

• res - Single, anonymous/composite constraint, LLVM pointer in address space 0 or LLVM pointer in address space 7

Return op name nvvm.match.sync as a bitstring.

nvvm.match.sync - Broadcast and compare a value across threads in warp

This op has support for result type inference.

Attributes

• kind - Single, MatchSyncKindAttr, NVVM match sync kind

Operands

• thread_mask - Single, I32, 32-bit signless integer
• val - Single, anonymous/composite constraint, 32-bit signless integer or 64-bit signless integer

Results

• res - Single, anonymous/composite constraint, 32-bit signless integer or LLVM struct type

Description

The match.sync op performs broadcast and compare of operand val across all non-exited threads in thread_mask and returns a mask depending on the kind and an optional predicate.

The matching operation kinds are:

• any: Returns a mask corresponding to the non-exited threads in the thread_mask that have the same value of operand val.
• all: Returns a mask and a predicate. If all non-exited threads in the thread_mask have the same value of operand val, the predicate is set to true and the mask corresponds to the non-exited threads in the thread_mask. Otherwise, the predicate is set to false and the mask is 0.

For more information, see PTX ISA

Return op name nvvm.mbarrier.arrive as a bitstring.

nvvm.mbarrier.arrive - MBarrier Arrive Operation

This op has support for result type inference.

Attributes

• scope - Single, MemScopeKindAttr, NVVM Memory Scope kind
• relaxed - Single, BoolAttr, bool attribute

Operands

• addr - Single, anonymous/composite constraint, LLVM pointer in address space 0 or LLVM pointer in address space 3 or LLVM pointer in address space 7
• count - Optional, I32, 32-bit signless integer

Results

• res - Optional, I64, 64-bit signless integer

Description

The nvvm.mbarrier.arrive operation performs an arrive-on operation on the mbarrier object at the specified address. Uses the default .release.cta semantics. This release pattern establishes memory ordering for operations occurring in program order before this arrive instruction by making operations from the current thread visible to subsequent operations in other threads within the CTA. When other threads perform corresponding acquire operations (like 'mbarrier.test.wait'), they synchronize with this release pattern.

This operation causes the executing thread to signal its arrival at the barrier.

• res: When the space is not shared_cluster, this operation returns an opaque 64-bit value capturing the phase of the mbarrier object prior to the arrive-on operation. The contents of this return value are implementation-specific. An mbarrier object located in the shared_cluster space cannot return a value.

The operation takes the following operands:

• addr: A pointer to the memory location of the mbarrier object. The addr must be a pointer to generic or shared_cta or shared_cluster memory. When it is generic, the underlying address must be within the shared_cta memory space; otherwise the behavior is undefined.
• count: This specifies the amount by which the pending arrival count is decremented. If the count argument is not specified, the pending arrival count is decremented by 1.
• scope: This specifies the set of threads that directly observe the memory synchronizing effect of the mbarrier.arrive operation.
• space: This indicates the memory space where the mbarrier object resides.
• relaxed: When set to true, the arrive operation has relaxed memory semantics and does not provide any ordering or visibility guarantees.

For more information, see PTX ISA

Return op name nvvm.mbarrier.arrive_drop as a bitstring.

nvvm.mbarrier.arrive_drop - MBarrier Arrive-Drop Operation

This op has support for result type inference.

Attributes

• scope - Single, MemScopeKindAttr, NVVM Memory Scope kind
• relaxed - Single, BoolAttr, bool attribute

Operands

• addr - Single, anonymous/composite constraint, LLVM pointer in address space 0 or LLVM pointer in address space 3 or LLVM pointer in address space 7
• count - Optional, I32, 32-bit signless integer

Results

• res - Optional, I64, 64-bit signless integer

Description

The nvvm.mbarrier.arrive_drop operation decrements the expected arrival count of the mbarrier object by count and then performs an arrive-on operation. When count is not specified, it defaults to 1. The decrement of the expected arrival count applies to all the subsequent phases of the mbarrier object. The remaining semantics are identical to those of the nvvm.mbarrier.arrive operation.

For more information, see PTX ISA

Return op name nvvm.mbarrier.arrive_drop.expect_tx as a bitstring.

nvvm.mbarrier.arrive_drop.expect_tx - MBarrier arrive_drop with expected transaction count

This op has support for result type inference.

Attributes

• scope - Single, MemScopeKindAttr, NVVM Memory Scope kind
• relaxed - Single, BoolAttr, bool attribute

Operands

• addr - Single, anonymous/composite constraint, LLVM pointer in address space 0 or LLVM pointer in address space 3 or LLVM pointer in address space 7
• txcount - Single, I32, 32-bit signless integer

Results

• res - Optional, I64, 64-bit signless integer

Description

The nvvm.mbarrier.arrive_drop.expect_tx operation is similar to the nvvm.mbarrier.arrive.expect_tx operation except that it performs an arrive_drop operation instead of only an arrive operation.

For more information, see PTX ISA

Return op name nvvm.mbarrier.arrive_drop.nocomplete as a bitstring.

nvvm.mbarrier.arrive_drop.nocomplete - MBarrier Arrive-Drop No-Complete Operation

This op has support for result type inference.

Operands

• addr - Single, anonymous/composite constraint, LLVM pointer in address space 0 or LLVM pointer in address space 3
• count - Single, I32, 32-bit signless integer

Results

• res - Single, I64, 64-bit signless integer

Description

The nvvm.mbarrier.arrive_drop.nocomplete operation decrements the expected arrival count of the mbarrier object by the amount count and then performs an arrive-on operation on the mbarrier object with the guarantee that it will not cause the barrier to complete its current phase.

For more information, see PTX ISA

Return op name nvvm.mbarrier.arrive.expect_tx as a bitstring.

nvvm.mbarrier.arrive.expect_tx - MBarrier Arrive with Expected Transaction Count

This op has support for result type inference.

Attributes

• scope - Single, MemScopeKindAttr, NVVM Memory Scope kind
• relaxed - Single, BoolAttr, bool attribute

Operands

• addr - Single, anonymous/composite constraint, LLVM pointer in address space 0 or LLVM pointer in address space 3 or LLVM pointer in address space 7
• txcount - Single, I32, 32-bit signless integer
• predicate - Optional, PtxPredicate, 1-bit signless integer

Results

• res - Optional, I64, 64-bit signless integer

Description

The nvvm.mbarrier.arrive.expect_tx operation performs an expect-tx operation followed by an arrive-on operation on the mbarrier object. Uses the default .release.cta semantics. This release pattern establishes memory ordering for operations occurring in program order before this arrive instruction by making operations from the current thread visible to subsequent operations in other threads within the CTA. When other threads perform corresponding acquire operations (like 'mbarrier.test.wait'), they synchronize with this release pattern.

This operation first performs an expect-tx operation with the specified transaction count, then performs an arrive-on operation with an implicit count of 1. The expect-tx operation increases the expect-count of the mbarrier object by the specified value (i.e. txcount), setting the current phase to expect and track the completion of additional asynchronous transactions.

The operation takes the following operands:

• addr: A pointer to the memory location of the mbarrier object. Uses generic addressing, but the address must still be in the shared memory space.
• txcount: An unsigned integer specifying the expected transaction count for the expect-tx operation. This represents the number of asynchronous transactions expected to complete before the barrier phase completes.
• scope: This specifies the set of threads that directly observe the memory synchronizing effect of the mbarrier.test.wait operation.
• relaxed: When set to true, the arrive operation has relaxed memory semantics and does not provide any ordering or visibility guarantees.
• predicate: Optional predicate for conditional execution used only when lowering to inline-ptx.

For more information, see PTX ISA

Return op name nvvm.mbarrier.arrive.nocomplete as a bitstring.

nvvm.mbarrier.arrive.nocomplete - MBarrier Arrive No-Complete Operation

This op has support for result type inference.

Operands

• addr - Single, anonymous/composite constraint, LLVM pointer in address space 0 or LLVM pointer in address space 3
• count - Single, I32, 32-bit signless integer

Results

• res - Single, I64, 64-bit signless integer

Description

The nvvm.mbarrier.arrive.nocomplete operation performs an arrive-on operation on the mbarrier object with the guarantee that it will not cause the barrier to complete its current phase. Uses the default .release.cta semantics. This release pattern establishes memory ordering for operations occurring in program order before this arrive instruction by making operations from the current thread visible to subsequent operations in other threads within the CTA. When other threads perform corresponding acquire operations (like 'mbarrier.test.wait'), they synchronize with this release pattern.

This operation causes the executing thread to signal its arrival at the barrier with a specified count, but ensures that the barrier phase will not complete as a result of this operation. The operation returns an opaque value that captures the phase of the mbarrier object prior to the arrive-on operation.

The operation takes the following operands:

• addr: A pointer to the memory location of the mbarrier object. The addr must be a pointer to generic or shared::cta memory. When it is generic, the underlying address must be within the shared::cta memory space; otherwise the behavior is undefined.
• count: Integer specifying the count argument to the arrive-on operation. Must be in the valid range as specified in the mbarrier object contents.

For more information, see PTX ISA

Return op name nvvm.mbarrier.complete_tx as a bitstring.

nvvm.mbarrier.complete_tx - MBarrier complete-tx Operation

Attributes

• scope - Single, MemScopeKindAttr, NVVM Memory Scope kind

Operands

• addr - Single, anonymous/composite constraint, LLVM pointer in address space 3 or LLVM pointer in address space 7
• txcount - Single, I32, 32-bit signless integer

Description

The nvvm.mbarrier.complete_tx operation decrements the transaction count of the mbarrier object at addr by txcount. It also signals the completion of asynchronous transactions that were tracked by the current phase. The scope specifies the set of threads that can directly observe the memory synchronizing effect of the mbarrier.complete_tx operation. CTA and CLUSTER are the only allowed values for scope.

For more information, see PTX ISA

Return op name nvvm.mbarrier.expect_tx as a bitstring.

nvvm.mbarrier.expect_tx - MBarrier expect-tx Operation

Attributes

• scope - Single, MemScopeKindAttr, NVVM Memory Scope kind

Operands

• addr - Single, anonymous/composite constraint, LLVM pointer in address space 3 or LLVM pointer in address space 7
• txcount - Single, I32, 32-bit signless integer

Description

The nvvm.mbarrier.expect_tx operation increases the transaction count of the mbarrier located at addr by txcount amount. The scope specifies the set of threads that can directly observe the memory synchronizing effect of the mbarrier.expect_tx operation. CTA and CLUSTER are the only allowed values for scope.

For more information, see PTX ISA

Return op name nvvm.mbarrier.init as a bitstring.

nvvm.mbarrier.init - MBarrier Initialization Op

Operands

• addr - Single, anonymous/composite constraint, LLVM pointer in address space 0 or LLVM pointer in address space 3
• count - Single, I32, 32-bit signless integer
• predicate - Optional, PtxPredicate, 1-bit signless integer

Description

The nvvm.mbarrier.init operation initializes an mbarrier object at the specified memory location.

This operation initializes the mbarrier object with the following state:

• Current phase: 0
• Expected arrival count: count
• Pending arrival count: count
• Transaction count (tx-count): 0

The operation takes the following operands:

• addr: A pointer to the memory location of the mbarrier object. The addr must be a pointer to generic or shared::cta memory. When it is generic, the underlying address must be within the shared::cta memory space; otherwise the behavior is undefined.
• count: Integer specifying the number of threads that will participate in barrier synchronization. Must be in the range [1, 2²⁰ - 1].
• predicate: Optional predicate for conditional execution.

For more information, see PTX ISA

Return op name nvvm.mbarrier.inval as a bitstring.

nvvm.mbarrier.inval - MBarrier Invalidation Operation

Operands

• addr - Single, anonymous/composite constraint, LLVM pointer in address space 0 or LLVM pointer in address space 3

Description

The nvvm.mbarrier.inval operation invalidates an mbarrier object at the specified memory location.

This operation marks the mbarrier object as invalid, making it safe to repurpose the memory location for other uses or to reinitialize it as a new mbarrier object. It is undefined behavior if the mbarrier object is already invalid.

The operation takes the following operand:

• addr: A pointer to the memory location of the mbarrier object. The addr must be a pointer to generic or shared::cta memory. When it is generic, the underlying address must be within the shared::cta memory space; otherwise the behavior is undefined.

For more information, see PTX ISA

Return op name nvvm.mbarrier.test.wait as a bitstring.

nvvm.mbarrier.test.wait - MBarrier Non-Blocking Test Wait Operation

This op has support for result type inference.

Attributes

• scope - Single, MemScopeKindAttr, NVVM Memory Scope kind
• relaxed - Single, BoolAttr, bool attribute

Operands

• addr - Single, anonymous/composite constraint, LLVM pointer in address space 0 or LLVM pointer in address space 3
• stateOrPhase - Single, anonymous/composite constraint, 64-bit signless integer or 32-bit signless integer

Results

• res - Single, I1, 1-bit signless integer

Description

The nvvm.mbarrier.test.wait operation performs a non-blocking test for the completion of a specific phase of an mbarrier object. It uses the default .acquire.cta semantics. This acquire pattern establishes memory ordering for operations occurring in program order after this wait instruction by making operations from other threads in the CTA visible to subsequent operations in the current thread. When this wait completes, it synchronizes with the corresponding release pattern from the mbarrier.arrive operation, establishing memory ordering within the CTA.

This operation tests whether the mbarrier phase specified by the state operand has completed. It is a non-blocking instruction that immediately returns the completion status without suspending the executing thread.

The operation takes the following operands:

• addr: A pointer to the memory location of the mbarrier object. Uses generic addressing, but the address must still be in the shared memory space.
• stateOrPhase: This argument represents a state when it is a 64-bit value and represents a phase when it is a 32-bit value. The state is an opaque value returned by a previous mbarrier.arrive operation on the same mbarrier object during the current or immediately preceding phase. The phase is an integer specifying the phase parity (0 or 1). Even phases have parity 0, odd phases have parity 1.
• scope: This specifies the set of threads that directly observe the memory synchronizing effect of the mbarrier.test.wait operation.
• relaxed: When set to true, the arrive operation has relaxed memory semantics and does not provide any ordering or visibility guarantees.

The operation returns a boolean value indicating whether the specified phase has completed:

• true: The immediately preceding phase has completed
• false: The phase is still incomplete (current phase)

Memory ordering guarantees: When this wait returns true, the following ordering guarantees hold:

1. All memory accesses (except async operations) requested prior to mbarrier.arrive having release semantics by participating CTA threads are visible to the executing thread.
2. All cp.async operations requested prior to cp.async.mbarrier.arrive by participating CTA threads are visible to the executing thread.
3. All cp.async.bulk operations using the same mbarrier object requested prior to mbarrier.arrive having release semantics by participating CTA threads are visible to the executing thread.
4. Memory accesses requested after this wait are not visible to memory accesses performed prior to mbarrier.arrive by other participating threads.
5. No ordering guarantee exists for memory accesses by the same thread between mbarrier.arrive and this wait.

For more information, see PTX ISA

Return op name nvvm.mbarrier.try_wait as a bitstring.

nvvm.mbarrier.try_wait - MBarrier try wait on state or phase with an optional timelimit

This op has support for result type inference.

Attributes

• scope - Single, MemScopeKindAttr, NVVM Memory Scope kind
• relaxed - Single, BoolAttr, bool attribute

Operands

• addr - Single, anonymous/composite constraint, LLVM pointer in address space 0 or LLVM pointer in address space 3
• stateOrPhase - Single, anonymous/composite constraint, 64-bit signless integer or 32-bit signless integer
• ticks - Optional, I32, 32-bit signless integer

Results

• res - Single, I1, 1-bit signless integer

Description

The nvvm.mbarrier.try_wait operation checks whether the specified mbarrier object at addr has completed the given phase. Note that unlike the nvvm.mbarrier.test.wait operation, the try_wait operation is a potentially-blocking one. If the phase is not yet complete, the calling thread may be suspended. A suspended thread resumes execution once the phase completes or when a system-defined timeout occurs. Optionally, the ticks operand can be used to provide a custom timeout (in nanoseconds), overriding the system-defined one. The semantics of this operation and its operands are otherwise similar to those of the nvvm.mbarrier.test.wait Op.

For more information, see PTX ISA

Return op name nvvm.mbarrier.try_wait.parity as a bitstring.

nvvm.mbarrier.try_wait.parity - MBarrier Potentially-Blocking Try Wait with Phase Parity

Operands

• addr - Single, anonymous/composite constraint, LLVM pointer in address space 0 or LLVM pointer in address space 3
• phase - Single, I32, 32-bit signless integer
• ticks - Single, I32, 32-bit signless integer

Description

The nvvm.mbarrier.try_wait.parity operation performs a potentially-blocking test for the completion of a specific phase of an mbarrier object using phase parity. It uses the default .acquire.cta semantics. This acquire pattern establishes memory ordering for operations occurring in program order after this wait instruction by making operations from other threads in the CTA visible to subsequent operations in the current thread. When this wait completes, it synchronizes with the corresponding release pattern from the mbarrier.arrive operation, establishing memory ordering within the CTA.

This operation waits for the completion of the mbarrier phase indicated by the phase parity. While it uses the underlying PTX mbarrier.try_wait.parity instruction, this MLIR operation generates a loop that enforces the test to complete before continuing execution, ensuring the barrier phase is actually completed rather than potentially timing out.

The operation takes the following operands:

• addr: A pointer to the memory location of the mbarrier object. Uses generic addressing, but the address must still be in the shared memory space.
• phase: An integer specifying the phase parity (0 or 1). Even phases have parity 0, odd phases have parity 1.
• ticks: An unsigned integer specifying the suspend time hint in nanoseconds. This may be used instead of the system-dependent time limit.

Memory ordering guarantees: When this wait returns true, the following ordering guarantees hold:

1. All memory accesses (except async operations) requested prior to mbarrier.arrive having release semantics by participating CTA threads are visible to the executing thread.
2. All cp.async operations requested prior to cp.async.mbarrier.arrive by participating CTA threads are visible to the executing thread.
3. All cp.async.bulk operations using the same mbarrier object requested prior to mbarrier.arrive having release semantics by participating CTA threads are visible to the executing thread.
4. Memory accesses requested after this wait are not visible to memory accesses performed prior to mbarrier.arrive by other participating threads.
5. No ordering guarantee exists for memory accesses by the same thread between mbarrier.arrive and this wait.

Implementation behavior: This operation generates a PTX loop that repeatedly calls the underlying mbarrier.try_wait.parity instruction until the barrier phase completes. Unlike the raw PTX instruction which may return without completion after a timeout, this MLIR operation guarantees completion by continuing to loop until the specified phase is reached.

For more information, see PTX ISA

Return op name nvvm.memory.barrier as a bitstring.

nvvm.memory.barrier - Memory barrier operation

Attributes

• scope - Single, MemScopeKindAttr, NVVM Memory Scope kind

Description

membar operation guarantees that prior memory accesses requested by this thread are performed at the specified scope, before later memory operations requested by this thread following the membar instruction.

For more information, see PTX ISA

Return op name nvvm.mma.block_scale as a bitstring.

nvvm.mma.block_scale - cooperative matrix-multiply and accumulate with block scaling

Attributes

• shape - Single, NVVM_MMAShapeAttr, Attribute for MMA operation shape.
• multiplicandAPtxType - Optional, MMATypesAttr, NVVM MMA types
• multiplicandBPtxType - Optional, MMATypesAttr, NVVM MMA types
• scaleVecSize - Single, ScaleVecSizeAttr, MMA Scale Vector Sizes
• blockScaleFormat - Single, BlockScaleFormatAttr, MMA Block Scale Format
• kind - Single, MMABlockScaleKindAttr, Block Scale Kind

Operands

• operandA - Variadic, LLVM_Type, variadic of LLVM dialect-compatible type
• operandB - Variadic, LLVM_Type, variadic of LLVM dialect-compatible type
• operandC - Variadic, LLVM_Type, variadic of LLVM dialect-compatible type
• scaleAData - Single, I32, 32-bit signless integer
• byteIdA - Single, I16, 16-bit signless integer
• threadIdA - Single, I16, 16-bit signless integer
• scaleBData - Single, I32, 32-bit signless integer
• byteIdB - Single, I16, 16-bit signless integer
• threadIdB - Single, I16, 16-bit signless integer

Results

• res - Single, LLVM_AnyStruct, LLVM structure type

Description

The nvvm.mma.block_scale operation collectively performs the operation D = matmul(A * SF_A, B * SF_B) + C using all threads in a warp.

A, B, C and D are dense matrices and SF_A and SF_B are scaling factors. Dimensions of SF_A and SF_B are based on scale vector sizes (x1, x2, x4), and the data type must be either ue8m0 or ue4m3.

All the threads in the warp must execute the same mma.block_scale operation.

This operation follows the same design pattern as nvvm.mma.sync, with additional scaling operands for both A and B matrices.

Example:

%d = nvvm.mma.block_scale A[%a0, %a1] B[%b0, %b1] C[%c0, %c1]
                          scaleA[%scaleAData, %byteIdA, %threadIdA]
                          scaleB[%scaleBData, %byteIdB, %threadIdB]
                          {shape = #nvvm.shape<m = 16, n = 8, k = 64>,
                           multiplicandAPtxType = #nvvm.mma_type<e2m1>,
                           multiplicandBPtxType = #nvvm.mma_type<e2m1>,
                           scaleVecSize = #nvvm.scale_vec_size<x2>,
                           blockScaleFormat = #nvvm.block_scale_format<ue8m0>,
                           kind = #nvvm.block_scale_kind<mxf4nvf4>}
    : (vector<4xf16>, vector<2xf16>, vector<2xf32>) -> !llvm.struct<(f32, f32)>

Return op name nvvm.mma.sp.block_scale as a bitstring.

nvvm.mma.sp.block_scale - cooperative sparse matrix-multiply and accumulate with block scaling

Attributes

• shape - Single, NVVM_MMAShapeAttr, Attribute for MMA operation shape.
• multiplicandAPtxType - Optional, MMATypesAttr, NVVM MMA types
• multiplicandBPtxType - Optional, MMATypesAttr, NVVM MMA types
• scaleVecSize - Single, ScaleVecSizeAttr, MMA Scale Vector Sizes
• blockScaleFormat - Single, BlockScaleFormatAttr, MMA Block Scale Format
• kind - Single, MMABlockScaleKindAttr, Block Scale Kind
• orderedMetadata - Optional, UnitAttr, unit attribute

Operands

• operandA - Variadic, LLVM_Type, variadic of LLVM dialect-compatible type
• operandB - Variadic, LLVM_Type, variadic of LLVM dialect-compatible type
• operandC - Variadic, LLVM_Type, variadic of LLVM dialect-compatible type
• sparseMetadata - Single, I32, 32-bit signless integer
• sparsitySelector - Single, I32, 32-bit signless integer
• scaleAData - Single, I32, 32-bit signless integer
• byteIdA - Single, I16, 16-bit signless integer
• threadIdA - Single, I16, 16-bit signless integer
• scaleBData - Single, I32, 32-bit signless integer
• byteIdB - Single, I16, 16-bit signless integer
• threadIdB - Single, I16, 16-bit signless integer

Results

• res - Single, LLVM_AnyStruct, LLVM structure type

Description

The nvvm.mma.sp.block_scale operation collectively performs the operation D = matmul(A_sparse * SF_A, B * SF_B) + C using all threads in a warp.

A is a sparse matrix, and B, C and D are dense matrices. SF_A and SF_B are scaling factors. Dimensions of SF_A and SF_B are based on scale vector sizes (x1, x2, x4), and the data type must be either ue8m0 or ue4m3.

This operation is similar to nvvm.mma.block_scale but with structured sparsity in the A operand. The sparsity follows the 2:4 structured sparse pattern where 2 out of every 4 elements are non-zero.

All the threads in the warp must execute the same mma.sp.block_scale operation.

The sparseMetadata operand provides the sparsity indices that indicate which elements in the A operand are non-zero. The sparsitySelector controls how the indices are distributed among threads in the warp and should typically be 0 or 1.

This operation follows the same design pattern as nvvm.mma.sp.sync, with additional scaling operands for both A and B matrices. Note that sparse block scale operations always use ordered metadata (sm_90+).

Example:

%d = nvvm.mma.sp.block_scale A[%a0, %a1] B[%b0, %b1] C[%c0, %c1]
                             sparseMetadata[%meta] selector[%sel]
                             scaleA[%scaleAData, %byteIdA, %threadIdA]
                             scaleB[%scaleBData, %byteIdB, %threadIdB]
                             {shape = #nvvm.shape<m = 16, n = 8, k = 128>,
                              multiplicandAPtxType = #nvvm.mma_type<e2m1>,
                              multiplicandBPtxType = #nvvm.mma_type<e2m1>,
                              scaleVecSize = #nvvm.scale_vec_size<x2>,
                              blockScaleFormat = #nvvm.block_scale_format<ue8m0>,
                              kind = #nvvm.block_scale_kind<mxf4>}
    : (vector<2xf16>, vector<2xf16>, vector<2xf32>) -> !llvm.struct<(f32, f32)>

Return op name nvvm.mma.sp.sync as a bitstring.

nvvm.mma.sp.sync - cooperative sparse matrix-multiply and accumulate

Attributes

• shape - Single, NVVM_MMAShapeAttr, Attribute for MMA operation shape.
• intOverflowBehavior - Optional, MMAIntOverflowAttr, MMA overflow options
• multiplicandAPtxType - Optional, MMATypesAttr, NVVM MMA types
• multiplicandBPtxType - Optional, MMATypesAttr, NVVM MMA types
• orderedMetadata - Optional, UnitAttr, unit attribute
• kind - Optional, MMAKindAttr, MMA operation kind

Operands

• operandA - Variadic, LLVM_Type, variadic of LLVM dialect-compatible type
• operandB - Variadic, LLVM_Type, variadic of LLVM dialect-compatible type
• operandC - Variadic, LLVM_Type, variadic of LLVM dialect-compatible type
• sparseMetadata - Single, I32, 32-bit signless integer
• sparsitySelector - Single, I32, 32-bit signless integer

Results

• res - Single, LLVM_AnyStruct, LLVM structure type

Description

The nvvm.mma.sp.sync operation collectively performs the sparse operation D = matmul(A_sparse, B) + C using all threads in a warp.

This operation is similar to nvvm.mma.sync but with structured sparsity in the A operand. The sparsity follows the 2:4 structured sparse pattern where 2 out of every 4 elements are non-zero.

All the threads in the warp must execute the same mma.sp.sync operation.

The sparseMetadata operand provides the sparsity indices that indicate which elements in the A operand are non-zero. The sparsitySelector controls how the indices are distributed among threads in the warp and should typically be 0 or 1.

The optional orderedMetadata attribute specifies the metadata ordering:

• Absence (default): Uses standard sparse metadata ordering
• Presence: Uses ordered metadata (PTX ISA 8.5+, sm_90+)

The optional kind attribute specifies mixed-precision modes for FP8 operations:

• f8f6f4: Enables e3m2, e2m3, e2m1 FP8 types and f16 accumulator (PTX ISA 8.7+, sm_90+)
• Only valid with ordered metadata and m16n8k64 shape

The shapes, layouts, and data types follow the same constraints as the regular nvvm.mma.sync operation, but the A operand contains only the non-zero elements in compressed format.

Example:

%d = nvvm.mma.sp.sync A[%a0, %a1] B[%b0, %b1] C[%c0, %c1]
                      sparseMetadata[%meta] selector[%sel]
                      {shape = {k = 32 : i32, m = 16 : i32, n = 8 : i32}}
    : (vector<2xf16>, vector<2xf16>, vector<2xf16>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>)>

// With ordered metadata:
%d = nvvm.mma.sp.sync A[%a0, %a1] B[%b0, %b1] C[%c0, %c1]
                      sparseMetadata[%meta] selector[%sel]
                      {orderedMetadata, shape = {k = 32 : i32, m = 16 : i32, n = 8 : i32}}
    : (vector<2xf16>, vector<2xf16>, vector<2xf16>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>)>

Return op name nvvm.mma.sync as a bitstring.

nvvm.mma.sync - cooperative matrix-multiply and accumulate

Attributes

• shape - Single, NVVM_MMAShapeAttr, Attribute for MMA operation shape.
• b1Op - Optional, MMAB1OpAttr, MMA binary operations
• intOverflowBehavior - Optional, MMAIntOverflowAttr, MMA overflow options
• layoutA - Single, MMALayoutAttr, NVVM MMA layout
• layoutB - Single, MMALayoutAttr, NVVM MMA layout
• multiplicandAPtxType - Optional, MMATypesAttr, NVVM MMA types
• multiplicandBPtxType - Optional, MMATypesAttr, NVVM MMA types