__asc_cvt_float2_to_fp8x2

产品支持情况

功能说明

将float2类型数据的两个分量，按照CAST_RINT模式转换为指定类型（float8_e4m3x2_t和float8_e5m2x2_t）的8位浮点数，并根据指定的饱和模式（饱和或非饱和）进行溢出处理。转换结果以位级打包形式存储为__asc_fp8x2_storage_t类型，该类型为16位无符号整数unsigned short int，用于存储float8_e4m3x2_t或float8_e5m2x2_t类型的数据。

函数原型

1

__asc_fp8x2_storage_t __asc_cvt_float2_to_fp8x2(const float2 x, const __asc_saturation_t saturate, const __asc_fp8_interpretation_t fp8_interpretation)

参数说明

返回值说明

输入的两个分量遵循CAST_RINT模式，根据指定的8位浮点数类型和指定的饱和模式，转换成的__asc_fp8x2_storage_t类型数据。

• float8_e4m3_t数据类型没有inf值，非饱和模式下输入超出该类型所能表示的范围时，其溢出结果为nan，饱和模式下溢出结果为该类型表示的最大值或最小值。
• float8_e5m2_t类型则有inf值，非饱和模式下输入超出该类型所能表示的范围，转换结果为对应符号的inf值，饱和模式下溢出结果为该类型表示的最大值或最小值。
• 饱和模式下，对于float8_e5m2_t和float8_e4m3_t数据类型，当寄存器CTRL[50]=0时，nan值会被转换为0，当寄存器CTRL[50]=1时，nan值被转化为fp8类型下的nan值。CTRL寄存器值可通过GetCtrlSpr(ISASI)接口获取。

约束说明

使用此接口前需将CTRL[60]寄存器设置为0，否则饱和行为控制不生效。设置方式请参见控制饱和行为的方式。

SIMT编程场景当前不支持使用该接口。

需要包含的头文件

使用该接口需要包含"simt_api/asc_fp8.h"头文件。

1

#include "simt_api/asc_fp8.h"

调用示例

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14

// 使用小向量可提升数据搬运效率
__simt_vf__ __launch_bounds__(1024) inline void simt_asc_cvt_float2_to_fp8x2(__gm__ float2* input, __gm__ uint16_t* output, uint32_t input_total_length)
{
    uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
    // 每个线程处理1个float2类型的数据，即2个float类型的数据，因此idx >= input_total_length / 2的线程不处理数据
    if (idx >= input_total_length / 2) {
        return;
    }
    output[idx] = __asc_cvt_float2_to_fp8x2(input[idx], __ASC_NOSAT, __ASC_E4M3);
}
__global__ __vector__ void cast_kernel(__gm__ float* input, __gm__ uint16_t* output, uint32_t input_total_length)
{
    asc_vf_call<simt_asc_cvt_float2_to_fp8x2>(dim3(1024), (__gm__ float2*)input, output, input_total_length);
}