Enum InstructionOpcode 

pub enum InstructionOpcode {
    Abs,
    ActiveMask,
    Add,
    AddCc,
    Addc,
    Alloca,
    And,
    ApplyPriority,
    Atom,
    Bar,
    BarWarpSync,
    Barrier,
    BarrierCluster,
    Bfe,
    Bfi,
    Bfind,
    Bmsk,
    Bra,
    Brev,
    Brkpt,
    BrxIdx,
    Call,
    Clmad,
    ClusterLaunchControlQueryCancel,
    ClusterLaunchControlTryCancel,
    Clz,
    Cnot,
    Copysign,
    Cos,
    CpAsyncMbarrierArrive,
    CreatePolicy,
    Cvt,
    CvtPack,
    Cvta,
    Discard,
    Div,
    Dp2A,
    Dp4A,
    ElectSync,
    Ex2,
    Exit,
    FabricSubmit,
    FabricTryGet,
    FabricTryPullred,
    FabricTryPut,
    FabricTryRed,
    FabricWait,
    Fence,
    Fma,
    Fns,
    GetCtaRank,
    GridDepControl,
    IsSpaceP,
    IsTypeP,
    Ld,
    LdGlobalNc,
    LdMatrix,
    Ldu,
    Lg2,
    Lop3,
    Mad,
    MadCc,
    Mad24,
    Madc,
    Mapa,
    MatchSync,
    Max,
    MbarrierArrive,
    MbarrierArriveDrop,
    MbarrierCheckLayout,
    MbarrierCompleteTx,
    MbarrierExpectTx,
    MbarrierInit,
    MbarrierInval,
    MbarrierPendingCount,
    MbarrierTestWait,
    MbarrierTryWait,
    Membar,
    Min,
    Mma,
    MmaSp,
    MmaSpOrderedMetadata,
    Mov,
    MovMatrix,
    Mul,
    Mul24,
    MultimemLdReduce,
    MultimemRed,
    MultimemRedAsync,
    MultimemSt,
    MultimemStAsync,
    Nanosleep,
    Neg,
    Not,
    Or,
    PmEvent,
    Popc,
    Prefetch,
    Prefetchu,
    Prmt,
    Rcp,
    RcpApproxFtzF64,
    Red,
    RedAsync,
    ReduxSync,
    Rem,
    Ret,
    Rsqrt,
    RsqrtApproxFtzF64,
    Sad,
    Selp,
    Set,
    SetMaxNReg,
    Setp,
    Shf,
    Shfl,
    ShflSync,
    Shl,
    Shr,
    Sin,
    Slct,
    Sqrt,
    St,
    StAsync,
    StBulk,
    StackRestore,
    StackSave,
    StMatrix,
    Sub,
    SubCc,
    Subc,
    Suld,
    Suq,
    Sured,
    Sust,
    Szext,
    Tanh,
    Tcgen05Alloc,
    Tcgen05Commit,
    Tcgen05Cp,
    Tcgen05Dealloc,
    Tcgen05Fence,
    Tcgen05Ld,
    Tcgen05RelinquishAllocPermit,
    Tcgen05Shift,
    Tcgen05St,
    Tcgen05Wait,
    TensorMapCpFenceProxy,
    TensormapReplace,
    Testp,
    Tex,
    Tld4,
    Trap,
    Txq,
    VAbsDiff,
    VAbsDiff2,
    VAbsDiff4,
    VAdd,
    VAdd2,
    VAdd4,
    VAvrg2,
    VAvrg4,
    VMad,
    VMax,
    VMax2,
    VMax4,
    VMin,
    VMin2,
    VMin4,
    Vote,
    VoteSync,
    VSet,
    VSet2,
    VSet4,
    VShl,
    VShr,
    VSub,
    VSub2,
    VSub4,
    WgmmaCommitGroup,
    WgmmaFence,
    WgmmaMmaAsync,
    WgmmaMmaAsyncSp,
    WgmmaWaitGroup,
    WmmaLoad,
    WmmaMma,
    WmmaStore,
    Xor,
}

Variants

Abs

Absolute value.

ActiveMask

Queries the active threads within a warp.

Add

Add two values.

AddCc

Add two values with carry-out.

Addc

Add two values with carry-in and optional carry-out.

Alloca

Dynamically allocate memory on stack.

And

Bitwise AND.

ApplyPriority

Apply the cache eviction priority to the specified address in the specified cache level.

Atom

Atomic reduction operations for thread-to-thread communication.

Bar

Barrier synchronization.

BarWarpSync

Barrier synchronization for threads in a warp.

Barrier

Barrier synchronization.

BarrierCluster

Barrier synchronization within a cluster.

Bfe

Bit Field Extract.

Bfi

Bit Field Insert.

Bfind

Find most significant non-sign bit.

Bmsk

Bit Field Mask.

Bra

Branch to a target and continue execution there.

Brev

Bit reverse.

Brkpt

Breakpoint.

BrxIdx

Branch to a label indexed from a list of potential branch targets.

Call

Call a function, recording the return location.

Clmad

Perform carryless multiply and add operation.

ClusterLaunchControlQueryCancel

Queries response of clusterlaunchcontrol.try_cancel operation.

ClusterLaunchControlTryCancel

Requests cancellation of cluster which is not launched yet.

Clz

Count leading zeros.

Cnot

C/C++ style logical negation.

Copysign

Copy sign of one input to another.

Cos

Find the cosine of a value.

CpAsyncMbarrierArrive

Makes the mbarrier object track all prior cp.async operations initiated by the executing thread.

CreatePolicy

Create a cache eviction policy for the specified cache level.

Cvt

Convert a value from one type to another.

CvtPack

Convert two integer values from one integer type to another and pack the results.

Cvta

Convert address from .const, Kernel Function Parameters ( .param ), .global, .local, or .shared state space to generic, or vice-versa. Take the generic address of a variable declared in .const, Kernel Function Parameters ( .param ), .global, .local, or .shared state space.

Discard

Discard the data at the specified address range and cache level.

Div

Divide one value by another.

Dp2A

Two-way dot product-accumulate.

Dp4A

Four-way byte dot product-accumulate.

ElectSync

Elect a leader thread from a set of threads.

Ex2

Find the base-2 exponential of a value.

Exit

Terminate a thread.

FabricSubmit

Submits prior fabric operations issued by the current thread.

FabricTryGet

Asynchronous copy from fabric handle.

FabricTryPullred

Asynchronous copy from fabric handle with pull-reduction.

FabricTryPut

Asynchronous copy to fabric handle.

FabricTryRed

Asynchronous copy to fabric handle with reduction.

FabricWait

Wait on local shared-memory reads of submitted fabric operations.

Fence

Enforce an ordering of memory operations.

Fma

Fused multiply-add.

Fns

Find the n-th set bit

GetCtaRank

Generate the CTA rank of the address.

GridDepControl

Control execution of dependent grids.

IsSpaceP

Query whether a generic address falls within a specified state space window.

IsTypeP

Query whether a register points to an opaque variable of a specified type.

Ld

Load a register variable from an addressable state space variable.

LdGlobalNc

Load a register variable from global state space via non-coherent cache.

LdMatrix

Collectively load one or more matrices from shared memory for mma instruction

Ldu

Load read-only data from an address that is common across threads in the warp.

Lg2

Find the base-2 logarithm of a value.

Lop3

Arbitrary logical operation on 3 inputs.

Mad

Multiply two values, optionally extract the high or low half of the intermediate result, and add a third value.

MadCc

Multiply two values, extract high or low half of result, and add a third value with carry-out.

Mad24

Multiply two 24-bit integer values and add a third value.

Madc

Multiply two values, extract high or low half of result, and add a third value with carry-in and optional carry-out.

Mapa

Map the address of the shared variable in the target CTA.

MatchSync

Broadcast and compare a value across threads in warp.

Max

Find the maximum of two values.

MbarrierArrive

Performs arrive-on operation on the mbarrier object .

MbarrierArriveDrop

Decrements the expected count of the mbarrier object and performs arrive-on operation .

MbarrierCheckLayout

Check the layout of the mbarrier object .

MbarrierCompleteTx

Performs complete-tx operation on the mbarrier object .

MbarrierExpectTx

Performs expect-tx operation on the mbarrier object .

MbarrierInit

Initialize the mbarrier object .

MbarrierInval

Invalidates the mbarrier object .

MbarrierPendingCount

Query the pending arrival count from the opaque mbarrier state.

MbarrierTestWait

Checks whether the mbarrier object has completed the phase.

MbarrierTryWait

Checks whether the mbarrier object has completed the phase.

Membar

Enforce an ordering of memory operations.

Min

Find the minimum of two values.

Mma

Perform matrix multiply-and-accumulate operation

MmaSp

Perform matrix multiply-and-accumulate operation with sparse matrix A

MmaSpOrderedMetadata

Perform matrix multiply-and-accumulate operation with sparse matrix A

Mov

Set a register variable with the value of a register variable or an immediate value. Take the non-generic address of a variable in global, local, or shared state space.

MovMatrix

Transpose a matrix in registers across the warp.

Mul

Multiply two values.

Mul24

Multiply two 24-bit integer values.

MultimemLdReduce

The multimem.* operations operate on multimem addresses and accesses all of the multiple memory locations which the multimem address points to.

MultimemRed

The multimem.* operations operate on multimem addresses and accesses all of the multiple memory locations which the multimem address points to.

MultimemRedAsync

Perform asynchronous reduction with release ordering on the multimem address.

MultimemSt

The multimem.* operations operate on multimem addresses and accesses all of the multiple memory locations which the multimem address points to.

MultimemStAsync

Perform asynchronous store with release ordering on the multimem address.

Nanosleep

Suspend the thread for an approximate delay given in nanoseconds.

Neg

Arithmetic negate.

Not

Bitwise negation; one’s complement.

Or

Biwise OR.

PmEvent

Trigger one or more Performance Monitor events.

Popc

Population count.

Prefetch

Prefetch line containing a generic address at a specified level of memory hierarchy, in specified state space.

Prefetchu

Prefetch line containing a generic address at a specified level of memory hierarchy, in specified state space.

Prmt

Permute bytes from register pair.

Rcp

Take the reciprocal of a value.

RcpApproxFtzF64

Compute a fast, gross approximation to the reciprocal of a value.

Red

Reduction operations on global and shared memory.

RedAsync

Asynchronous reduction operation.

ReduxSync

Perform reduction operation on the data from each predicated active thread in the thread group.

Rem

The remainder of integer division.

Ret

Return from function to instruction after call.

Rsqrt

Take the reciprocal of the square root of a value.

RsqrtApproxFtzF64

Compute an approximation of the square root reciprocal of a value.

Sad

Sum of absolute differences.

Selp

Select between source operands, based on the value of the predicate source operand.

Set

Compare two numeric values with a relational operator, and optionally combine this result with a predicate value by applying a Boolean operator.

SetMaxNReg

Hint to change the number of registers owned by the warp.

Setp

Compare two numeric values with a relational operator, and (optionally) combine this result with a predicate value by applying a Boolean operator.

Shf

Funnel shift.

Shfl

Register data shuffle within threads of a warp.

ShflSync

Register data shuffle within threads of a warp.

Shl

Shift bits left, zero-fill on right.

Shr

Shift bits right, sign or zero-fill on left.

Sin

Find the sine of a value.

Slct

Select one source operand, based on the sign of the third operand.

Sqrt

Take the square root of a value.

St

Store data to an addressable state space variable.

StAsync

Asynchronous store operation.

StBulk

Initializes a region of memory as specified by state space.

StackRestore

Update the stack pointer with a new value.

StackSave

Save the value of stack pointer into a register.

StMatrix

Collectively store one or more matrices to shared memory.

Sub

Subtract one value from another.

SubCc

Subtract one value from another, with borrow-out.

Subc

Subtract one value from another, with borrow-in and optional borrow-out.

Suld

Load from surface memory.

Suq

Query a surface attribute.

Sured

Reduce surface memory.

Sust

Store to surface memory.

Szext

Sign-extend or Zero-extend.

Tanh

Find the hyperbolic tangent of a value (in radians)

Tcgen05Alloc

Dynamic Tensor Memory allocation management instructions

Tcgen05Commit

Makes the mbarrier object track the completion of all prior async-tcgen05 operations initiated by the executing thread.

Tcgen05Cp

Initiates an asynchronous copy operation from shared memory to the Tensor Memory .

Tcgen05Dealloc

Dynamic Tensor Memory allocation management instructions

Tcgen05Fence

Specialized fence for the asynchronous tcgen05 operations.

Tcgen05Ld

Asynchronous collective load from tensor memory into registers.

Tcgen05RelinquishAllocPermit

Dynamic Tensor Memory allocation management instructions

Tcgen05Shift

Asynchronously shift down the rows of the matrix in the Tensor Memory for a warp.

Tcgen05St

Asynchronous collective store to tensor memory from registers.

Tcgen05Wait

Waits for the completion of all prior asynchronous tcgen05.ld / tcgen05.st instructions.

TensorMapCpFenceProxy

A fused copy and fence operation.

TensormapReplace

Modifies the field of a tensor-map object.

Testp

Test floating-point property.

Tex

Perform a texture memory lookup.

Tld4

Perform a texture fetch of the 4-texel bilerp footprint.

Trap

Perform trap operation.

Txq

Query texture and sampler attributes.

VAbsDiff

Integer byte/half-word/word addition/subtraction.

VAbsDiff2

Integer dual half-word SIMD addition/subtraction.

VAbsDiff4

Integer quad byte SIMD addition/subtraction.

VAdd

Integer byte/half-word/word addition/subtraction.

VAdd2

Integer dual half-word SIMD addition/subtraction.

VAdd4

Integer quad byte SIMD addition/subtraction.

VAvrg2

Integer dual half-word SIMD addition/subtraction.

VAvrg4

Integer quad byte SIMD addition/subtraction.

VMad

Integer byte/half-word/word multiply-accumulate.

VMax

Integer byte/half-word/word addition/subtraction.

VMax2

Integer dual half-word SIMD addition/subtraction.

VMax4

Integer quad byte SIMD addition/subtraction.

VMin

Integer byte/half-word/word addition/subtraction.

VMin2

Integer dual half-word SIMD addition/subtraction.

VMin4

Integer quad byte SIMD addition/subtraction.

Vote

Vote across thread group.

VoteSync

Vote across thread group.

VSet

Integer byte/half-word/word comparison.

VSet2

Integer dual half-word SIMD comparison.

VSet4

Integer quad byte SIMD comparison.

VShl

Integer byte/half-word/word left/right shift.

VShr

Integer byte/half-word/word left/right shift.

VSub

Integer byte/half-word/word addition/subtraction.

VSub2

Integer dual half-word SIMD addition/subtraction.

VSub4

Integer quad byte SIMD addition/subtraction.

WgmmaCommitGroup

Commits all prior uncommitted wgmma.mma_async operations into a wgmma-group .

WgmmaFence

Enforce an ordering of register accesses between wgmma.mma_async and other operations.

WgmmaMmaAsync

Perform matrix multiply-and-accumulate operation across warpgroup

WgmmaMmaAsyncSp

Perform matrix multiply-and-accumulate operation with sparse matrix A across warpgroup

WgmmaWaitGroup

Signal the completion of a preceding warpgroup operation.

WmmaLoad

Collectively load a matrix from memory for WMMA

WmmaMma

Perform a single matrix multiply-and-accumulate operation across a warp

WmmaStore

Collectively store a matrix into memory for WMMA

Xor

Bitwise exclusive-OR (inequality).

Trait Implementations

impl Clone for InstructionOpcode

fn clone(&self) -> InstructionOpcode

Returns a duplicate of the value.

fn clone_from(&mut self, source: &Self)

Performs copy-assignment from source.

impl Copy for InstructionOpcode

impl Debug for InstructionOpcode

fn fmt(&self, f: &mut Formatter<'_>) -> Result

Formats the value using the given formatter.

impl Eq for InstructionOpcode

impl PartialEq for InstructionOpcode

fn eq(&self, other: &InstructionOpcode) -> bool

Tests for self and other values to be equal, and is used by ==.

fn ne(&self, other: &Rhs) -> bool

Tests for !=. The default implementation is almost always sufficient, and should not be overridden without very good reason.

impl StructuralPartialEq for InstructionOpcode