__low2bfloat16

产 品 支 持 情 况

• Ascend 950PR/Ascend 950DT：支 持

• Atlas A3 训 练 系 列 产 品/Atlas A3 推 理 系 列 产 品：不 支 持

• Atlas A2 训 练 系 列 产 品/Atlas A2 推 理 系 列 产 品：不 支 持

• Atlas 200I/500 A2 推 理 产 品：不 支 持

• Atlas 推 理 系 列 产 品AI Core：不 支 持
• Atlas 推 理 系 列 产 品Vector Core：不 支 持

• Atlas 训 练 系 列 产 品：不 支 持

功 能 说 明

返 回 输 入 数 据 的 低16位。

函 数 原 型

Text

inline bfloat16_t __low2bfloat16(const bfloat16x2_t x)

参 数 说 明

表1 参 数 说 明

参 数 名	输 入/输 出	描 述
x	输 入	源 操 作 数。

返 回 值 说 明

输 入 数 据 的 低16位。特 殊 值 如 下：

x低16位 值	返 回 值
0	0
-0	-0
nan	nan
inf	inf
-inf	-inf
ASCRT_MAX_NORMAL_BF16	ASCRT_MAX_NORMAL_BF16
-ASCRT_MAX_NORMAL_BF16	-ASCRT_MAX_NORMAL_BF16

约 束 说 明

无

需 要 包 含 的 头 文 件

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

Text

#include "simt_api/asc_bf16.h"

调 用 示 例

• SIMT编 程 场 景：

  Text

  // 使用短向量可提升数据搬运效率
  __aicore__ void simt_low2bfloat16(bfloat16x2_t* input, bfloat16_t* output, uint32_t input_total_length)
  {
      uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
      // 每个线程处理1个bfloat16x2_t类型的数据，即2个bfloat16_t类型的数据，因此idx >= input_total_length / 2的线程不处理数据
      if (idx > input_total_length / 2) {
          return;
      }
      output[idx] = __low2bfloat16(input[idx]);
  }
  __global__ __launch_bounds__(1024) void cast_kernel(bfloat16_t* input, bfloat16_t* output, uint32_t input_total_length)
  {
      simt_low2bfloat16((bfloat16x2_t*)input, output, input_total_length);
  }
  

• SIMD与SIMT混 合 编 程 场 景：

  Text

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