开发者
资源

__halves2bfloat162

产品支持情况

产品

是否支持

Ascend 950PR/Ascend 950DT

Atlas A3 训练系列产品/Atlas A3 推理系列产品

x

Atlas A2 训练系列产品/Atlas A2 推理系列产品

x

Atlas 200I/500 A2 推理产品

x

Atlas 推理系列产品AI Core

x

Atlas 推理系列产品Vector Core

x

Atlas 训练系列产品

x

功能说明

将输入的数据分别填充为bfloat16x2前后两个分量,返回填充后数据。

函数原型

1
inline bfloat16x2_t __halves2bfloat162(const bfloat16_t x, const bfloat16_t y)

参数说明

表1 参数说明

参数名

输入/输出

描述

x

输入

源操作数。

y

输入

源操作数。

返回值说明

输入的数据分别填充为bfloat16x2前后两个分量,返回填充后数据。

约束说明

需要包含的头文件

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

1
#include "simt_api/asc_bf16.h"

调用示例

  • SIMT编程场景:
     1
     2
     3
     4
     5
     6
     7
     8
     9
    10
    11
    12
    13
    // 使用小向量可提升数据搬运效率
    __aicore__ void simt_halves2bfloat162(bfloat16_t* input1, bfloat16_t* input2, bfloat16x2_t* output, uint32_t input_total_length)
    {
        uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
        if (idx > input_total_length) {
            return;
        }
        output[idx] = __halves2bfloat162(input1[idx], input2[idx]);
    }
    __global__ __launch_bounds__(1024) void cast_kernel(bfloat16_t* input1,  bfloat16_t* input2, bfloat16_t* output, uint32_t input_total_length)
    {
        simt_halves2bfloat162(input1, input2, (bfloat16x2_t*)output, input_total_length);
    }
    
  • SIMD与SIMT混合编程场景:
     1
     2
     3
     4
     5
     6
     7
     8
     9
    10
    11
    12
    13
    // 使用小向量可提升数据搬运效率
    __simt_vf__ __launch_bounds__(1024) inline void simt_halves2bfloat162(__gm__ bfloat16_t* input1, __gm__ bfloat16_t* input2, __gm__ bfloat16x2_t* output, uint32_t input_total_length)
    {
        uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
        if (idx > input_total_length) {
            return;
        }
        output[idx] = __halves2bfloat162(input1[idx], input2[idx]);
    }
    __global__ __vector__ void cast_kernel(__gm__ bfloat16_t* input1,  __gm__ bfloat16_t* input2, __gm__ bfloat16_t* output, uint32_t input_total_length)
    {
        asc_vf_call<simt_halves2bfloat162>(dim3(1024), input1, input2, (__gm__ bfloat16x2_t*)output, input_total_length);
    }