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编程场景中，接口形参不携带UB/GM地址空间标识，编译器需先完成入参指针到UB/GM地址空间指针的转换，再完成对应的原子操作实现。
| API名 | 非软仿数据类型 | 软仿数据类型 |
|---|---|---|
| asc_atomic_add | int32_t、uint32_t、int64_t、uint64_t、half、bfloat16_t | float、half2、bfloat16x2_t |
| asc_atomic_and | - | 所有数据类型 |
| asc_atomic_cas | uint32_t、uint64_t | int32_t、int64_t、float、half、bfloat16_t、half2、bfloat16x2_t |
| asc_atomic_exch | uint32_t、uint64_t | int32_t、int64_t、float、half、bfloat16_t、half2、bfloat16x2_t |
| asc_atomic_dec | - | 所有数据类型 |
| asc_atomic_inc | - | 所有数据类型 |
| asc_atomic_max | uint32_t、int32_t、int64_t、uint64_t、half、bfloat16_t | float、half2、bfloat16x2_t |
| asc_atomic_min | uint32_t、int32_t、int64_t、uint64_t、half、bfloat16_t | float、half2、bfloat16x2_t |
| asc_atomic_or | - | 所有数据类型 |
| asc_atomic_sub | - | 所有数据类型 |
| asc_atomic_xor | - | 所有数据类型 |
同步与内存栅栏
内存栅栏接口
| 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类型比较函数
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类型精度转换函数
float类型数学库函数
fp8类型
fp8数据转换函数
half类型
half2类型数学库函数
| API名 | API是否软仿 |
|---|---|
| h2ceil | 是 |
| h2cos | 是 |
| h2exp | 是 |
| h2exp10 | 是 |
| h2exp2 | 是 |
| h2floor | 是 |
| h2log | 是 |
| h2log10 | 是 |
| h2log2 | 是 |
| h2rcp | 是 |
| h2rint | 是 |
| h2rsqrt | 是 |
| h2sin | 是 |
| h2sqrt | 否 |
| h2tanh | 是 |
| h2trunc | 是 |
half2类型比较函数
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是否软仿 |
|---|---|
| asc_dcci_entire | 否 |
| asc_dcci_single | 否 |
| asc_ldca | 否 |
| asc_ldcg | 否 |
| asc_stcg | 否 |
| asc_stwt | 否 |