Ascend C
搜索

Appearance

本页内容

API列表

同步与内存栅栏​

表1 同步接口

接口名功能描述
asc_syncthreads等待当前thread block内所有thread代码都执行到该函数位置。

表2 内存栅栏接口

接口名功能描述
asc_threadfence用于保证不同核对同一份全局、共享内存的访问过程中，写入操作的时序性。
asc_threadfence_block用于协调同一线程块（Thread Block）内线程之间的内存操作顺序，确保某一线程在调用asc_threadfence_block()之前的所有内存读写操作对同一线程块内的其他线程可见。

原子操作​

表3 原子操作

接口名功能描述
asc_atomic_add对Unified Buffer或Global Memory上的数据与指定数据执行原子加操作，即将指定数据累加到Unified Buffer或Global Memory的数据中。
asc_atomic_sub对Unified Buffer或Global Memory上的数据与指定数据执行原子减操作，即在Unified Buffer或Global Memory的数据上减去指定数据。
asc_atomic_exch对Unified Buffer或Global Memory地址做原子赋值操作，即将指定数据赋值到Unified Buffer或Global Memory地址中。
asc_atomic_max对Unified Buffer或Global Memory数据做原子求最大值操作，即将Unified Buffer或Global Memory的数据与指定数据中的最大值赋值到Unified Buffer或Global Memory地址中。
asc_atomic_min对Unified Buffer或Global Memory数据做原子求最小值操作，即将Unified Buffer或Global Memory的数据与指定数据中的最小值赋值到Unified Buffer或Global Memory地址中。
asc_atomic_inc对Unified Buffer或Global Memory上address的数值进行原子加1操作，如果address上的数值大于等于指定数值val，则对address赋值为0，否则将address上数值加1。
asc_atomic_dec对Unified Buffer或Global Memory上address的数值进行原子减1操作，如果address上的数值等于0或大于指定数值val，则对address赋值为val，否则将address上数值减1。
asc_atomic_cas对Unified Buffer或Global Memory上address的数值进行原子比较赋值操作，如果address上的数值等于指定数值compare，则对address赋值为指定数值val，否则address的数值不变。
asc_atomic_and对Unified Buffer或Global Memory上address的数值与指定数值val进行原子与（&）操作，即将address数值与（&）val的结果赋值到Unified Buffer或Global Memory上。
asc_atomic_or对Unified Buffer或Global Memory上address的数值与指定数值val进行原子或（|）操作，即将address数值或（|）val的结果赋值到Unified Buffer或Global Memory上。
asc_atomic_xor对Unified Buffer或Global Memory上address的数值与指定数值val进行原子异或（^）操作，即将address数值异或（^）val的结果赋值到Unified Buffer或Global Memory上。

Warp函数​

表4 Warp Vote类函数

接口名功能描述
asc_all判断是否所有活跃线程的输入均不为0。
asc_any判断是否有活跃线程的输入不为0。
asc_ballot判断Warp内每个活跃线程的输入是否不为0。
asc_activemask查看Warp内所有线程是否为活跃状态。

表5 Warp Shfl类函数

接口名功能描述
asc_shfl获取Warp内指定线程srcLane输入的用于交换的var值。
asc_shfl_up获取Warp内当前线程向前偏移delta（当前线程Lane ID - delta）的线程输入的用于交换的var值。
asc_shfl_down获取Warp内当前线程向后偏移delta（当前线程Lane ID + delta）的线程输入的用于交换的var值。
asc_shfl_xor获取Warp内当前线程Lane ID与输入laneMask做异或操作（Lane ID ^ laneMask）得到的dstLaneId对应线程输入的用于交换的var值。

表6 Warp Reduce类函数

接口名功能描述
asc_reduce_add对Warp内所有活跃线程输入的val求和。
asc_reduce_max对Warp内所有活跃线程输入的val求最大值。
asc_reduce_min对Warp内所有活跃线程输入val求最小值。

表7 Lane ID类函数

接口名功能描述
laneid获取线程在其线程束内的索引。
lanemask_eq提供当前线程的一个32位掩码，在当前线程所属的Warp中，只有当前线程所在的lane位被置为1，其余位为0。
lanemask_le提供当前线程的一个32位掩码，在当前线程所属的Warp中，将“Lane ID小于等于当前线程”的线程的对应位设为1，其余位为0。
lanemask_lt提供当前线程的一个32位掩码，在当前线程所属的Warp中，将“Lane ID严格小于当前线程”的线程的对应位设为1，其余位为0。
lanemask_ge提供当前线程的一个32位掩码，在当前线程所属的Warp中，将“Lane ID大于等于当前线程”的线程的对应位设为1，其余位为0。
lanemask_gt提供当前线程的一个32位掩码，在当前线程所属的Warp中，将“Lane ID严格大于当前线程”的线程的对应位设为1，其余位为0。

数学函数​

表8 fp8数据转换函数

接口名功能描述
__float22hif82_rna获取输入的两个分量遵循CAST_ROUND模式转换成的hifloat8x2_t类型数据。
__float22hif82_rna_sat饱和模式下获取输入的两个分量遵循CAST_ROUND模式转换成的hifloat8x2_t类型数据。
__float22hif82_rh获取输入的两个分量遵循CAST_HYBRID模式转换成的hifloat8x2_t类型数据。
__float22hif82_rh_sat饱和模式下获取输入的两个分量遵循CAST_HYBRID模式转换成的hifloat8x2_t类型数据。
__asc_cvt_float2_to_fp8x2输入的两个分量遵循CAST_RINT模式，根据指定的8位浮点数类型和指定的饱和模式，转换成__asc_fp8x2_storage_t类型数据。
__half22hif82_rna获取输入的两个分量遵循CAST_ROUND模式转换成的hifloat8x2_t类型数据。
__half22hif82_rna_sat饱和模式下获取输入的两个分量遵循CAST_ROUND模式转换成的hifloat8x2_t类型数据。
__half22hif82_rh获取输入的两个分量遵循CAST_HYBRID模式转换成的hifloat8x2_t类型数据。
__half22hif82_rh_sat饱和模式下获取输入的两个分量遵循CAST_HYBRID模式转换成的hifloat8x2_t类型数据。
__hif822float2获取hifloat8x2_t类型输入转换成的float2类型数据。
__hif822half2获取hifloat8x2_t类型输入转换成的half2类型数据。
__e4m3x22float2获取float8_e4m3x2_t类型输入转换成的float2类型数据。
__e5m2x22float2获取float8_e5m2x2_t类型输入转换成的float2类型数据。

表9 half类型算术函数

接口名功能描述
__habs获取输入数据的绝对值。
__hfma对输入数据x、y、z，计算x与y相乘加上z的结果。
__hadd计算两个half类型数据的相加结果，并遵循CAST_RINT模式舍入。
__hsub计算两个half类型数据的相减结果，并遵循CAST_RINT模式舍入。
__hmul计算两个half类型数据的相乘结果，并遵循CAST_RINT模式舍入。
__hdiv计算两个half类型数据的相除结果，并遵循CAST_RINT模式舍入。
__hneg获取输入half类型数据的负值。
__hfma_relu对输入half类型数据x、y、z，计算x与y相乘加上z的结果，并遵循CAST_RINT模式舍入。负数结果置为0。

表10 half类型比较函数

接口名功能描述
__hmax获取两个输入数据中的最大值。
__hmin获取两个输入数据中的最小值。
__hisnan判断浮点数是否为nan。
__hisinf判断浮点数是否为无穷。
__heq比较两个half类型数据是否相等，相等时返回true。
__hne比较两个half类型数据是否不相等，不相等时返回true。
__hle比较两个half类型数据，仅当第一个数小于或等于第二个数时返回true。
__hge比较两个half类型数据，仅当第一个数大于或等于第二个数时返回true。
__hlt比较两个half类型数据，仅当第一个数小于第二个数时返回true。
__hgt比较两个half类型数据，仅当第一个数大于第二个数时返回true。
__hequ比较两个half类型数据是否相等，相等时返回true。若任一输入为nan，返回true。
__hneu比较两个half类型数据是否不相等，不相等时返回true。若任一输入为nan，返回true。
__hleu比较两个half类型数据，当第一个数小于或等于第二个数时返回true。若任一输入为nan，返回true。
__hgeu比较两个half类型数据，当第一个数大于或等于第二个数时返回true。若任一输入为nan，返回true。
__hltu比较两个half类型数据，当第一个数小于第二个数时返回true。若任一输入为nan，返回true。
__hgtu比较两个half类型数据，当第一个数大于第二个数时返回true。若任一输入为nan，返回true。
__hmax_nan获取两个输入数据中的最大值。任一输入为nan时返回nan。
__hmin_nan获取两个输入数据中的最小值。任一输入为nan时返回nan。

表11 half类型数学库函数

接口名功能描述
htanh获取输入数据的三角函数双曲正切值。
hexp指定输入x，获取e的x次方。
hexp2指定输入x，获取2的x次方。
hexp10指定输入x，获取10的x次方。
hlog获取以e为底，输入数据的对数。
hlog2获取以2为底，输入数据的对数。
hlog10获取以10为底，输入数据的对数。
hcos获取输入数据的三角函数余弦值。
hsin获取输入数据的三角函数正弦值。
hsqrt获取输入数据x的平方根。
hrsqrt获取输入数据x的平方根的倒数。
hrcp获取输入数据x的倒数。
hrint获取与输入数据最接近的整数，若存在两个同样接近的整数，则获取其中的偶数。
hfloor获取小于或等于输入数据的最大整数值。
hceil获取大于或等于输入数据的最小整数值。
htrunc获取对输入数据的浮点数截断后的整数。

表12 half类型精度转换函数

接口名功能描述
__float2half获取输入遵循CAST_RINT模式转换成的半精度浮点数。
__float2half_rn获取输入遵循CAST_RINT模式转换成的半精度浮点数。
__float2half_rn_sat饱和模式下获取输入遵循CAST_RINT模式转换成的半精度浮点数。
__float22half2_rn_sat饱和模式下获取输入的两个分量遵循CAST_RINT模式转换成的half2类型数据。
__float2half_rz获取输入遵循CAST_TRUNC模式转换成的半精度浮点数。
__float2half_rz_sat饱和模式下获取输入遵循CAST_TRUNC模式转换成的半精度浮点数。
__float22half2_rz获取输入的两个分量遵循CAST_TRUNC模式转换成的half2类型数据。
__float22half2_rz_sat饱和模式下获取输入的两个分量遵循CAST_TRUNC模式转换成的half2类型数据。
__float2half_rd获取输入遵循CAST_FLOOR模式转换成的半精度浮点数。
__float2half_rd_sat饱和模式下获取输入遵循CAST_FLOOR模式转换成的半精度浮点数。
__float22half2_rd获取输入的两个分量遵循CAST_FLOOR模式转换成的half2类型数据。
__float22half2_rd_sat饱和模式下获取输入的两个分量遵循CAST_FLOOR模式转换成的half2类型数据。
__float2half_ru获取输入遵循CAST_CEIL模式转换成的半精度浮点数。
__float2half_ru_sat饱和模式下获取输入遵循CAST_CEIL模式转换成的半精度浮点数。
__float22half2_ru获取输入的两个分量遵循CAST_CEIL模式转换成的half2类型数据。
__float22half2_ru_sat饱和模式下获取输入的两个分量遵循CAST_CEIL模式转换成的half2类型数据。
__float2half_rna获取输入遵循CAST_ROUND模式转换成的半精度浮点数。
__float2half_rna_sat饱和模式下获取输入遵循CAST_ROUND模式转换成的半精度浮点数。
__float22half2_rna获取输入的两个分量遵循CAST_ROUND模式转换成的half2类型数据。
__float22half2_rna_sat饱和模式下获取输入的两个分量遵循CAST_ROUND模式转换成的half2类型数据。
__float2half_ro获取输入遵循CAST_ODD模式转换成的半精度浮点数。
__float2half_ro_sat饱和模式下获取输入遵循CAST_ODD模式转换成的半精度浮点数。
__float22half2_ro获取输入的两个分量遵循CAST_ODD模式转换成的half2类型数据。
__float22half2_ro_sat饱和模式下获取输入的两个分量遵循CAST_ODD模式转换成的half2类型数据。
__half2float获取输入转换成的浮点数。
__half2half_rn获取输入遵循CAST_RINT模式取整后的half类型数据。
__half2half_rz获取输入遵循CAST_TRUNC模式取整后的half类型数据。
__half2half_rd获取输入遵循CAST_FLOOR模式取整后的half类型数据。
__half2half_ru获取输入遵循CAST_CEIL模式取整后的half类型数据。
__half2half_rna获取输入遵循CAST_ROUND模式取整后的half类型数据。
__half2uint_rn获取输入遵循CAST_RINT模式转换成的无符号整数。
__half2uint_rz获取输入遵循CAST_TRUNC模式转换成的无符号整数。
__half2uint_rd获取输入遵循CAST_FLOOR模式转换成的无符号整数。
__half2uint_ru获取输入遵循CAST_CEIL模式转换成的无符号整数。
__half2uint_rna获取输入遵循CAST_ROUND模式转换成的无符号整数。
__half2int_rn获取输入遵循CAST_RINT模式转换成的有符号整数。
__half2int_rz获取输入遵循CAST_TRUNC模式转换成的有符号整数。
__half2int_rd获取输入遵循CAST_FLOOR模式转换成的有符号整数。
__half2int_ru获取输入遵循CAST_CEIL模式转换成的有符号整数。
__half2int_rna获取输入遵循CAST_ROUND模式转换成的有符号整数。
__half2ull_rn获取输入遵循CAST_RINT模式转换成的64位无符号整数。
__half2ull_rz获取输入遵循CAST_TRUNC模式转换成的64位无符号整数。
__half2ull_rd获取输入遵循CAST_FLOOR模式转换成的64位无符号整数。
__half2ull_ru获取输入遵循CAST_CEIL模式转换成的64位无符号整数。
__half2ull_rna获取输入遵循CAST_ROUND模式转换成的64位无符号整数。
__half2ll_rn获取输入遵循CAST_RINT模式转换成的64位有符号整数。
__half2ll_rz获取输入遵循CAST_TRUNC模式转换成的64位有符号整数。
__half2ll_rd获取输入遵循CAST_FLOOR模式转换成的64位有符号整数。
__half2ll_ru获取输入遵循CAST_CEIL模式转换成的64位有符号整数。
__half2ll_rna获取输入遵循CAST_ROUND模式转换成的64位有符号整数。
__bfloat162half_rn获取输入遵循CAST_RINT模式转换成的half类型数据。
__bfloat162half_rn_sat饱和模式下获取输入遵循CAST_RINT模式转换成的half类型数据。
__bfloat162half_rz获取输入遵循CAST_TRUNC模式转换成的half类型数据。
__bfloat162half_rz_sat饱和模式下获取输入遵循CAST_TRUNC模式转换成的half类型数据。
__bfloat162half_rd获取输入遵循CAST_FLOOR模式转换成的half类型数据。
__bfloat162half_rd_sat饱和模式下获取输入遵循CAST_FLOOR模式转换成的half类型数据。
__bfloat162half_ru获取输入遵循CAST_CEIL模式转换成的half类型数据。
__bfloat162half_ru_sat饱和模式下获取输入遵循CAST_CEIL模式转换成的half类型数据。
__bfloat162half_rna获取输入遵循CAST_ROUND模式转换成的half类型数据。
__bfloat162half_rna_sat饱和模式下获取输入遵循CAST_ROUND模式转换成的half类型数据。
__uint2half_rn获取输入遵循CAST_RINT模式转换成的half类型数据。
__uint2half_rn_sat饱和模式下获取输入的uint32数据转换成的half数据，并遵循CAST_RINT模式。
__uint2half_rz获取输入遵循CAST_TRUNC模式转换成的half类型数据。
__uint2half_rz_sat饱和模式下获取输入的uint32数据转换成的half数据，并遵循CAST_TRUNC模式。
__uint2half_rd获取输入遵循CAST_FLOOR模式转换成的half类型数据。
__uint2half_rd_sat饱和模式下获取输入的uint32数据转换成的half数据，并遵循CAST_FLOOR模式。
__uint2half_ru获取输入遵循CAST_CEIL模式转换成的half类型数据。
__uint2half_ru_sat饱和模式下获取输入的uint32数据转换成的half数据，并遵循CAST_CEIL模式。
__uint2half_rna获取输入遵循CAST_ROUND模式转换成的half类型数据。
__uint2half_rna_sat饱和模式下获取输入的uint32数据转换成的half数据，并遵循CAST_ROUND模式。
__int2half_rn获取输入遵循CAST_RINT模式转换成的half类型数据。
__int2half_rn_sat饱和模式下获取输入的int32数据转换成的half数据，并遵循CAST_RINT模式。
__int2half_rz获取输入遵循CAST_TRUNC模式转换成的half类型数据。
__int2half_rz_sat饱和模式下获取输入的int32数据转换成的half数据，并遵循CAST_TRUNC模式。
__int2half_rd获取输入遵循CAST_FLOOR模式转换成的half类型数据。
__int2half_rd_sat饱和模式下获取输入的int32数据转换成的half数据，并遵循CAST_FLOOR模式。
__int2half_ru获取输入遵循CAST_CEIL模式转换成的half类型数据。
__int2half_ru_sat饱和模式下获取输入的int32数据转换成的half数据，并遵循CAST_CEIL模式。
__int2half_rna获取输入遵循CAST_ROUND模式转换成的half类型数据。
__int2half_rna_sat饱和模式下获取输入的int32数据转换成的half数据，并遵循CAST_ROUND模式。
__ull2half_rn获取输入遵循CAST_RINT模式转换成的half类型数据。
__ull2half_rz获取输入遵循CAST_TRUNC模式转换成的half类型数据。
__ull2half_rd获取输入遵循CAST_FLOOR模式转换成的half类型数据。
__ull2half_ru获取输入遵循CAST_CEIL模式转换成的half类型数据。
__ull2half_rna获取输入遵循CAST_ROUND模式转换成的half类型数据。
__ll2half_rn获取输入遵循CAST_RINT模式转换成的half类型数据。
__ll2half_rz获取输入遵循CAST_TRUNC模式转换成的half类型数据。
__ll2half_rd获取输入遵循CAST_FLOOR模式转换成的half类型数据。
__ll2half_ru获取输入遵循CAST_CEIL模式转换成的half类型数据。
__ll2half_rna获取输入遵循CAST_ROUND模式转换成的half类型数据。
__floats2half2_rn将输入的数据x，y遵循CAST_RINT模式分别转换为half类型并填充到half2的前后两部分，返回转换后的half2类型数据。
__float22half2_rn将float2类型数据遵循CAST_RINT模式转换为half2类型，返回转换后的half2类型数据。
__low2half返回输入数据的低16位。
__low2half2将输入数据的低16位填充到half2并返回。
__low2float将输入数据的低16位转换为浮点数并返回结果。
__lowhigh2highlow将输入数据的高低16位进行交换并返回。
__high2half提取输入half2的高16位，并返回
__high2half2将输入数据的高16位填充到half2并返回结果。
__high2float将输入数据的高16位转换为float类型并返回结果。
__highs2half2分别提取两个half2输入的高16位，并填充到half2中。返回填充后的数据。
__lows2half2分别提取两个half2输入的低16位，并填充到half2中。返回填充后的数据。
__halves2half2将输入的数据分别填充为half2前后两个分量，返回填充后数据。
__half22float2将half2的两个分量分别转换为float，并填充到float2返回。
__ushort_as_half将unsigned short int的按位重新解释为half，即将unsigned short int的数据存储的位按照half的格式进行读取。

表13 half2类型算术函数

接口名功能描述
__haddx2计算两个half2类型数据各分量的相加结果，并遵循CAST_RINT模式舍入。
__hsubx2计算两个half2类型数据各分量的相减结果，并遵循CAST_RINT模式舍入。
__hmulx2计算两个half2类型数据各分量的相乘结果，并遵循CAST_RINT模式舍入。
__hdivx2计算两个half2类型数据各分量的相除结果，并遵循CAST_RINT模式舍入。
__habsx2计算输入half2类型数据各分量的绝对值。
__hfmax2计算两个half2类型数据各分量的乘加的结果（前两个输入相乘后与第三个输入相加），并遵循CAST_RINT模式舍入。
__hnegx2获取输入half2类型数据各分量的负值。
__hfmax2_relu计算两个half2类型数据各分量的乘加的结果（前两个输入相乘后与第三个输入相加），并遵循CAST_RINT模式舍入。负数结果置为0。
__hcmadd将三个half2输入视为复数（第一个分量为实部，第二个分量为虚部），执行复数乘加运算x*y+z。

表14 half2类型比较函数

接口名功能描述
__hbeqx2比较两个half2类型数据的两个分量是否相等，仅当两个分量均相等时返回true。
__hbnex2比较两个half2类型数据的两个分量是否不相等，仅当两个分量均不相等时返回true。
__hblex2比较两个half2类型数据的两个分量，仅当两个分量均满足第一个数小于或等于第二个数时返回true。
__hbgex2比较两个half2类型数据的两个分量，仅当两个分量均满足第一个数大于或等于第二个数时返回true。
__hbltx2比较两个half2类型数据的两个分量，仅当两个分量均满足第一个数小于第二个数时返回true。
__hbgtx2比较两个half2类型数据的两个分量，仅当两个分量均满足第一个数大于第二个数时返回true。
__hbequx2比较两个half2类型数据的两个分量是否相等，当两个分量均相等时返回true。若任一输入的分量为nan，该分量的比较结果为true。
__hbneux2比较两个half2类型数据的两个分量是否不相等，当两个分量均不相等时返回true。若任一输入的分量为nan，该分量的比较结果为true。
__hbleux2比较两个half2类型数据的两个分量，当两个分量均满足第一个数小于或等于第二个数时返回true。若任一输入的分量为nan，该分量的比较结果为true。
__hbgeux2比较两个half2类型数据的两个分量，当两个分量均满足第一个数大于或等于第二个数时返回true。若任一输入的分量为nan，该分量的比较结果为true。
__hbltux2比较两个half2类型数据的两个分量，当两个分量均满足第一个数小于第二个数时返回true。若任一输入的分量为nan，该分量的比较结果为true。
__hbgtux2比较两个half2类型数据的两个分量，当两个分量均满足第一个数大于第二个数时返回true。若任一输入的分量为nan，该分量的比较结果为true。
__heqx2比较两个half2类型数据的两个分量，如果分量相等，则对应比较结果为1.0，否则为0.0。
__hnex2比较两个half2类型数据的两个分量，如果分量不相等，则对应比较结果为1.0，否则为0.0。
__hlex2比较两个half2类型数据的两个分量，如果分量满足第一个数小于或等于第二个数，则对应比较结果为1.0，否则为0.0。
__hgex2比较两个half2类型数据的两个分量，如果分量满足第一个数大于或等于第二个数，则对应比较结果为1.0，否则为0.0。
__hltx2比较两个half2类型数据的两个分量，如果分量满足第一个数小于第二个数，则对应比较结果为1.0，否则为0.0。
__hgtx2比较两个half2类型数据的两个分量，如果分量满足第一个数大于第二个数，则对应比较结果为1.0，否则为0.0。
__hequx2比较两个half2类型数据的两个分量，如果分量相等，则对应比较结果为1.0，否则为0.0。若任一输入的分量为nan，该分量的比较结果为1.0。
__hneux2比较两个half2类型数据的两个分量，如果分量不相等，则对应比较结果为1.0，否则为0.0。若任一输入的分量为nan，该分量的比较结果为1.0。
__hleux2比较两个half2类型数据的两个分量，如果分量满足第一个数小于或等于第二个数，则对应比较结果为1.0，否则为0.0。若任一输入的分量为nan，该分量的比较结果为1.0。
__hgeux2比较两个half2类型数据的两个分量，如果分量满足第一个数大于或等于第二个数，则对应比较结果为1.0，否则为0.0。若任一输入的分量为nan，该分量的比较结果为1.0。
__hltux2比较两个half2类型数据的两个分量，如果分量满足第一个数小于第二个数，则对应比较结果为1.0，否则为0.0。若任一输入的分量为nan，该分量的比较结果为1.0。
__hgtux2比较两个half2类型数据的两个分量，如果分量满足第一个数大于第二个数，则对应比较结果为1.0，否则为0.0。若任一输入的分量为nan，该分量的比较结果为1.0。
__heqx2_mask比较两个half2类型数据的两个分量，结果以unsigned int形式返回，低16位为第一个分量的掩码结果，高16位为第二个分量的掩码结果。如果分量相等，则对应16位掩码为0xFFFF，否则为0x0。
__hnex2_mask比较两个half2类型数据的两个分量，结果以unsigned int形式返回，低16位为第一个分量的掩码结果，高16位为第二个分量的掩码结果。如果分量不相等，则对应16位掩码为0xFFFF，否则为0x0。
__hlex2_mask比较两个half2类型数据的两个分量，结果以unsigned int形式返回，低16位为第一个分量的掩码结果，高16位为第二个分量的掩码结果。如果分量满足第一个数小于或等于第二个数，则对应16位掩码为0xFFFF，否则为0x0。
__hgex2_mask比较两个half2类型数据的两个分量，结果以unsigned int形式返回，低16位为第一个分量的掩码结果，高16位为第二个分量的掩码结果。如果分量满足第一个数大于或等于第二个数，则对应16位掩码为0xFFFF，否则为0x0。
__hltx2_mask比较两个half2类型数据的两个分量，结果以unsigned int形式返回，低16位为第一个分量的掩码结果，高16位为第二个分量的掩码结果。如果分量满足第一个数小于第二个数，则对应16位掩码为0xFFFF，否则为0x0。
__hgtx2_mask比较两个half2类型数据的两个分量，结果以unsigned int形式返回，低16位为第一个分量的掩码结果，高16位为第二个分量的掩码结果。如果分量满足第一个数大于第二个数，则对应16位掩码为0xFFFF，否则为0x0。
__hequx2_mask比较两个half2类型数据的两个分量，结果以unsigned int形式返回，低16位为第一个分量的掩码结果，高16位为第二个分量的掩码结果。如果分量相等，则对应16位掩码为0xFFFF，否则为0x0。若任一输入的分量为nan，对应16位掩码为0xFFFF。
__hneux2_mask比较两个half2类型数据的两个分量，结果以unsigned int形式返回，低16位为第一个分量的掩码结果，高16位为第二个分量的掩码结果。如果分量不相等，则对应16位掩码为0xFFFF，否则为0x0。若任一输入的分量为nan，对应16位掩码为0xFFFF。
__hleux2_mask比较两个half2类型数据的两个分量，结果以unsigned int形式返回，低16位为第一个分量的掩码结果，高16位为第二个分量的掩码结果。如果分量满足第一个数小于或等于第二个数，则对应16位掩码为0xFFFF，否则为0x0。若任一输入的分量为nan，对应16位掩码为0xFFFF。
__hgeux2_mask比较两个half2类型数据的两个分量，结果以unsigned int形式返回，低16位为第一个分量的掩码结果，高16位为第二个分量的掩码结果。如果分量满足第一个数大于或等于第二个数，则对应16位掩码为0xFFFF，否则为0x0。若任一输入的分量为nan，对应16位掩码为0xFFFF。
__hltux2_mask比较两个half2类型数据的两个分量，结果以unsigned int形式返回，低16位为第一个分量的掩码结果，高16位为第二个分量的掩码结果。如果分量满足第一个数小于第二个数，则对应16位掩码为0xFFFF，否则为0x0。若任一输入的分量为nan，对应16位掩码为0xFFFF。
__hgtux2_mask比较两个half2类型数据的两个分量，结果以unsigned int形式返回，低16位为第一个分量的掩码结果，高16位为第二个分量的掩码结果。如果分量满足第一个数大于第二个数，则对应16位掩码为0xFFFF，否则为0x0。若任一输入的分量为nan，对应16位掩码为0xFFFF。
__isnanx2判断half2类型数据的两个分量是否为nan。
__hmaxx2获取两个half2类型数据各分量的最大值。
__hmaxx2_nan获取两个half2类型数据各分量的最大值。任一分量为nan时对应结果为nan。
__hminx2获取两个half2类型数据各分量的最小值。
__hminx2_nan获取两个half2类型数据各分量的最小值。任一分量为nan时对应结果为nan。

表15 half2类型数学库函数

接口名功能描述
h2tanh获取输入数据各元素的三角函数双曲正切值。
h2exp指定输入x，对x的各元素，获取e的该元素次方。
h2exp2指定输入x，对x的各元素，获取2的该元素次方。
h2exp10指定输入x，对x的各元素，获取10的该元素次方。
h2log获取以e为底，输入数据各元素的对数。
h2log2获取以2为底，输入数据各元素的对数。
h2log10获取以10为底，输入数据各元素的对数。
h2cos获取输入数据各元素的三角函数余弦值。
h2sin获取输入数据各元素的三角函数正弦值。
h2sqrt获取输入数据x各元素的平方根。
h2rsqrt获取输入数据x各元素的平方根的倒数。
h2rcp获取输入数据x各元素的倒数。
h2rint获取与输入数据各元素最接近的整数，若存在两个同样接近的整数，则获取其中的偶数。
h2floor获取小于或等于输入数据各元素的最大整数值。
h2ceil获取大于或等于输入数据各元素的最小整数值。
h2trunc获取对输入数据各元素的浮点数截断后的整数。

表16 bfloat16类型算术函数

接口名功能描述
__habs获取输入数据的绝对值。
__hfma对输入数据x、y、z，计算x与y相乘加上z的结果。
__hadd计算两个bfloat16类型数据的相加结果，并遵循CAST_RINT模式舍入。
__hsub计算两个bfloat16类型数据的相减结果，并遵循CAST_RINT模式舍入。
__hmul计算两个bfloat16类型数据的相乘结果，并遵循CAST_RINT模式舍入。
__hdiv计算两个bfloat16类型数据的相除结果，并遵循CAST_RINT模式舍入。
__hneg获取输入bfloat16类型数据的负值。
__hfma_relu对输入bfloat16类型数据x、y、z，计算x与y相乘加上z的结果，并遵循CAST_RINT模式舍入。负数结果置为0。

表17 bfloat16类型比较函数

接口名功能描述
__hmax获取两个输入数据中的最大值。
__hmin获取两个输入数据中的最小值。
__hisnan判断浮点数是否为nan。
__hisinf判断浮点数是否为无穷。
__heq比较两个bfloat16类型数据是否相等，相等时返回true。
__hne比较两个bfloat16类型数据是否不相等，不相等时返回true。
__hle比较两个bfloat16类型数据，仅当第一个数小于或等于第二个数时返回true。
__hge比较两个bfloat16类型数据，仅当第一个数大于或等于第二个数时返回true。
__hlt比较两个bfloat16类型数据，仅当第一个数小于第二个数时返回true。
__hgt比较两个bfloat16类型数据，仅当第一个数大于第二个数时返回true。
__hequ比较两个bfloat16类型数据是否相等，相等时返回true。若任一输入为nan，返回true。
__hneu比较两个bfloat16类型数据是否不相等，不相等时返回true。若任一输入为nan，返回true。
__hleu比较两个bfloat16类型数据，当第一个数小于或等于第二个数时返回true。若任一输入为nan，返回true。
__hgeu比较两个bfloat16类型数据，当第一个数大于或等于第二个数时返回true。若任一输入为nan，返回true。
__hltu比较两个bfloat16类型数据，当第一个数小于第二个数时返回true。若任一输入为nan，返回true。
__hgtu比较两个bfloat16类型数据，当第一个数大于第二个数时返回true。若任一输入为nan，返回true。
__hmax_nan获取两个输入数据中的最大值。任一输入为nan时返回nan。
__hmin_nan获取两个输入数据中的最小值。任一输入为nan时返回nan。

表18 bfloat16类型数学库函数

接口名功能描述
htanh获取输入数据的三角函数双曲正切值。
hexp指定输入x，获取e的x次方。
hexp2指定输入x，获取2的x次方。
hexp10指定输入x，获取10的x次方。
hlog获取以e为底，输入数据的对数。
hlog2获取以2为底，输入数据的对数。
hlog10获取以10为底，输入数据的对数。
hcos获取输入数据的三角函数余弦值。
hsin获取输入数据的三角函数正弦值。
hsqrt获取输入数据x的平方根。
hrsqrt获取输入数据x的平方根的倒数。
hrcp获取输入数据x的倒数。
hrint获取与输入数据最接近的整数，若存在两个同样接近的整数，则获取其中的偶数。
hfloor获取小于或等于输入数据的最大整数值。
hceil获取大于或等于输入数据的最小整数值。
htrunc获取对输入数据的浮点数截断后的整数。

表19 bfloat16类型精度转换函数

接口名功能描述
__float2bfloat16获取输入遵循CAST_RINT模式转换成的bfloat16类型数据。
__float2bfloat16_rn获取输入遵循CAST_RINT模式转换成的bfloat16类型数据。
__float2bfloat16_rn_sat饱和模式下获取输入遵循CAST_RINT模式转换成的bfloat16类型数据。
__float22bfloat162_rn_sat饱和模式下获取输入的两个分量遵循CAST_RINT模式转换成的bfloat16x2_t类型数据。
__float2bfloat16_rz获取输入遵循CAST_TRUNC模式转换成的bfloat16类型数据。
__float2bfloat16_rz_sat饱和模式下获取输入遵循CAST_TRUNC模式转换成的bfloat16类型数据。
__float22bfloat162_rz获取输入的两个分量遵循CAST_TRUNC模式转换成的bfloat16x2_t类型数据。
__float22bfloat162_rz_sat饱和模式下获取输入的两个分量遵循CAST_TRUNC模式转换成的bfloat16x2_t类型数据。
__float2bfloat16_rd获取输入遵循CAST_FLOOR模式转换成的bfloat16类型数据。
__float2bfloat16_rd_sat饱和模式下获取输入遵循CAST_FLOOR模式转换成的bfloat16类型数据。
__float22bfloat162_rd获取输入的两个分量遵循CAST_FLOOR模式转换成的bfloat16x2_t类型数据。
__float22bfloat162_rd_sat饱和模式下获取输入的两个分量遵循CAST_FLOOR模式转换成的bfloat16x2_t类型数据。
__float2bfloat16_ru获取输入遵循CAST_CEIL模式转换成的bfloat16类型数据。
__float2bfloat16_ru_sat饱和模式下获取输入遵循CAST_CEIL模式转换成的bfloat16类型数据。
__float22bfloat162_ru获取输入的两个分量遵循CAST_CEIL模式转换成的bfloat16x2_t类型数据。
__float22bfloat162_ru_sat饱和模式下获取输入的两个分量遵循CAST_CEIL模式转换成的bfloat16x2_t类型数据。
__float2bfloat16_rna获取输入遵循CAST_ROUND模式转换成的bfloat16类型数据。
__float2bfloat16_rna_sat饱和模式下获取输入遵循CAST_ROUND模式转换成的bfloat16类型数据。
__float22bfloat162_rna获取输入的两个分量遵循CAST_ROUND模式转换成的bfloat16x2_t类型数据。
__float22bfloat162_rna_sat饱和模式下获取输入的两个分量遵循CAST_ROUND模式转换成的bfloat16x2_t类型数据。
__half2bfloat16_rn获取输入遵循CAST_RINT模式转换成的bfloat16类型数据。
__half2bfloat16_rz获取输入遵循CAST_TRUNC模式转换成的bfloat16类型数据。
__half2bfloat16_rd获取输入遵循CAST_FLOOR模式转换成的bfloat16类型数据。
__half2bfloat16_ru获取输入遵循CAST_CEIL模式转换成的bfloat16类型数据。
__half2bfloat16_rna获取输入遵循CAST_ROUND模式转换成的bfloat16类型数据。
__bfloat162float获取输入转换为浮点数的结果。
__bfloat162bfloat16_rn获取输入遵循CAST_RINT模式取整后的bfloat16_t类型数据。
__bfloat162bfloat16_rz获取输入遵循CAST_TRUNC模式取整后的bfloat16_t类型数据。
__bfloat162bfloat16_rd获取输入遵循CAST_FLOOR模式取整后的bfloat16_t类型数据。
__bfloat162bfloat16_ru获取输入遵循CAST_CEIL模式取整后的bfloat16_t类型数据。
__bfloat162bfloat16_rna获取输入遵循CAST_ROUND模式取整后的bfloat16_t类型数据。
__bfloat162uint_rn获取输入遵循CAST_RINT模式转换成的无符号整数。
__bfloat162uint_rz获取输入遵循CAST_TRUNC模式转换成的无符号整数。
__bfloat162uint_rd获取输入遵循CAST_FLOOR模式转换成的无符号整数。
__bfloat162uint_ru获取输入遵循CAST_CEIL模式转换成的无符号整数。
__bfloat162uint_rna获取输入遵循CAST_ROUND模式转换成的无符号整数。
__bfloat162int_rn获取输入遵循CAST_RINT模式转换成的有符号整数。
__bfloat162int_rz获取输入遵循CAST_TRUNC模式转换成的有符号整数。
__bfloat162int_rd获取输入遵循CAST_FLOOR模式转换成的有符号整数。
__bfloat162int_ru获取输入遵循CAST_CEIL模式转换成的有符号整数。
__bfloat162int_rna获取输入遵循CAST_ROUND模式转换成的有符号整数。
__bfloat162ull_rn获取输入遵循CAST_RINT模式转换成的64位无符号整数。
__bfloat162ull_rz获取输入遵循CAST_TRUNC模式转换成的64位无符号整数。
__bfloat162ull_rd获取输入遵循CAST_FLOOR模式转换成的64位无符号整数。
__bfloat162ull_ru获取输入遵循CAST_CEIL模式转换成的64位无符号整数。
__bfloat162ull_rna获取输入遵循CAST_ROUND模式转换成的64位无符号整数。
__bfloat162ll_rn获取输入遵循CAST_RINT模式转换成的64位有符号整数。
__bfloat162ll_rz获取输入遵循CAST_TRUNC模式转换成的64位有符号整数。
__bfloat162ll_rd获取输入遵循CAST_FLOOR模式转换成的64位有符号整数。
__bfloat162ll_ru获取输入遵循CAST_CEIL模式转换成的64位有符号整数。
__bfloat162ll_rna获取输入遵循CAST_ROUND模式转换成的64位有符号整数。
__uint2bfloat16_rn获取输入遵循CAST_RINT模式转换成的bfloat16类型数据。
__uint2bfloat16_rz获取输入遵循CAST_TRUNC模式转换成的bfloat16类型数据。
__uint2bfloat16_rd获取输入遵循CAST_FLOOR模式转换成的bfloat16类型数据。
__uint2bfloat16_ru获取输入遵循CAST_CEIL模式转换成的bfloat16类型数据。
__uint2bfloat16_rna获取输入遵循CAST_ROUND模式转换成的bfloat16类型数据。
__int2bfloat16_rn获取输入遵循CAST_RINT模式转换成的bfloat16类型数据。
__int2bfloat16_rz获取输入遵循CAST_TRUNC模式转换成的bfloat16类型数据。
__int2bfloat16_rd获取输入遵循CAST_FLOOR模式转换成的bfloat16类型数据。
__int2bfloat16_ru获取输入遵循CAST_CEIL模式转换成的bfloat16类型数据。
__int2bfloat16_rna获取输入遵循CAST_ROUND模式转换成的bfloat16类型数据。
__ull2bfloat16_rn获取输入遵循CAST_RINT模式转换成的bfloat16类型数据。
__ull2bfloat16_rz获取输入遵循CAST_TRUNC模式转换成的bfloat16类型数据。
__ull2bfloat16_rd获取输入遵循CAST_FLOOR模式转换成的bfloat16类型数据。
__ull2bfloat16_ru获取输入遵循CAST_CEIL模式转换成的bfloat16类型数据。
__ull2bfloat16_rna获取输入遵循CAST_ROUND模式转换成的bfloat16类型数据。
__ll2bfloat16_rn获取输入遵循CAST_RINT模式转换成的bfloat16类型数据。
__ll2bfloat16_rz获取输入遵循CAST_TRUNC模式转换成的bfloat16类型数据。
__ll2bfloat16_rd获取输入遵循CAST_FLOOR模式转换成的bfloat16类型数据。
__ll2bfloat16_ru获取输入遵循CAST_CEIL模式转换成的bfloat16类型数据。
__ll2bfloat16_rna获取输入遵循CAST_ROUND模式转换成的bfloat16类型数据。
__float2bfloat162_rn将float类型数据遵循CAST_RINT模式转换为bfloat16类型并填充到bfloat16x2的前后两部分，返回填充后的bfloat16x2类型数据。
__floats2bfloat162_rn将输入的数据x，y遵循CAST_RINT模式分别转换为bfloat16类型并填充到bfloat16x2的前后两部分，返回转换后的bfloat16x2类型数据。
__float22bfloat162_rn将float2类型数据遵循CAST_RINT模式转换为bfloat16x2类型，返回转换后的bfloat16x2类型数据。
__bfloat162bfloat162将输入的数据填充为bfloat16x2前后两个分量，返回转换后的bfloat16x2类型数据。
__halves2bfloat162将输入的数据分别填充为bfloat16x2前后两个分量，返回填充后数据。
__high2bfloat16提取输入bfloat16x2的高16位，并返回。
__high2bfloat162将输入数据的高16位填充到bfloat16x2并返回结果。
__high2float将输入数据的高16位转换为float类型并返回结果。
__highs2bfloat162分别提取两个bfloat16x2输入的高16位，并填充到bfloat16x2中。返回填充后的数据。
__low2bfloat16返回输入数据的低16位。
__low2bfloat162将输入数据的低16位填充到bfloat16x2并返回。
__low2float将输入数据的低16位转换为浮点数并返回结果。
__lowhigh2highlow将输入数据的高低16位进行交换并返回。
__lows2bfloat162分别提取两个bfloat16x2输入的低16位，并填充到bfloat16x2中。返回填充后的数据。
__bfloat1622float2将bfloat16x2的两个分量分别转换为float，并填充到float2返回。
__ushort_as_bfloat16将unsigned short int的按位重新解释为bfloat16，即将unsigned short int的数据存储的位按照bfloat16的格式进行读取。

表20 bfloat16x2类型算术函数

接口名功能描述
__haddx2计算两个bfloat16x2_t类型数据各分量的相加结果，并遵循CAST_RINT模式舍入。
__hsubx2计算两个bfloat16x2_t类型数据各分量的相减结果，并遵循CAST_RINT模式舍入。
__hmulx2计算两个bfloat16x2_t类型数据各分量的相乘结果，并遵循CAST_RINT模式舍入。
__hdivx2计算两个bfloat16x2_t类型数据各分量的相除结果，并遵循CAST_RINT模式舍入。
__habsx2计算输入bfloat16x2_t类型数据各分量的绝对值。
__hfmax2计算两个bfloat16x2_t类型数据各分量的乘加的结果（前两个输入相乘后与第三个输入相加），并遵循CAST_RINT模式舍入。
__hnegx2获取输入bfloat16x2_t类型数据各分量的负值。
__hfmax2_relu计算两个bfloat16x2_t类型数据各分量的乘加的结果（前两个输入相乘后与第三个输入相加），并遵循CAST_RINT模式舍入。负数结果置为0。
__hcmadd将三个bfloat16x2_t输入视为复数（第一个分量为实部，第二个分量为虚部），执行复数乘加运算x*y+z。

表21 bfloat16x2类型比较函数

接口名功能描述
__hbeqx2比较两个bfloat16x2_t类型数据的两个分量是否相等，仅当两个分量均相等时返回true。
__hbnex2比较两个bfloat16x2_t类型数据的两个分量是否不相等，仅当两个分量均不相等时返回true。
__hblex2比较两个bfloat16x2_t类型数据的两个分量，仅当两个分量均满足第一个数小于或等于第二个数时返回true。
__hbgex2比较两个bfloat16x2_t类型数据的两个分量，仅当两个分量均满足第一个数大于或等于第二个数时返回true。
__hbltx2比较两个bfloat16x2_t类型数据的两个分量，仅当两个分量均满足第一个数小于第二个数时返回true。
__hbgtx2比较两个bfloat16x2_t类型数据的两个分量，仅当两个分量均满足第一个数大于第二个数时返回true。
__hbequx2比较两个bfloat16x2_t类型数据的两个分量是否相等，当两个分量均相等时返回true。若任一输入的分量为nan，该分量的比较结果为true。
__hbneux2比较两个bfloat16x2_t类型数据的两个分量是否不相等，当两个分量均不相等时返回true。若任一输入的分量为nan，该分量的比较结果为true。
__hbleux2比较两个bfloat16x2_t类型数据的两个分量，当两个分量均满足第一个数小于或等于第二个数时返回true。若任一输入的分量为nan，该分量的比较结果为true。
__hbgeux2比较两个bfloat16x2_t类型数据的两个分量，当两个分量均满足第一个数大于或等于第二个数时返回true。若任一输入的分量为nan，该分量的比较结果为true。
__hbltux2比较两个bfloat16x2_t类型数据的两个分量，当两个分量均满足第一个数小于第二个数时返回true。若任一输入的分量为nan，该分量的比较结果为true。
__hbgtux2比较两个bfloat16x2_t类型数据的两个分量，当两个分量均满足第一个数大于第二个数时返回true。若任一输入的分量为nan，该分量的比较结果为true。
__heqx2比较两个bfloat16x2_t类型数据的两个分量，如果分量相等，则对应比较结果为1.0，否则为0.0。
__hnex2比较两个bfloat16x2_t类型数据的两个分量，如果分量不相等，则对应比较结果为1.0，否则为0.0。
__hlex2比较两个bfloat16x2_t类型数据的两个分量，如果分量满足第一个数小于或等于第二个数，则对应比较结果为1.0，否则为0.0。
__hgex2比较两个bfloat16x2_t类型数据的两个分量，如果分量满足第一个数大于或等于第二个数，则对应比较结果为1.0，否则为0.0。
__hltx2比较两个bfloat16x2_t类型数据的两个分量，如果分量满足第一个数小于第二个数，则对应比较结果为1.0，否则为0.0。
__hgtx2比较两个bfloat16x2_t类型数据的两个分量，如果分量满足第一个数大于第二个数，则对应比较结果为1.0，否则为0.0。
__hequx2比较两个bfloat16x2_t类型数据的两个分量，如果分量相等，则对应比较结果为1.0，否则为0.0。若任一输入的分量为nan，该分量的比较结果为1.0。
__hneux2比较两个bfloat16x2_t类型数据的两个分量，如果分量不相等，则对应比较结果为1.0，否则为0.0。若任一输入的分量为nan，该分量的比较结果为1.0。
__hleux2比较两个bfloat16x2_t类型数据的两个分量，如果分量满足第一个数小于或等于第二个数，则对应比较结果为1.0，否则为0.0。若任一输入的分量为nan，该分量的比较结果为1.0。
__hgeux2比较两个bfloat16x2_t类型数据的两个分量，如果分量满足第一个数大于或等于第二个数，则对应比较结果为1.0，否则为0.0。若任一输入的分量为nan，该分量的比较结果为1.0。
__hltux2比较两个bfloat16x2_t类型数据的两个分量，如果分量满足第一个数小于第二个数，则对应比较结果为1.0，否则为0.0。若任一输入的分量为nan，该分量的比较结果为1.0。
__hgtux2比较两个bfloat16x2_t类型数据的两个分量，如果分量满足第一个数大于第二个数，则对应比较结果为1.0，否则为0.0。若任一输入的分量为nan，该分量的比较结果为1.0。
__heqx2_mask比较两个bfloat16x2_t类型数据的两个分量，结果以unsigned int形式返回，低16位为第一个分量的掩码结果，高16位为第二个分量的掩码结果。如果分量相等，则对应16位掩码为0xFFFF，否则为0x0。
__hnex2_mask比较两个bfloat16x2_t类型数据的两个分量，结果以unsigned int形式返回，低16位为第一个分量的掩码结果，高16位为第二个分量的掩码结果。如果分量不相等，则对应16位掩码为0xFFFF，否则为0x0。
__hlex2_mask比较两个bfloat16x2_t类型数据的两个分量，结果以unsigned int形式返回，低16位为第一个分量的掩码结果，高16位为第二个分量的掩码结果。如果分量满足第一个数小于或等于第二个数，则对应16位掩码为0xFFFF，否则为0x0。
__hgex2_mask比较两个bfloat16x2_t类型数据的两个分量，结果以unsigned int形式返回，低16位为第一个分量的掩码结果，高16位为第二个分量的掩码结果。如果分量满足第一个数大于或等于第二个数，则对应16位掩码为0xFFFF，否则为0x0。
__hltx2_mask比较两个bfloat16x2_t类型数据的两个分量，结果以unsigned int形式返回，低16位为第一个分量的掩码结果，高16位为第二个分量的掩码结果。如果分量满足第一个数小于第二个数，则对应16位掩码为0xFFFF，否则为0x0。
__hgtx2_mask比较两个bfloat16x2_t类型数据的两个分量，结果以unsigned int形式返回，低16位为第一个分量的掩码结果，高16位为第二个分量的掩码结果。如果分量满足第一个数大于第二个数，则对应16位掩码为0xFFFF，否则为0x0。
__hequx2_mask比较两个bfloat16x2_t类型数据的两个分量，结果以unsigned int形式返回，低16位为第一个分量的掩码结果，高16位为第二个分量的掩码结果。如果分量相等，则对应16位掩码为0xFFFF，否则为0x0。若任一输入的分量为nan，对应16位掩码为0xFFFF。
__hneux2_mask比较两个bfloat16x2_t类型数据的两个分量，结果以unsigned int形式返回，低16位为第一个分量的掩码结果，高16位为第二个分量的掩码结果。如果分量不相等，则对应16位掩码为0xFFFF，否则为0x0。若任一输入的分量为nan，对应16位掩码为0xFFFF。
__hleux2_mask比较两个bfloat16x2_t类型数据的两个分量，结果以unsigned int形式返回，低16位为第一个分量的掩码结果，高16位为第二个分量的掩码结果。如果分量满足第一个数小于或等于第二个数，则对应16位掩码为0xFFFF，否则为0x0。若任一输入的分量为nan，对应16位掩码为0xFFFF。
__hgeux2_mask比较两个bfloat16x2_t类型数据的两个分量，结果以unsigned int形式返回，低16位为第一个分量的掩码结果，高16位为第二个分量的掩码结果。如果分量满足第一个数大于或等于第二个数，则对应16位掩码为0xFFFF，否则为0x0。若任一输入的分量为nan，对应16位掩码为0xFFFF。
__hltux2_mask比较两个bfloat16x2_t类型数据的两个分量，结果以unsigned int形式返回，低16位为第一个分量的掩码结果，高16位为第二个分量的掩码结果。如果分量满足第一个数小于第二个数，则对应16位掩码为0xFFFF，否则为0x0。若任一输入的分量为nan，对应16位掩码为0xFFFF。
__hgtux2_mask比较两个bfloat16x2_t类型数据的两个分量，结果以unsigned int形式返回，低16位为第一个分量的掩码结果，高16位为第二个分量的掩码结果。如果分量满足第一个数大于第二个数，则对应16位掩码为0xFFFF，否则为0x0。若任一输入的分量为nan，对应16位掩码为0xFFFF。
__isnanx2判断bfloat16x2_t类型数据的两个分量是否为nan。
__hmaxx2获取两个bfloat16x2_t类型数据各分量的最大值。
__hmaxx2_nan获取两个bfloat16x2_t类型数据各分量的最大值。任一分量为nan时对应结果为nan。
__hminx2获取两个bfloat16x2_t类型数据各分量的最小值。
__hminx2_nan获取两个bfloat16x2_t类型数据各分量的最小值。任一分量为nan时对应结果为nan。

表22 bfloat16x2类型数学库函数

接口名功能描述
h2tanh获取输入数据各元素的三角函数双曲正切值。
h2exp指定输入x，对x的各元素，获取e的该元素次方。
h2exp2指定输入x，对x的各元素，获取2的该元素次方。
h2exp10指定输入x，对x的各元素，获取10的该元素次方。
h2log获取以e为底，输入数据各元素的对数。
h2log2获取以2为底，输入数据各元素的对数。
h2log10获取以10为底，输入数据各元素的对数。
h2cos获取输入数据各元素的三角函数余弦值。
h2sin获取输入数据各元素的三角函数正弦值。
h2sqrt获取输入数据x各元素的平方根。
h2rsqrt获取输入数据x各元素的平方根的倒数。
h2rcp获取输入数据x各元素的倒数。
h2rint获取与输入数据各元素最接近的整数，若存在两个同样接近的整数，则获取其中的偶数。
h2floor获取小于或等于输入数据各元素的最大整数值。
h2ceil获取大于或等于输入数据各元素的最小整数值。
h2trunc获取对输入数据各元素的浮点数截断后的整数。

表23 float类型数学库函数

接口名功能描述
tanf获取输入数据的三角函数正切值。
tanhf获取输入数据的三角函数双曲正切值。
tanpif获取输入数据与π相乘的正切值。
atanf获取输入数据的反正切值。
atan2f获取输入数据y/x的反正切值。
atanhf获取输入数据的反双曲正切值。
expf指定输入x，获取e的x次方。
exp2f指定输入x，获取2的x次方。
exp10f指定输入x，获取10的x次方。
expm1f指定输入x，获取e的x次方减1。
logf获取以e为底，输入数据的对数。
log2f获取以2为底，输入数据的对数。
log10f获取以10为底，输入数据的对数。
log1pf获取以e为底，输入数据加1的对数。
logbf计算以2为底，输入数据的对数，并对结果向下取整，返回浮点数。
ilogbf计算以2为底，输入数据的对数，并对结果向下取整，返回整数。
cosf获取输入数据的三角函数余弦值。
coshf获取输入数据的双曲余弦值。
cospif获取输入数据与π相乘的余弦值。
acosf获取输入数据的反余弦值。
acoshf获取输入数据的双曲反余弦值。
sinf获取输入数据的三角函数正弦值。
sinhf获取输入数据的双曲正弦值。
sinpif获取输入数据与π相乘的正弦值。
asinf获取输入数据的反正弦值。
asinhf获取输入数据的双曲反正弦值。
sincosf获取输入数据的三角函数正弦值和余弦值。
sincospif获取输入数据与π相乘的三角函数正弦值和余弦值。
frexpf将x转换为归一化[1/2, 1)的有符号数乘以2的积分幂。
ldexpf获取输入x乘以2的exp次幂的结果。
sqrtf获取输入数据x的平方根。
rsqrtf获取输入数据x的平方根的倒数。
hypotf获取输入数据x、y的平方和x^2 + y^2的平方根。
rhypotf获取输入数据x、y的平方和x^2 + y^2的平方根的倒数。
powf获取输入数据x的y次幂。
norm3df获取输入数据a、b、c的平方和a^2 + b^2 + c^2的平方根。
rnorm3df获取输入数据a、b、c的平方和a^2 + b^2 + c^2的平方根的倒数。
norm4df获取输入数据a、b、c、d的平方和a^2 + b^2+ c^2+ d^2的平方根。
rnorm4df获取输入数据a、b、c、d的平方和a^2 + b^2 + c^2 + d^2的平方根的倒数。
normf获取输入数据a中前n个元素的平方和a[0]^2 + a[1]^2 +...+ a[n-1]^2的平方根。
rnormf获取输入数据a中前n个元素的平方和a[0]^2 + a[1]^2 + ...+ a[n-1]^2的平方根的倒数。
cbrtf获取输入数据x的立方根。
rcbrtf获取输入数据x的立方根的倒数。
erff获取输入数据的误差函数值。
erfcf获取输入数据的互补误差函数值。
erfinvf获取输入数据的逆误差函数值。
erfcinvf获取输入数据的逆互补误差函数值。
erfcxf获取输入数据的缩放互补误差函数值。
tgammaf获取输入数据x的伽马函数值。
lgammaf获取输入数据x伽马值的绝对值并求自然对数。
cyl_bessel_i0f获取输入数据x的0阶常规修正圆柱贝塞尔函数的值。
cyl_bessel_i1f获取输入数据x的1阶常规修正圆柱贝塞尔函数的值。
normcdff获取输入数据x的标准正态分布的累积分布函数值。
normcdfinvf获取输入数据x的标准正态累积分布的逆函数
j0f获取输入数据x的0阶第一类贝塞尔函数j0的值。
j1f获取输入数据x的1阶第一类贝塞尔函数j1的值。
jnf获取输入数据x的n阶第一类贝塞尔函数jn的值。
y0f获取输入数据x的0阶第二类贝塞尔函数y0的值。
y1f获取输入数据x的1阶第二类贝塞尔函数y1的值。
ynf获取输入数据x的n阶第二类贝塞尔函数yn的值。
fabsf获取输入数据的绝对值。
fmaf对输入数据x、y、z，计算x与y相乘加上z的结果。
fmaxf获取两个输入数据中的最大值。
fminf获取两个输入数据中的最小值。
fdimf获取输入数据的差值，差值小于0时，返回0。
remquof获取输入数据x除以y的余数。求余数时，商取最接近x除以y浮点数结果的整数，当x除以y的浮点数结果与左右最接近的整数距离相等时，商取偶数，同时将商赋值给指针变量quo。
fmodf获取输入数据x除以y的余数。求余数时，商取x除以y浮点数结果的整数部分。
remainderf获取输入数据x除以y的余数。求余数时，商取最接近x除以y浮点数结果的整数，当x除以y的浮点数结果与左右最接近的整数距离相等时，商取偶数。
copysignf获取由第一个输入x的数值部分和第二个输入y的符号部分拼接得到的浮点数。
nearbyintf获取与输入浮点数最接近的整数，输入浮点数与左右整数的距离相等时，返回偶数。
nextafterf如果y大于x，返回比x大的下一个可表示的浮点值，即浮点数二进制最低位加1。

如果y小于x，返回比x小的下一个可表示的浮点值，即浮点数二进制最低位减1。

如果y等于x，返回x。
scalbnf获取输入数据x与2的n次方的乘积。
scalblnf获取输入数据x与2的n次方的乘积。
modff将输入数据分解为小数部分和整数部分。
fdividef获取两个输入数据相除的结果。
signbit获取输入数据的符号位。
__saturatef将输入数据钳位到[0.0, 1.0]区间。
__fdividef获取两个输入数据相除的结果。
rintf获取与输入数据最接近的整数，若存在两个同样接近的整数，则获取其中的偶数。
lrintf获取与输入数据最接近的整数，若存在两个同样接近的整数，则获取其中的偶数。
llrintf获取与输入数据最接近的整数，若存在两个同样接近的整数，则获取其中的偶数。
roundf获取对输入数据四舍五入后的整数。
lroundf获取对输入数据四舍五入后的整数。
llroundf获取对输入数据四舍五入后的整数。
floorf获取小于或等于输入数据的最大整数值。
ceilf获取大于或等于输入数据的最小整数值。
truncf获取对输入数据的浮点数截断后的整数。
isfinite判断浮点数是否为有限数（非inf、非nan）。
isnan判断浮点数是否为nan。
isinf判断浮点数是否为无穷。

表24 类型转换函数

接口名功能描述
__float2float_rn获取输入遵循CAST_RINT模式取整后的浮点数。
__float2float_rz获取输入遵循CAST_TRUNC模式取整后的浮点数。
__float2float_rd获取输入遵循CAST_FLOOR模式取整后的浮点数。
__float2float_ru获取输入遵循CAST_CEIL模式取整后的浮点数。
__float2float_rna获取输入遵循CAST_ROUND模式取整后的浮点数。
__float2uint_rn获取输入遵循CAST_RINT模式转换成的无符号整数。
__float2uint_rz获取输入遵循CAST_TRUNC模式转换成的无符号整数。
__float2uint_rd获取输入遵循CAST_FLOOR模式转换成的无符号整数。
__float2uint_ru获取输入遵循CAST_CEIL模式转换成的无符号整数。
__float2uint_rna获取输入遵循CAST_ROUND模式转换成的无符号整数。
__float2int_rn获取输入遵循CAST_RINT模式转换成的有符号整数。
__float2int_rz获取输入遵循CAST_TRUNC模式转换成的有符号整数。
__float2int_rd获取输入遵循CAST_FLOOR模式转换成的有符号整数。
__float2int_ru获取输入遵循CAST_CEIL模式转换成的有符号整数。
__float2int_rna获取输入遵循CAST_ROUND模式转换成的有符号整数。
__float2ull_rn获取输入遵循CAST_RINT模式转换成的64位无符号整数。
__float2ull_rz获取输入遵循CAST_TRUNC模式转换成的64位无符号整数。
__float2ull_rd获取输入遵循CAST_FLOOR模式转换成的64位无符号整数。
__float2ull_ru获取输入遵循CAST_CEIL模式转换成的64位无符号整数。
__float2ull_rna获取输入遵循CAST_ROUND模式转换成的64位无符号整数。
__float2ll_rn获取输入遵循CAST_RINT模式转换成的64位有符号整数。
__float2ll_rz获取输入遵循CAST_TRUNC模式转换成的64位有符号整数。
__float2ll_rd获取输入遵循CAST_FLOOR模式转换成的64位有符号整数。
__float2ll_ru获取输入遵循CAST_CEIL模式转换成的64位有符号整数。
__float2ll_rna获取输入遵循CAST_ROUND模式转换成的64位有符号整数。
__uint2float_rn获取输入遵循CAST_RINT模式转换成的浮点数。
__uint2float_rz获取输入遵循CAST_TRUNC模式转换成的浮点数。
__uint2float_rd获取输入遵循CAST_FLOOR模式转换成的浮点数。
__uint2float_ru获取输入遵循CAST_CEIL模式转换成的浮点数。
__uint2float_rna获取输入遵循CAST_ROUND模式转换成的浮点数。
__int2float_rn获取输入遵循CAST_RINT模式转换成的浮点数。
__int2float_rz获取输入遵循CAST_TRUNC模式转换成的浮点数。
__int2float_rd获取输入遵循CAST_FLOOR模式转换成的浮点数。
__int2float_ru获取输入遵循CAST_CEIL模式转换成的浮点数。
__int2float_rna获取输入遵循CAST_ROUND模式转换成的浮点数。
__ull2float_rn获取输入遵循CAST_RINT模式转换成的浮点数。
__ull2float_rz获取输入遵循CAST_TRUNC模式转换成的浮点数。
__ull2float_rd获取输入遵循CAST_FLOOR模式转换成的浮点数。
__ull2float_ru获取输入遵循CAST_CEIL模式转换成的浮点数。
__ull2float_rna获取输入遵循CAST_ROUND模式转换成的浮点数。
__ll2float_rn获取输入遵循CAST_RINT模式转换成的浮点数。
__ll2float_rz获取输入遵循CAST_TRUNC模式转换成的浮点数。
__ll2float_rd获取输入遵循CAST_FLOOR模式转换成的浮点数。
__ll2float_ru获取输入遵循CAST_CEIL模式转换成的浮点数。
__ll2float_rna获取输入遵循CAST_ROUND模式转换成的浮点数。
__int_as_float将整数中的位重新解释为浮点数。
__uint_as_float将无符号整数中的位重新解释为浮点数。
__float_as_int将浮点数中的位重新解释为有符号整数。
__float_as_uint将浮点数中的位重新解释为无符号整数。

表25 整型数学库函数

接口名功能描述
labs获取输入数据的绝对值。
llabs获取输入数据的绝对值。
llmax获取两个输入数据中的最大值。
ullmax获取两个输入数据中的最大值。
umax获取两个输入数据中的最大值。
llmin获取两个输入数据中的最小值。
ullmin获取两个输入数据中的最小值。
umin获取两个输入数据中的最小值。
__mulhi获取输入int32类型数据x和y乘积的高32位。
__umulhi获取输入uint32类型数据x和y乘积的高32位。
__mul64hi获取输入int64类型数据x和y乘积的高64位。
__umul64hi获取输入uint64类型数据x和y乘积的高64位。
__mul_i32toi64计算输入32位整数x和y的乘积，返回64位结果。
__brev将输入数据的位序反转，返回反转后的值。
__clz从输入数据的二进制最高有效位开始，返回连续的前导零的位数。
__ffs从二进制输入数据的最低位开始，查找第一个值为1的比特位的位置，并返回该位置的索引，索引从1开始计数；如果二进制数据中没有1，则返回0。
__popc统计输入数据从二进制的高位到低位比特位为1的数量。
__byte_perm由输入的两个4字节的uint32_t类型数据组成一个8个字节的64比特位的整数，通过选择器s指定选取其中的4个字节，将这4个字节从低位到高位拼成一个uint32_t类型的整数。
__sad对输入数据x、y、z，计算
__usad对输入数据x、y、z，计算
__mul24获取输入int32类型数据x和y低24位乘积的低32位结果。x和y的高8位被忽略。
__umul24获取输入uint32类型数据x和y低24位乘积的低32位结果。x和y的高8位被忽略。
__hadd获取输入int32类型数据x和y的平均值，避免中间求和溢出。
__rhadd获取输入int32类型数据x和y的向上取整平均值，避免中间求和溢出。
__uhadd获取输入uint32类型数据x和y的平均值，避免中间求和溢出。
__urhadd获取输入uint32类型数据x和y的向上取整平均值，避免中间求和溢出。
max获取两个输入数据中的最大值。
min获取两个输入数据中的最小值。

访存函数​

表26 访存函数

接口名功能描述
asc_ldcg从L2 Cache加载缓存的数据，如果缓存命中，则直接返回数据。若未命中，则从Global Memory地址预加载数据缓存至L2 Cache，并返回数据。
asc_ldca首先从Data Cache加载缓存数据，若未命中，则尝试从L2 Cache加载。如果Data Cache和L2 Cache中均未找到所需数据，则从Global Memory中读取数据，然后将其缓存到L2 Cache和Data Cache中。
asc_stcg将指定数据存储到Global Memory的地址address中，并缓存到L2 Cache，但不缓存至Data Cache。
asc_stwt将指定数据存储到Global Memory的地址address中，并缓存至Data Cache和L2 Cache。
asc_dcci_single刷新指定地址所在的Cache Line，保证数据读取时Cache的一致性。
asc_dcci_entire刷新核内的整个Data Cache，保证数据读取时Cache的一致性。
asc_nop生成一条空操作指令，占用当前线程15个时钟周期，不执行任何实际计算与访存操作。

协作组​

表27 thread_block接口

接口名功能描述
this_thread_block获取当前线程块协作组对象。
sync同步线程块内所有线程。
size获取线程块内线程总数。
num_threads获取线程块内线程总数。
thread_rank获取当前线程在线程块内的排名。
thread_index获取当前线程在线程块内的三维索引。
group_index获取当前线程块在网格中的三维索引。
group_dim获取线程块的维度。
dim_threads获取线程块内线程的三维维度。

表28 coalesced_group接口

接口名功能描述
coalesced_threads获取当前活跃线程组成的协作组对象。
sync同步组内所有线程。
size获取组内线程总数。
num_threads获取组内线程总数。
thread_rank获取当前线程在组内的排名。
meta_group_rank获取当前协作组在父组的排名。
meta_group_size获取父组被划分时创建的子组数量。
shfl组内线程的数据交换，直接读取组内指定线程的数据。
shfl_up获取组内当前线程向前偏移delta的线程的数据。
shfl_down获取组内当前线程向后偏移delta的线程的数据。
ballot判断组内每个活跃线程的输入是否非零。
any判断是否有组内线程的输入不为0。
all判断是否所有组内线程的输入均不为0。

表29 thread_block_tile接口

接口名功能描述
tiled_partition创建指定大小的thread_block_tile协作组。
sync同步组内所有线程。
size获取组内线程总数。
num_threads获取组内线程总数。
thread_rank获取当前线程在组内的排名。
meta_group_rank获取当前协作组在父组的排名。
meta_group_size获取父组被划分时创建的子组数量。
shfl组内线程的数据交换，直接读取组内指定线程的数据。
shfl_up获取组内当前线程向前偏移delta的线程的数据。
shfl_down获取组内当前线程向后偏移delta的线程的数据。
shfl_xor获取组内与当前线程rank做异或运算后的线程的数据。
ballot判断组内每个活跃线程的输入是否非零。
any判断是否有组内线程的输入不为0。
all判断是否所有组内线程的输入均不为0。

表30 协作组划分接口

接口名功能描述
tiled_partition将一个线程组按指定大小划分为多个子组。
binary_partition根据一个标签（0或1）将父组划分为两个子组。

免责声明：本站内容由 asc-devkit 仓 master 分支自动编译生成，属于持续开发版本，可能存在缺陷，仅供预览与参考。如需稳定及商用资料，请查阅官方 昇腾社区。

Copyright Huawei Technologies Co., Ltd.