PairReduce

函数功能

PairReduceSum：相邻两个（奇偶）元素求和，例如（a1, a2, a3, a4, a5, a6...），相邻两个数据求和为（a1+a2, a3+a4, a5+a6, ......）。

函数原型

参数说明

表1 模板参数说明

参数名

描述

T

操作数数据类型。

isSetMask

是否在接口内部设置mask。

  • true，表示在接口内部设置mask。
  • false，表示在接口外部设置mask，开发者需要使用SetVectorMask接口设置mask值。这种模式下，本接口入参中的mask值必须设置为MASK_PLACEHOLDER。
表2 参数说明

参数名称

输入/输出

含义

dstLocal

输出

目的操作数。

类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。

Atlas 训练系列产品，支持的数据类型为：half

Atlas推理系列产品AI Core，支持的数据类型为：half/float

Atlas A2训练系列产品/Atlas 800I A2推理产品，支持的数据类型为：half/float

Atlas 200I/500 A2推理产品，支持的数据类型为：half/float

srcLocal

输入

源操作数。

类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。

Atlas 训练系列产品，支持的数据类型为：half

Atlas推理系列产品AI Core，支持的数据类型为：half/float

Atlas A2训练系列产品/Atlas 800I A2推理产品，支持的数据类型为：half/float

Atlas 200I/500 A2推理产品，支持的数据类型为：half/float

repeat

输入

重复迭代次数。取值范围为[0, 255]。

矢量计算单元，每次读取连续的256 Bytes数据进行计算，为完成对输入数据的处理，必须通过多次迭代（repeat）才能完成所有数据的读取与计算。repeatTimes表示迭代的次数。

关于该参数的具体描述请参考基础API通用说明。

mask[2]/ maskCount

输入

mask用于控制每次迭代内参与计算的元素。

  • 连续模式：表示前面连续的多少个元素参与计算。数据类型为uint64。取值范围和操作数的数据类型有关，数据类型不同，每次迭代内能够处理的元素个数最大值不同。当操作数为16位时，mask∈[1, 128]；当操作数为32位时，mask∈[1, 64]。
  • 逐bit模式：可以按位控制哪些元素参与计算，bit位的值为1表示参与计算，0表示不参与。参数类型为长度为2的uint64_t类型数组。

    例如，mask=[8, 0]，8=0b1000，表示仅第4个元素参与计算。

    参数取值范围和操作数的数据类型有关，数据类型不同，每次迭代内能够处理的元素个数最大值不同。当操作数为16位时，mask[0]、mask[1]∈[0, 264-1]；当操作数为32位时，mask[1]为0，mask[0]∈[0, 264-1]。

dstRepStride

输入

表示相邻迭代间，矢量目的操作数相同block地址步长。RepStride单位为128Byte。

注意，此参数值Atlas 训练系列产品不支持配置0。

srcBlkStride

输入

表示单次迭代内，矢量源操作数不同block间地址步长。

srcRepStride

输入

表示相邻迭代间，矢量源操作数相同block地址步长。

返回值

无

支持的型号

Atlas 训练系列产品

Atlas推理系列产品AI Core

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

Atlas 200I/500 A2推理产品

注意事项

调用示例

本样例中展示Compute流程中的部分代码。如果您需要运行样例代码，请将该代码段拷贝并替换下方主体代码的Compute函数指令执行部分即可。

#include "kernel_operator.h"

namespace AscendC {
class KernelReduce {
public:
    __aicore__ inline KernelReduce() {}
    __aicore__ inline void Init(__gm__ uint8_t* src, __gm__ uint8_t* dstGm)
    {
        srcGlobal.SetGlobalBuffer((__gm__ half*)src);
        dstGlobal.SetGlobalBuffer((__gm__ half*)dstGm);

        pipe.InitBuffer(inQueueSrc, 1, srcDataSize * sizeof(half));
        pipe.InitBuffer(outQueueDst, 1, dstDataSize * sizeof(half));
    }
    __aicore__ inline void Process()
    {
        CopyIn();
        Compute();
        CopyOut();
    }

private:
    __aicore__ inline void CopyIn()
    {
        LocalTensor<half> srcLocal = inQueueSrc.AllocTensor<half>();
        DataCopy(srcLocal, srcGlobal, srcDataSize);
        inQueueSrc.EnQue(srcLocal);
    }
    __aicore__ inline void Compute()
    {
        LocalTensor<half> srcLocal = inQueueSrc.DeQue<half>();
        LocalTensor<half> dstLocal = outQueueDst.AllocTensor<half>();

        half zero(0);
        Duplicate(dstLocal, zero, dstDataSize);

        //指令执行部分（替换成上述代码）

        outQueueDst.EnQue<half>(dstLocal);
        inQueueSrc.FreeTensor(srcLocal);
    }
    __aicore__ inline void CopyOut()
    {
        LocalTensor<half> dstLocal = outQueueDst.DeQue<half>();
        DataCopy(dstGlobal, dstLocal, dstDataSize);
        outQueueDst.FreeTensor(dstLocal);
    }

private:
    TPipe pipe;
    TQue<QuePosition::VECIN, 1> inQueueSrc;
    TQue<QuePosition::VECOUT, 1> outQueueDst;
    GlobalTensor<half> srcGlobal, dstGlobal;
    int srcDataSize = 128;
    int dstDataSize = 64;
};
} // namespace AscendC

extern "C" __global__ __aicore__ void reduce_simple_kernel(__gm__ uint8_t* src, __gm__ uint8_t* dstGm)
{
    AscendC::KernelReduce op;
    op.Init(src, dstGm);
    op.Process();
}
父主题： 归约指令