开发者
下载

__hmaxx2

产品支持情况

产品

是否支持

Atlas 350 加速卡

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

功能说明

获取两个half2类型数据各分量的最大值。

函数原型

1
half2 __hmaxx2(const half2 x, const half2 y)

参数说明

表1 参数说明

参数名

输入/输出

描述

x

输入

源操作数。

y

输入

源操作数。

返回值说明

输入数据各分量的最大值。比较的分量a和b满足:

  • a为+0,b为-0时,返回a。
  • a为-0,b为+0时,返回b。
  • a为nan时,返回值为b。
  • b为nan时,返回值为a。
  • a,b同时为nan时,返回值为nan。
  • a,b任意一个为inf时,返回值为inf。
  • a为-inf时,返回值为b。
  • b为-inf时,返回值为a。

约束说明

需要包含的头文件

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

1
#include "simt_api/asc_fp16.h"

调用示例

  • SIMT编程场景:
     1
     2
     3
     4
     5
     6
     7
     8
     9
    10
    11
    12
    13
    // 使用短向量可提升数据搬运效率
    __global__ __launch_bounds__(1024) void simt_hmaxx2(half* x, half* y, half* dst, uint32_t input_total_length)
    {
        uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
        // 每个线程处理1个half2类型的数据,即2个half类型的数据,因此idx >= input_total_length / 2的线程不处理数据
        if (idx >= input_total_length / 2) {
            return;
        }
        half2* input1 = (half2*)x;
        half2* input2 = (half2*)y;
        half2* out = (half2*)dst;    
        out[idx] = __hmaxx2(input1[idx], input2[idx]);
    }
    
  • SIMD与SIMT混合编程场景:
     1
     2
     3
     4
     5
     6
     7
     8
     9
    10
    11
    12
    13
    14
    15
    // 使用短向量可提升数据搬运效率
    __simt_vf__ __launch_bounds__(1024) inline void simt_hmaxx2(__gm__ half2* x, __gm__ half2* y, __gm__ half2* dst, uint32_t input_total_length)
    {
        uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
        // 每个线程处理1个half2类型的数据,即2个half类型的数据,因此idx >= input_total_length / 2的线程不处理数据
        if (idx >= input_total_length / 2) {
            return;
        } 
        dst[idx] = __hmaxx2(x[idx], y[idx]);
    }
    
    __global__ __vector__ void compare_kernel(__gm__ half* x, __gm__ half* y, __gm__ half* dst, uint32_t input_total_length)
    {
        asc_vf_call<simt_hmaxx2>(dim3(1024), (__gm__ half2*)x, (__gm__ half2*)y, (__gm__ half2*)dst, input_total_length);
    }