PTX Instructions

• ld
• st
• shr
• shl
• bmsk
• elect.sync
• prmt
• barrier.cluster
• bfind
• clusterlaunchcontrol
• cp.async.bulk
• cp.async.bulk.commit_group
• cp.async.bulk.wait_group
• cp.async.bulk.tensor
• cp.async.mbarrier.arrive
• cp.reduce.async.bulk
• cp.reduce.async.bulk.tensor
• exit
• fence
• getctarank
• mapa
• mbarrier.init
• mbarrier.inval
• mbarrier.arrive
• mbarrier.expect_tx
• mbarrier.test_wait / mbarrier.try_wait
• multimem.ld_reduce
• multimem.red
• multimem.st
• red.async
• shfl.sync
• st.async
• st.bulk
• tcgen05.alloc
• tcgen05.commit
• tcgen05.cp
• tcgen05.fence
• tcgen05.ld
• tcgen05.mma
• tcgen05.mma.ws
• tcgen05.shift
• tcgen05.st
• tcgen05.wait
• tensormap.replace
• tensormap.cp_fenceproxy
• trap
• setmaxnreg
• Special registers
• applypriority.async.bulk
• cp.async.bulk.prefetch
• cp.async.bulk.prefetch.tensor
• cp.async.mbarrier.arrive.noinc
• fabric.submit
• fabric.try_get
• fabric.try_pullred
• fabric.try_put
• fabric.try_red
• fabric.wait
• ldmatrix
• mbarrier.check_layout
• mbarrier.complete_tx
• mbarrier.pending_count
• prefetch
• stmatrix

Instructions by section

Integer Arithmetic Instructions

Instruction	Available in libcu++	Header
sad	No
div	No
rem	No
abs	No
neg	No
min	No
max	No
popc	No
clz	No
bfind	CCCL 3.0.0	<cuda/ptxs/bfind.h>
fns	No
brev	No
bfe	No
bfi	No
szext	No
bmsk	Yes, CCCL 3.0.0 / CUDA 13.0	<cuda/ptxs/bmsk.h>
dp4a	No
dp2a	No

Extended-Precision Integer Arithmetic Instructions

Instruction	Available in libcu++	Header
add.cc	No
addc	No
sub.cc	No
subc	No
mad.cc	No
madc	No

Floating-Point Instructions

Instruction	Available in libcu++	Header
testp	No
copysign	No
add	No
sub	No
mul	No
fma	No
mad	No
div	No
abs	No
neg	No
min	No
max	No
rcp	No
rcp.approx.ftz.f64	No
sqrt	No
rsqrt	No
rsqrt.approx.ftz.f64	No
sin	No
cos	No
lg2	No
ex2	No
tanh	No

Half Precision Floating-Point Instructions

Instruction	Available in libcu++	Header
add	No
sub	No
mul	No
fma	No
neg	No
abs	No
min	No
max	No
tanh	No
ex2	No

Comparison and Selection Instructions

Instruction	Available in libcu++	Header
set	No
setp	No
selp	No
slct	No

Half Precision Comparison Instructions

Instruction	Available in libcu++	Header
set	No
setp	No

Logic and Shift Instructions

Instruction	Available in libcu++	Header
and	No
or	No
xor	No
not	No
cnot	No
lop3	No
shf	No
shl	Yes, CCCL 3.0.0 / CUDA 13.0	<cuda/ptxs/shl.h>
shr	Yes, CCCL 3.0.0 / CUDA 13.0	<cuda/ptxs/shr.h>

Data Movement and Conversion Instructions

Instruction	Available in libcu++	Header
mov	No
shfl	No
shfl.sync	Yes, CCCL 2.9.0 / CUDA 12.9	<cuda/ptxs/shfl_sync.h>
prmt	Yes, CCCL 3.0.0 / CUDA 13.0	<cuda/ptxs/prmt.h>
ld	Yes, CCCL 3.0.0 / CUDA 13.0	<cuda/ptxs/ld.h>
ld.global.nc	Yes, CCCL 3.0.0 / CUDA 13.0	<cuda/ptxs/ld.h>
ldu	No
st	Yes, CCCL 3.0.0 / CUDA 13.0	<cuda/ptxs/st.h>
st.async	CCCL 2.3.0 / CUDA 12.4	<cuda/ptxs/st_async.h>
st.bulk	CCCL 2.8 / CUDA 12.9	<cuda/ptxs/st_bulk.h>
multimem.ld_reduce	CCCL 2.8 / CUDA 12.9	<cuda/ptxs/multimem_ld_reduce.h>
multimem.st	CCCL 2.8 / CUDA 12.9	<cuda/ptxs/multimem_st.h>
multimem.red	CCCL 2.8 / CUDA 12.9	<cuda/ptxs/multimem_red.h>
prefetch, prefetchu	CCCL 3.4.0 / CUDA 13.4	<cuda/ptxs/prefetch.h>
applypriority	No
discard	No
createpolicy	No
isspacep	No
cvta	No
cvt	No
cvt.pack	No
mapa	No
getctarank	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/getctarank.h>

Data Movement and Conversion Instructions: Asynchronous copy

Instruction	Available in libcu++	Header
cp.async	No
cp.async.commit_group	No
cp.async.wait_group	No
cp.async.bulk	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/cp_async_bulk.h>
cp.reduce.async.bulk	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/cp_reduce_async_bulk.h>
cp.async.bulk.prefetch	CCCL 3.4.0 / CUDA 13.4	<cuda/ptxs/cp_async_bulk_prefetch.h>
cp.async.bulk.tensor	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/cp_async_bulk_tensor.h>
cp.reduce.async.bulk.tensor	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/cp_reduce_async_bulk_tensor.h>
cp.async.bulk.prefetch.tensor	CCCL 3.4.0 / CUDA 13.4	<cuda/ptxs/cp_async_bulk_prefetch_tensor.h>
cp.async.bulk.commit_group	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/cp_async_bulk_commit_group.h>
cp.async.bulk.wait_group	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/cp_async_bulk_wait_group.h>
applypriority.async.bulk	CCCL 3.4.0 / CUDA 13.4	<cuda/ptxs/applypriority_async_bulk.h>
tensormap.replace	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/tensormap_replace.h>

Texture Instructions

Instruction	Available in libcu++	Header
tex	No
tld4	No
txq	No
istypep	No

Surface Instructions

Instruction	Available in libcu++	Header
suld	No
sust	No
sured	No
suq	No

Control Flow Instructions

Instruction	Available in libcu++	Header
{}	No
@	No
bra	No
bra	No
call	No
ret	No
exit	CCCL 3.0.0	<cuda/ptxs/exit.h>

Parallel Synchronization and Communication Instructions

Instruction	Available in libcu++	Header
bar, barrier	No
bar.warp.sync	No
barrier.cluster	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/barrier_cluster.h>
membar	No
fence	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/fence.h>
atom	No
red	No
red.async	CCCL 2.3.0 / CUDA 12.4	<cuda/ptxs/red_async.h>
vote	No
vote.sync	No
match.sync	No
activemask	No
redux.sync	No
griddepcontrol	No
elect.sync	CCCL 3.1.0 / CUDA 13.1	<cuda/ptxs/elect_sync.h>
fabric.submit	CCCL 3.4.0 / CUDA 13.4	<cuda/ptxs/fabric_submit.h>
fabric.try_get	CCCL 3.4.0 / CUDA 13.4	<cuda/ptxs/fabric_try_get.h>
fabric.try_pullred	CCCL 3.4.0 / CUDA 13.4	<cuda/ptxs/fabric_try_pullred.h>
fabric.try_put	CCCL 3.4.0 / CUDA 13.4	<cuda/ptxs/fabric_try_put.h>
fabric.try_red	CCCL 3.4.0 / CUDA 13.4	<cuda/ptxs/fabric_try_red.h>
fabric.wait	CCCL 3.4.0 / CUDA 13.4	<cuda/ptxs/fabric_wait.h>

Parallel Synchronization and Communication Instructions: mbarrier

Instruction	Available in libcu++	Header
mbarrier.init	CCCL 2.5.0 / CUDA Future	<cuda/ptxs/mbarrier_init.h>
mbarrier.inval	CCCL 3.2.0 / CUDA 13.2	<cuda/ptxs/mbarrier_inval.h>
mbarrier.complete_tx	CCCL 3.4.0 / CUDA 13.4	<cuda/ptxs/mbarrier_complete_tx.h>
mbarrier.arrive	CCCL 2.3.0 / CUDA 12.4	<cuda/ptxs/mbarrier_arrive.h>
mbarrier.arrive_drop	CCCL 3.4.0 / CUDA 13.4	<cuda/ptxs/mbarrier_arrive.h>
mbarrier.arrive.expect_tx	CCCL 3.4.0 / CUDA 13.4	<cuda/ptxs/mbarrier_arrive.h>
mbarrier.arrive.noComplete	CCCL 3.4.0 / CUDA 13.4	<cuda/ptxs/mbarrier_arrive.h>
cp.async.mbarrier.arrive	CCCL 2.8 / CUDA 12.9	<cuda/ptxs/cp_async_mbarrier_arrive.h>
cp.async.mbarrier.arrive.noinc	CCCL 3.4.0 / CUDA 13.4	<cuda/ptxs/cp_async_mbarrier_arrive.h>
mbarrier.check_layout	CCCL 3.4.0 / CUDA 13.4	<cuda/ptxs/mbarrier_check_layout.h>
mbarrier.expect_tx	CCCL 2.8 / CUDA 12.9	<cuda/ptxs/mbarrier_expect_tx.h>
mbarrier.test_wait	CCCL 2.3.0 / CUDA 12.4	<cuda/ptxs/mbarrier_wait.h>
mbarrier.test_wait.parity	CCCL 3.4.0 / CUDA 13.4	<cuda/ptxs/mbarrier_wait.h>
mbarrier.try_wait	CCCL 2.3.0 / CUDA 12.4	<cuda/ptxs/mbarrier_wait.h>
mbarrier.try_wait.parity	CCCL 3.4.0 / CUDA 13.4	<cuda/ptxs/mbarrier_wait.h>
mbarrier.pending_count	CCCL 3.4.0 / CUDA 13.4	<cuda/ptxs/mbarrier_pending_count.h>
tensormap.cp_fenceproxy	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/tensormap_cp_fenceproxy.h>
clusterlaunchcontrol.try_cancel	CCCL 2.8 / CUDA 12.9	<cuda/ptxs/clusterlaunchcontrol.h>
clusterlaunchcontrol.query_cancel	CCCL 2.8 / CUDA 12.9	<cuda/ptxs/clusterlaunchcontrol.h>

Warp Level Matrix Multiply-Accumulate Instructions

Instruction	Available in libcu++	Header
wmma.load	No
wmma.store	No
wmma.mma	No
mma	No
ldmatrix	CCCL 3.4.0 / CUDA 13.4	<cuda/ptxs/ldmatrix.h>
stmatrix	CCCL 3.4.0 / CUDA 13.4	<cuda/ptxs/stmatrix.h>
movmatrix	No
mma.sp	No

Asynchronous Warpgroup Level Matrix Multiply-Accumulate Instructions

Instruction	Available in libcu++	Header
wgmma.mma_async	No
wgmma.mma_async.sp	No
wgmma.fence	No
wgmma.commit_group	No
wgmma.wait_group	No

TensorCore 5th Generation Family Instructions

Instruction	Available in libcu++	Header
tcgen05.alloc	CCCL 2.8 / CUDA 12.9	<cuda/ptxs/tcgen05_alloc.h>
tcgen05.commit	CCCL 2.8 / CUDA 12.9	<cuda/ptxs/tcgen05_commit.h>
tcgen05.cp	CCCL 2.8 / CUDA 12.9	<cuda/ptxs/tcgen05_cp.h>
tcgen05.fence	CCCL 2.8 / CUDA 12.9	<cuda/ptxs/tcgen05_fence.h>
tcgen05.ld	CCCL 2.8 / CUDA 12.9	<cuda/ptxs/tcgen05_ld.h>
tcgen05.mma	CCCL 2.8 / CUDA 12.9	<cuda/ptxs/tcgen05_mma.h>
tcgen05.mma.ws	CCCL 2.8 / CUDA 12.9	<cuda/ptxs/tcgen05_mma_ws.h>
tcgen05.shift	CCCL 2.8 / CUDA 12.9	<cuda/ptxs/tcgen05_shift.h>
tcgen05.st	CCCL 2.8 / CUDA 12.9	<cuda/ptxs/tcgen05_st.h>
tcgen05.wait	CCCL 2.8 / CUDA 12.9	<cuda/ptxs/tcgen05_wait.h>

Stack Manipulation Instructions

Instruction	Available in libcu++	Header
stacksave	No
stackrestore	No
alloca	No

Video Instructions

Instruction	Available in libcu++	Header
vadd, vsub, vabsdiff, vmin, vmax	No
vshl, vshr	No
vmad	No
vset	No

SIMD Video Instructions

Instruction	Available in libcu++	Header
vadd2, vsub2, vavrg2, vabsdiff2, vmin2, vmax2	No
vset2	No
vadd4, vsub4, vavrg4, vabsdiff4, vmin4, vmax4	No
vset4	No

Miscellaneous Instructions

Instruction	Available in libcu++	Header
brkpt	No
nanosleep	No
pmevent	No
trap	CCCL 3.0.0	<cuda/ptxs/trap.h>
setmaxnreg	CCCL 3.2.0 / CUDA 13.2	<cuda/ptxs/setmaxnreg.h>

Special registers <libcudacxx-ptx-instructions-special-registers>

Instruction	Available in libcu++	Header
tid	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>
ntid	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>
laneid	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>
warpid	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>
nwarpid	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>
ctaid	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>
nctaid	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>
smid	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>
nsmid	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>
gridid	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>
is_explicit_cluster	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>
clusterid	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>
nclusterid	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>
cluster_ctaid	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>
cluster_nctaid	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>
cluster_ctarank	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>
cluster_nctarank	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>
lanemask_eq	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>
lanemask_le	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>
lanemask_lt	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>
lanemask_ge	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>
lanemask_gt	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>
clock, clock_hi	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>
clock64	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>
pm0	No
pm0_64	No
envreg	No
globaltimer, globaltimer_lo, globaltimer_hi	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>
reserved_smem_offset_begin, reserved_smem_offset_end, reserved_smem_offset_cap, reserved_smem_offset_2	No
total_smem_size	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>
aggr_smem_size	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>
dynamic_smem_size	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>
current_graph_exec	CCCL 2.4.0 / CUDA 12.5	<cuda/ptxs/get_sreg.h>