tvm_tirx.tirx.script.ptx#

Public canonical TVMScript dialect namespace.

Functions#

`pred`()	Tag an operand as a `.pred` register -- see `PredArg`.
`addr`(byte_offset)	Form `[base+byte_offset]` for an eligible PTX address operand.
`add`
`sub`
`fma`
`mul`
`set`
`ld`
`prefetch`
`cvt`
`cp`
`mbarrier`
`tcgen05`
`applypriority`
`spcompress`
`spdecompress`
`atom`
`red`
`multimem_cp`
`ldmatrix`
`max`
`min`
`fns`
`mad`
`clmad`
`mul24`
`mad24`
`sad`
`div`
`rem`
`abs`
`neg`
`popc`
`clz`
`bfind`
`brev`
`bfe`
`bfi`
`szext`
`bmsk`
`dp4a`
`dp2a`
`rcp`
`sqrt`
`rsqrt`
`ex2`
`sin`
`cos`
`lg2`
`tanh`
`testp`
`copysign`
`setp`
`selp`
`slct`
`and`
`or`
`xor`
`not`
`cnot`
`lop3`
`shf`
`shl`
`shr`
`mov`
`st`
`st_bulk`
`st_async`
`multimem_st_async`
`multimem_ld_reduce`
`multimem_st`
`multimem_red`
`createpolicy`
`tensormap_replace`
`prmt`
`ldu`
`prefetchu`
`discard`
`isspacep`
`getctarank`
`cvt_pack`
`shfl_sync`
`cvta`
`mapa`
`fabric`
`bar`
`barrier`
`vote_sync`
`match`
`activemask`
`elect_sync`
`redux_sync`
`fence`
`griddepcontrol`
`tensormap_cp_fenceproxy`
`red_async`
`multimem_red_async`
`clusterlaunchcontrol`
`mma`
`movmatrix`
`stmatrix`
`wgmma`
`setmaxnreg`