Special registers#

tid.x#

// mov.u32 sreg_value, %%tid.x; // PTX ISA 20
template <typename = void>
__device__ static inline uint32_t get_sreg_tid_x();

tid.y#

// mov.u32 sreg_value, %%tid.y; // PTX ISA 20
template <typename = void>
__device__ static inline uint32_t get_sreg_tid_y();

tid.z#

// mov.u32 sreg_value, %%tid.z; // PTX ISA 20
template <typename = void>
__device__ static inline uint32_t get_sreg_tid_z();

ntid.x#

// mov.u32 sreg_value, %%ntid.x; // PTX ISA 20
template <typename = void>
__device__ static inline uint32_t get_sreg_ntid_x();

ntid.y#

// mov.u32 sreg_value, %%ntid.y; // PTX ISA 20
template <typename = void>
__device__ static inline uint32_t get_sreg_ntid_y();

ntid.z#

// mov.u32 sreg_value, %%ntid.z; // PTX ISA 20
template <typename = void>
__device__ static inline uint32_t get_sreg_ntid_z();

laneid#

// mov.u32 sreg_value, %%laneid; // PTX ISA 13
template <typename = void>
__device__ static inline uint32_t get_sreg_laneid();

warpid#

// mov.u32 sreg_value, %%warpid; // PTX ISA 13
template <typename = void>
__device__ static inline uint32_t get_sreg_warpid();

nwarpid#

// mov.u32 sreg_value, %%nwarpid; // PTX ISA 20, SM_35
template <typename = void>
__device__ static inline uint32_t get_sreg_nwarpid();

ctaid.x#

// mov.u32 sreg_value, %%ctaid.x; // PTX ISA 20
template <typename = void>
__device__ static inline uint32_t get_sreg_ctaid_x();

ctaid.y#

// mov.u32 sreg_value, %%ctaid.y; // PTX ISA 20
template <typename = void>
__device__ static inline uint32_t get_sreg_ctaid_y();

ctaid.z#

// mov.u32 sreg_value, %%ctaid.z; // PTX ISA 20
template <typename = void>
__device__ static inline uint32_t get_sreg_ctaid_z();

nctaid.x#

// mov.u32 sreg_value, %%nctaid.x; // PTX ISA 20
template <typename = void>
__device__ static inline uint32_t get_sreg_nctaid_x();

nctaid.y#

// mov.u32 sreg_value, %%nctaid.y; // PTX ISA 20
template <typename = void>
__device__ static inline uint32_t get_sreg_nctaid_y();

nctaid.z#

// mov.u32 sreg_value, %%nctaid.z; // PTX ISA 20
template <typename = void>
__device__ static inline uint32_t get_sreg_nctaid_z();

smid#

// mov.u32 sreg_value, %%smid; // PTX ISA 13
template <typename = void>
__device__ static inline uint32_t get_sreg_smid();

nsmid#

// mov.u32 sreg_value, %%nsmid; // PTX ISA 20, SM_35
template <typename = void>
__device__ static inline uint32_t get_sreg_nsmid();

gridid#

// mov.u64 sreg_value, %%gridid; // PTX ISA 30
template <typename = void>
__device__ static inline uint64_t get_sreg_gridid();

is_explicit_cluster#

// mov.pred sreg_value, %%is_explicit_cluster; // PTX ISA 78, SM_90
template <typename = void>
__device__ static inline bool get_sreg_is_explicit_cluster();

clusterid.x#

// mov.u32 sreg_value, %%clusterid.x; // PTX ISA 78, SM_90
template <typename = void>
__device__ static inline uint32_t get_sreg_clusterid_x();

clusterid.y#

// mov.u32 sreg_value, %%clusterid.y; // PTX ISA 78, SM_90
template <typename = void>
__device__ static inline uint32_t get_sreg_clusterid_y();

clusterid.z#

// mov.u32 sreg_value, %%clusterid.z; // PTX ISA 78, SM_90
template <typename = void>
__device__ static inline uint32_t get_sreg_clusterid_z();

nclusterid.x#

// mov.u32 sreg_value, %%nclusterid.x; // PTX ISA 78, SM_90
template <typename = void>
__device__ static inline uint32_t get_sreg_nclusterid_x();

nclusterid.y#

// mov.u32 sreg_value, %%nclusterid.y; // PTX ISA 78, SM_90
template <typename = void>
__device__ static inline uint32_t get_sreg_nclusterid_y();

nclusterid.z#

// mov.u32 sreg_value, %%nclusterid.z; // PTX ISA 78, SM_90
template <typename = void>
__device__ static inline uint32_t get_sreg_nclusterid_z();

cluster_ctaid.x#

// mov.u32 sreg_value, %%cluster_ctaid.x; // PTX ISA 78, SM_90
template <typename = void>
__device__ static inline uint32_t get_sreg_cluster_ctaid_x();

cluster_ctaid.y#

// mov.u32 sreg_value, %%cluster_ctaid.y; // PTX ISA 78, SM_90
template <typename = void>
__device__ static inline uint32_t get_sreg_cluster_ctaid_y();

cluster_ctaid.z#

// mov.u32 sreg_value, %%cluster_ctaid.z; // PTX ISA 78, SM_90
template <typename = void>
__device__ static inline uint32_t get_sreg_cluster_ctaid_z();

cluster_nctaid.x#

// mov.u32 sreg_value, %%cluster_nctaid.x; // PTX ISA 78, SM_90
template <typename = void>
__device__ static inline uint32_t get_sreg_cluster_nctaid_x();

cluster_nctaid.y#

// mov.u32 sreg_value, %%cluster_nctaid.y; // PTX ISA 78, SM_90
template <typename = void>
__device__ static inline uint32_t get_sreg_cluster_nctaid_y();

cluster_nctaid.z#

// mov.u32 sreg_value, %%cluster_nctaid.z; // PTX ISA 78, SM_90
template <typename = void>
__device__ static inline uint32_t get_sreg_cluster_nctaid_z();

cluster_ctarank#

// mov.u32 sreg_value, %%cluster_ctarank; // PTX ISA 78, SM_90
template <typename = void>
__device__ static inline uint32_t get_sreg_cluster_ctarank();

cluster_nctarank#

// mov.u32 sreg_value, %%cluster_nctarank; // PTX ISA 78, SM_90
template <typename = void>
__device__ static inline uint32_t get_sreg_cluster_nctarank();

lanemask_eq#

// mov.u32 sreg_value, %%lanemask_eq; // PTX ISA 20, SM_35
template <typename = void>
__device__ static inline uint32_t get_sreg_lanemask_eq();

lanemask_le#

// mov.u32 sreg_value, %%lanemask_le; // PTX ISA 20, SM_35
template <typename = void>
__device__ static inline uint32_t get_sreg_lanemask_le();

lanemask_lt#

// mov.u32 sreg_value, %%lanemask_lt; // PTX ISA 20, SM_35
template <typename = void>
__device__ static inline uint32_t get_sreg_lanemask_lt();

lanemask_ge#

// mov.u32 sreg_value, %%lanemask_ge; // PTX ISA 20, SM_35
template <typename = void>
__device__ static inline uint32_t get_sreg_lanemask_ge();

lanemask_gt#

// mov.u32 sreg_value, %%lanemask_gt; // PTX ISA 20, SM_35
template <typename = void>
__device__ static inline uint32_t get_sreg_lanemask_gt();

clock#

// mov.u32 sreg_value, %%clock; // PTX ISA 10
template <typename = void>
__device__ static inline uint32_t get_sreg_clock();

clock_hi#

// mov.u32 sreg_value, %%clock_hi; // PTX ISA 50, SM_35
template <typename = void>
__device__ static inline uint32_t get_sreg_clock_hi();

clock64#

// mov.u64 sreg_value, %%clock64; // PTX ISA 20, SM_35
template <typename = void>
__device__ static inline uint64_t get_sreg_clock64();

globaltimer#

// mov.u64 sreg_value, %%globaltimer; // PTX ISA 31, SM_35
template <typename = void>
__device__ static inline uint64_t get_sreg_globaltimer();

globaltimer_lo#

// mov.u32 sreg_value, %%globaltimer_lo; // PTX ISA 31, SM_35
template <typename = void>
__device__ static inline uint32_t get_sreg_globaltimer_lo();

globaltimer_hi#

// mov.u32 sreg_value, %%globaltimer_hi; // PTX ISA 31, SM_35
template <typename = void>
__device__ static inline uint32_t get_sreg_globaltimer_hi();

total_smem_size#

// mov.u32 sreg_value, %%total_smem_size; // PTX ISA 41, SM_35
template <typename = void>
__device__ static inline uint32_t get_sreg_total_smem_size();

aggr_smem_size#

// mov.u32 sreg_value, %%aggr_smem_size; // PTX ISA 81, SM_90
template <typename = void>
__device__ static inline uint32_t get_sreg_aggr_smem_size();

dynamic_smem_size#

// mov.u32 sreg_value, %%dynamic_smem_size; // PTX ISA 41, SM_35
template <typename = void>
__device__ static inline uint32_t get_sreg_dynamic_smem_size();

current_graph_exec#

// mov.u64 sreg_value, %%current_graph_exec; // PTX ISA 80, SM_50
template <typename = void>
__device__ static inline uint64_t get_sreg_current_graph_exec();