SIMT接口软仿情况汇总

本节按照SIMT API文档所在目录分级汇总各接口的软仿情况。

非软仿接口是指使用单一硬件指令完成操作的接口；软仿接口是指由编译器实现，或由上层接口通过拼接多条硬件指令实现功能的接口。

Warp函数

LaneID类函数

API名	API是否软仿
laneid	否
lanemask_eq	否
lanemask_ge	否
lanemask_gt	否
lanemask_le	否
lanemask_lt	否

Warp Reduce类函数

API名	API是否软仿
asc_reduce_add	否
asc_reduce_max	否
asc_reduce_min	否

Warp Shfl类函数

API名	API是否软仿
asc_shfl	否
asc_shfl_down	否
asc_shfl_up	否
asc_shfl_xor	否

Warp Vote类函数

API名	API是否软仿
asc_activemask	否
asc_all	否
asc_any	否
asc_ballot	否

协作组

API名	API是否软仿
binary_partition	是
tiled_partition	是

coalesced_group

API名	API是否软仿
all	是
any	是
ballot	是
coalesced_group构造函数	是
meta_group_rank	是
meta_group_size	是
num_threads	是
shfl	是
shfl_down	是
shfl_up	是
size	是
sync	是
thread_rank	是

thread_block

API名	API是否软仿
dim_threads	是
group_dim	是
group_index	是
num_threads	是
size	是
sync	是
thread_block构造函数	是
thread_index	是
thread_rank	是

thread_block_tile

API名	API是否软仿
all	是
any	是
ballot	是
meta_group_rank	是
meta_group_size	是
num_threads	是
shfl	是
shfl_down	是
shfl_up	是
shfl_xor	是
size	是
sync	是
thread_block_tile构造函数	是
thread_rank	是

原子操作

下表汇总原子操作接口在SIMD与SIMT编程场景下的软仿情况。SIMT编程场景中，接口形参不携带Unified Buffer（UB）/GM地址空间标识，编译器需先完成入参指针到UB/GM地址空间指针的转换，再完成对应的原子操作实现。是否使用接口返回值会影响部分接口的指令生成，接口软仿情况分“使用返回值”与“不使用返回值”两张表列出。

表1 原子操作（使用返回值）

API名	非软仿数据类型	软仿数据类型
asc_atomic_add	int32_t、uint32_t、int64_t、uint64_t	float、half2、bfloat16x2_t
asc_atomic_and	-	int32_t、uint32_t、int64_t、uint64_t
asc_atomic_cas	uint32_t、uint64_t	int32_t、int64_t、float、half2、bfloat16x2_t
asc_atomic_dec	-	uint32_t、uint64_t
asc_atomic_exch	uint32_t、uint64_t	int32_t、int64_t、float、half2、bfloat16x2_t
asc_atomic_inc	-	uint32_t、uint64_t
asc_atomic_max	uint32_t、int32_t、int64_t、uint64_t	float、half2、bfloat16x2_t
asc_atomic_min	uint32_t、int32_t、int64_t、uint64_t	float、half2、bfloat16x2_t
asc_atomic_or	-	int32_t、uint32_t、int64_t、uint64_t
asc_atomic_sub	-	int32_t、uint32_t、float、int64_t、uint64_t、half2、bfloat16x2_t
asc_atomic_xor	-	int32_t、uint32_t、int64_t、uint64_t

表2 原子操作（不使用返回值）

API名	非软仿数据类型	软仿数据类型
asc_atomic_add	int32_t、uint32_t、int64_t、uint64_t、float、half、bfloat16_t、half2、bfloat16x2_t	-
asc_atomic_and	-	int32_t、uint32_t、int64_t、uint64_t
asc_atomic_cas	uint32_t、uint64_t	int32_t、int64_t、float、half2、bfloat16x2_t
asc_atomic_dec	-	uint32_t、uint64_t
asc_atomic_exch	uint32_t、uint64_t	int32_t、int64_t、float、half2、bfloat16x2_t
asc_atomic_inc	-	uint32_t、uint64_t
asc_atomic_max	int32_t、uint32_t、int64_t、uint64_t、float、half、bfloat16_t、half2、bfloat16x2_t	-
asc_atomic_min	int32_t、uint32_t、int64_t、uint64_t、float、half、bfloat16_t、half2、bfloat16x2_t	-
asc_atomic_or	-	int32_t、uint32_t、int64_t、uint64_t
asc_atomic_sub	half、bfloat16_t	int32_t、uint32_t、float、int64_t、uint64_t、half2、bfloat16x2_t
asc_atomic_xor	-	int32_t、uint32_t、int64_t、uint64_t

同步与内存栅栏

内存栅栏接口

API名	API是否软仿
asc_threadfence	否
asc_threadfence_block	否

同步接口

API名	API是否软仿
asc_syncthreads	否

地址空间谓词函数

API名	API是否软仿
__isGlobal	否
__isLocal	否
__isUbuf	否

地址空间转换函数

API名	API是否软仿
__cvta_generic_to_global	否
__cvta_generic_to_local	否
__cvta_generic_to_ubuf	否
__cvta_global_to_generic	否
__cvta_local_to_generic	否
__cvta_ubuf_to_generic	否

数学函数

bfloat16类型

bfloat16x2类型数学库函数

API名	API是否软仿
h2ceil	否
h2cos	是
h2exp	是
h2exp10	是
h2exp2	是
h2floor	否
h2log	是
h2log10	是
h2log2	是
h2rcp	是
h2rint	否
h2rsqrt	是
h2sin	是
h2sqrt	是
h2tanh	是
h2trunc	是

bfloat16x2类型比较函数

API名	API是否软仿
__hbequx2	是
__hbeqx2	是
__hbgeux2	是
__hbgex2	是
__hbgtux2	是
__hbgtx2	是
__hbleux2	是
__hblex2	是
__hbltux2	是
__hbltx2	是
__hbneux2	是
__hbnex2	是
__hequx2	是
__hequx2_mask	是
__heqx2	是
__heqx2_mask	是
__hgeux2	是
__hgeux2_mask	是
__hgex2	是
__hgex2_mask	是
__hgtux2	是
__hgtux2_mask	是
__hgtx2	是
__hgtx2_mask	是
__hleux2	是
__hleux2_mask	是
__hlex2	是
__hlex2_mask	是
__hltux2	是
__hltux2_mask	是
__hltx2	是
__hltx2_mask	是
__hmaxx2	否
__hmaxx2_nan	是
__hminx2	否
__hminx2_nan	是
__hneux2	是
__hneux2_mask	是
__hnex2	是
__hnex2_mask	是
__isnanx2	是

bfloat16x2类型算术函数

API名	API是否软仿
__habsx2	否
__haddx2	否
__hcmadd	是
__hdivx2	否
__hfmax2	否
__hfmax2_relu	是
__hmulx2	否
__hnegx2	否
__hsubx2	否

bfloat16类型数学库函数

API名	API是否软仿
hceil	否
hcos	是
hexp	是
hexp10	是
hexp2	是
hfloor	否
hlog	是
hlog10	是
hlog2	是
hrcp	是
hrint	否
hrsqrt	是
hsin	是
hsqrt	是
htanh	是
htrunc	是

bfloat16类型比较函数

API名	API是否软仿
__heq	否
__hequ	否
__hge	否
__hgeu	否
__hgt	否
__hgtu	否
__hisinf	否
__hisnan	否
__hle	否
__hleu	否
__hlt	否
__hltu	否
__hmax	是
__hmax_nan	是
__hmin	是
__hmin_nan	是
__hne	否
__hneu	否

bfloat16类型算术函数

API名	API是否软仿
__habs	是
__hadd	否
__hdiv	否
__hfma	否
__hfma_relu	是
__hmul	否
__hneg	否
__hsub	否

bfloat16类型精度转换函数

API名	API是否软仿
__bfloat1622float2	是
__bfloat162bfloat162	是
__bfloat162bfloat16_rd	否
__bfloat162bfloat16_rn	否
__bfloat162bfloat16_rna	否
__bfloat162bfloat16_ru	否
__bfloat162bfloat16_rz	否
__bfloat162float	是
__bfloat162int_rd	否
__bfloat162int_rn	否
__bfloat162int_rna	否
__bfloat162int_ru	否
__bfloat162int_rz	否
__bfloat162ll_rd	是
__bfloat162ll_rn	是
__bfloat162ll_rna	是
__bfloat162ll_ru	是
__bfloat162ll_rz	是
__bfloat162uint_rd	否
__bfloat162uint_rn	否
__bfloat162uint_rna	否
__bfloat162uint_ru	否
__bfloat162uint_rz	否
__bfloat162ull_rd	是
__bfloat162ull_rn	是
__bfloat162ull_rna	是
__bfloat162ull_ru	是
__bfloat162ull_rz	是
__float22bfloat162_rd	否
__float22bfloat162_rd_sat	否
__float22bfloat162_rn	否
__float22bfloat162_rn_sat	否
__float22bfloat162_rna	否
__float22bfloat162_rna_sat	否
__float22bfloat162_ru	否
__float22bfloat162_ru_sat	否
__float22bfloat162_rz	否
__float22bfloat162_rz_sat	否
__float2bfloat16	否
__float2bfloat162_rn	是
__float2bfloat16_rd	否
__float2bfloat16_rd_sat	否
__float2bfloat16_rn	否
__float2bfloat16_rn_sat	否
__float2bfloat16_rna	否
__float2bfloat16_rna_sat	否
__float2bfloat16_ru	否
__float2bfloat16_ru_sat	否
__float2bfloat16_rz	否
__float2bfloat16_rz_sat	否
__floats2bfloat162_rn	是
__half2bfloat16_rd	否
__half2bfloat16_rn	否
__half2bfloat16_rna	否
__half2bfloat16_ru	否
__half2bfloat16_rz	否
__halves2bfloat162	是
__high2bfloat16	否
__high2bfloat162	是
__high2float	否
__highs2bfloat162	是
__int2bfloat16_rd	否
__int2bfloat16_rn	否
__int2bfloat16_rna	否
__int2bfloat16_ru	否
__int2bfloat16_rz	否
__ll2bfloat16_rd	是
__ll2bfloat16_rn	是
__ll2bfloat16_rna	是
__ll2bfloat16_ru	是
__ll2bfloat16_rz	是
__low2bfloat16	否
__low2bfloat162	是
__low2float	否
__lowhigh2highlow	是
__lows2bfloat162	是
__uint2bfloat16_rd	否
__uint2bfloat16_rn	否
__uint2bfloat16_rna	否
__uint2bfloat16_ru	否
__uint2bfloat16_rz	否
__ull2bfloat16_rd	是
__ull2bfloat16_rn	是
__ull2bfloat16_rna	是
__ull2bfloat16_ru	是
__ull2bfloat16_rz	是
__ushort_as_bfloat16	是

float类型数学库函数

API名	API是否软仿
__fdividef	否
__saturatef	是
acosf	是
acoshf	是
asinf	是
asinhf	是
atan2f	是
atanf	是
atanhf	是
cbrtf	是
ceilf	否
copysignf	是
cosf	是
coshf	是
cospif	是
cyl_bessel_i0f	是
cyl_bessel_i1f	是
erfcf	是
erfcinvf	是
erfcxf	是
erff	是
erfinvf	是
exp10f	否
exp2f	否
expf	否
expm1f	是
fabsf	否
fdimf	是
fdividef	是
floorf	否
fmaf	否
fmaxf	是
fminf	是
fmodf	是
frexpf	是
hypotf	是
ilogbf	是
isfinite	否
isinf	否
isnan	否
j0f	是
j1f	是
jnf	是
ldexpf	是
lgammaf	是
llrintf	是
llroundf	是
log10f	是
log1pf	是
log2f	是
logbf	是
logf	是
lrintf	是
lroundf	是
modff	是
nearbyintf	是
nextafterf	是
norm3df	是
norm4df	是
normcdff	是
normcdfinvf	是
normf	是
powf	否
rcbrtf	是
remainderf	是
remquof	是
rhypotf	是
rintf	是
rnorm3df	是
rnorm4df	是
rnormf	是
roundf	否
rsqrtf	是
scalblnf	是
scalbnf	是
signbit	是
sincosf	是
sincospif	是
sinf	是
sinhf	是
sinpif	是
sqrtf	否
tanf	是
tanhf	是
tanpif	是
tgammaf	是
truncf	否
y0f	是
y1f	是
ynf	是

fp8类型

fp8数据转换函数

API名	API是否软仿
__asc_cvt_float2_to_fp8x2	是
__e4m3x22float2	否
__e5m2x22float2	否
__float22hif82_rh	否
__float22hif82_rh_sat	否
__float22hif82_rna	否
__float22hif82_rna_sat	否
__half22hif82_rh	否
__half22hif82_rh_sat	否
__half22hif82_rna	否
__half22hif82_rna_sat	否
__hif822float2	否
__hif822half2	否

half类型

half2类型数学库函数

API名	API是否软仿
h2ceil	是
h2cos	是
h2exp	是
h2exp10	是
h2exp2	是
h2floor	是
h2log	是
h2log10	是
h2log2	是
h2rcp	是
h2rint	是
h2rsqrt	是
h2sin	是
h2sqrt	否
h2tanh	是
h2trunc	是

half2类型比较函数

API名	API是否软仿
__hbequx2	否
__hbeqx2	否
__hbgeux2	否
__hbgex2	否
__hbgtux2	否
__hbgtx2	否
__hbleux2	否
__hblex2	否
__hbltux2	否
__hbltx2	否
__hbneux2	否
__hbnex2	否
__hequx2	否
__hequx2_mask	否
__heqx2	否
__heqx2_mask	否
__hgeux2	否
__hgeux2_mask	否
__hgex2	否
__hgex2_mask	否
__hgtux2	否
__hgtux2_mask	否
__hgtx2	否
__hgtx2_mask	否
__hleux2	否
__hleux2_mask	否
__hlex2	否
__hlex2_mask	否
__hltux2	否
__hltux2_mask	否
__hltx2	否
__hltx2_mask	否
__hmaxx2	否
__hmaxx2_nan	是
__hminx2	否
__hminx2_nan	是
__hneux2	否
__hneux2_mask	否
__hnex2	否
__hnex2_mask	否
__isnanx2	否

half2类型算术函数

API名	API是否软仿
__habsx2	否
__haddx2	否
__hcmadd	是
__hdivx2	否
__hfmax2	否
__hfmax2_relu	是
__hmulx2	否
__hnegx2	否
__hsubx2	否

half类型数学库函数

API名	API是否软仿
hceil	是
hcos	是
hexp	否
hexp10	是
hexp2	是
hfloor	是
hlog	否
hlog10	是
hlog2	是
hrcp	是
hrint	否
hrsqrt	是
hsin	是
hsqrt	否
htanh	是
htrunc	是

half类型比较函数

API名	API是否软仿
__heq	否
__hequ	否
__hge	否
__hgeu	否
__hgt	否
__hgtu	否
__hisinf	否
__hisnan	否
__hle	否
__hleu	否
__hlt	否
__hltu	否
__hmax	是
__hmax_nan	是
__hmin	是
__hmin_nan	是
__hne	否
__hneu	否

half类型算术函数

API名	API是否软仿
__habs	是
__hadd	否
__hdiv	否
__hfma	否
__hfma_relu	是
__hmul	否
__hneg	否
__hsub	否

half类型精度转换函数

API名	API是否软仿
__bfloat162half_rd	否
__bfloat162half_rd_sat	否
__bfloat162half_rn	否
__bfloat162half_rn_sat	否
__bfloat162half_rna	否
__bfloat162half_rna_sat	否
__bfloat162half_ru	否
__bfloat162half_ru_sat	否
__bfloat162half_rz	否
__bfloat162half_rz_sat	否
__float22half2_rd	否
__float22half2_rd_sat	否
__float22half2_rn	否
__float22half2_rn_sat	否
__float22half2_rna	否
__float22half2_rna_sat	否
__float22half2_ro	否
__float22half2_ro_sat	否
__float22half2_ru	否
__float22half2_ru_sat	否
__float22half2_rz	否
__float22half2_rz_sat	否
__float2half	否
__float2half_rd	否
__float2half_rd_sat	否
__float2half_rn	否
__float2half_rn_sat	否
__float2half_rna	否
__float2half_rna_sat	否
__float2half_ro	否
__float2half_ro_sat	否
__float2half_ru	是
__float2half_ru_sat	否
__float2half_rz	否
__float2half_rz_sat	否
__floats2half2_rn	是
__half22float2	否
__half2float	是
__half2half_rd	否
__half2half_rn	否
__half2half_rna	否
__half2half_ru	否
__half2half_rz	否
__half2int_rd	否
__half2int_rn	否
__half2int_rna	否
__half2int_ru	否
__half2int_rz	否
__half2ll_rd	是
__half2ll_rn	是
__half2ll_rna	是
__half2ll_ru	是
__half2ll_rz	是
__half2uint_rd	否
__half2uint_rn	否
__half2uint_rna	否
__half2uint_ru	否
__half2uint_rz	否
__half2ull_rd	是
__half2ull_rn	是
__half2ull_rna	是
__half2ull_ru	是
__half2ull_rz	是
__halves2half2	是
__high2float	否
__high2half	否
__high2half2	是
__highs2half2	是
__int2half_rd	否
__int2half_rd_sat	否
__int2half_rn	否
__int2half_rn_sat	否
__int2half_rna	否
__int2half_rna_sat	否
__int2half_ru	否
__int2half_ru_sat	否
__int2half_rz	否
__int2half_rz_sat	否
__ll2half_rd	是
__ll2half_rn	是
__ll2half_rna	是
__ll2half_ru	是
__ll2half_rz	是
__low2float	否
__low2half	是
__low2half2	是
__lowhigh2highlow	是
__lows2half2	是
__uint2half_rd	否
__uint2half_rd_sat	否
__uint2half_rn	否
__uint2half_rn_sat	否
__uint2half_rna	否
__uint2half_rna_sat	否
__uint2half_ru	否
__uint2half_ru_sat	否
__uint2half_rz	否
__uint2half_rz_sat	否
__ull2half_rd	是
__ull2half_rn	是
__ull2half_rna	是
__ull2half_ru	是
__ull2half_rz	是
__ushort_as_half	是

数据类型转换

类型转换函数

API名	API是否软仿
__float2float_rd	否
__float2float_rn	否
__float2float_rna	否
__float2float_ru	否
__float2float_rz	否
__float2int_rd	否
__float2int_rn	否
__float2int_rna	否
__float2int_ru	否
__float2int_rz	否
__float2ll_rd	否
__float2ll_rn	否
__float2ll_rna	否
__float2ll_ru	否
__float2ll_rz	否
__float2uint_rd	否
__float2uint_rn	否
__float2uint_rna	否
__float2uint_ru	否
__float2uint_rz	否
__float2ull_rd	否
__float2ull_rn	否
__float2ull_rna	否
__float2ull_ru	否
__float2ull_rz	否
__float_as_int	是
__float_as_uint	是
__int2float_rd	否
__int2float_rn	否
__int2float_rna	否
__int2float_ru	否
__int2float_rz	否
__int_as_float	是
__ll2float_rd	否
__ll2float_rn	否
__ll2float_rna	否
__ll2float_ru	否
__ll2float_rz	否
__uint2float_rd	否
__uint2float_rn	否
__uint2float_rna	否
__uint2float_ru	否
__uint2float_rz	否
__uint_as_float	是
__ull2float_rd	否
__ull2float_rn	否
__ull2float_rna	否
__ull2float_ru	否
__ull2float_rz	否

整型数学库函数

API名	API是否软仿
__brev	否
__byte_perm	否
__clz	否
__ffs	否
__hadd	否
__mul24	否
__mul64hi	否
__mul_i32toi64	否
__mulhi	否
__popc	否
__rhadd	否
__sad	否
__uhadd	否
__umul24	否
__umul64hi	否
__umulhi	否
__urhadd	否
__usad	否
labs	是
llabs	是
llmax	是
llmin	是
max	否
min	否
ullmax	是
ullmin	是
umax	否
umin	否

访存函数

API名	API是否软仿
asc_dcci_entire	否
asc_dcci_single	否
asc_ldca	否
asc_ldcg	否
asc_stcg	否
asc_stwt	否