norm4df

产品支持情况

• Ascend 950PR/Ascend 950DT：支持

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

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

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

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

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

功能说明

获取输入数据a、b、c、d的平方和a^2 + b^2+ c^2+ d^2的平方根。

函数原型

C++

inline float norm4df(float a, float b, float c, float d)

参数说明

表1 参数说明

参数名	输入/输出	描述
a	输入	源操作数。
b	输入	源操作数。
c	输入	源操作数。
d	输入	源操作数。

返回值说明

a^2 + b^2 + c^2 + d^2的平方根。

• 若a^2 + b^2 + c^2 + d^2的平方根超出float最大范围，返回值为inf。
• 若a, b, c, d任意一个或多个为±inf，返回值为inf。
• 若a, b, c, d任意一个或多个为nan同时不是±inf，返回值为nan。

约束说明

针对Ascend 950PR/Ascend 950DT，本接口不支持Subnormal场景：本接口内部实现使用到了除法运算符，由于除法运算符不支持Subnormal场景，当所有输入均为Subnormal数据时，本接口最终返回nan。

需要包含的头文件

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

C++

#include "simt_api/math_functions.h"

调用示例

• SIMT编程场景：

  C++

  __global__ __launch_bounds__(256) void compute_norm4df(float *result, const float *x, const float *y, const float *z, const float *w, uint32_t count)
  {
      const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
      if (idx >= count) {
          return;
      }
      result[idx] = norm4df(x[idx], y[idx], z[idx], w[idx]);
  }
  

• SIMD与SIMT混合编程场景：

  C++

  __simt_vf__ __launch_bounds__(256) inline void compute_norm4df_vf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, __gm__ const float *z, __gm__ const float *w, uint32_t count)
  {
      const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
      if (idx >= count) {
          return;
      }
      result[idx] = norm4df(x[idx], y[idx], z[idx], w[idx]);
  }
  
  __global__ __vector__ void run_norm4df(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, __gm__ const float *z, __gm__ const float *w, uint32_t count)
  {
      asc_vf_call<compute_norm4df_vf>(dim3(256), result, x, y, z, w, count);
  }
  

输入输出示例如下：

Text

x：0.25, 0.75, 1.25, 1.75
y：1.5, 2.5, 3.5, 4.5
z：-0.5, 0.5, 1.5, 2.5
w：2, 3, 4, 5
result: 2.561738 4.007805 5.662376 7.386643