BlockReduceMax
函数功能
对每个block内所有元素求最大值。
函数原型
- mask参数使用逐bit模式,该模式的具体介绍请参考参数说明中的mask参数说明:
template <typename T, bool isSetMask = true> __aicore__ inline void BlockReduceMax(const LocalTensor<T>& dstLocal, const LocalTensor<T>& srcLocal,const int32_t repeat, const uint64_t mask[2], const int32_t dstRepStride, const int32_t srcBlkStride,const int32_t srcRepStride)
- mask参数使用连续模式,该模式的具体介绍请参考参数说明中的mask参数说明:
template <typename T, bool isSetMask = true> __aicore__ inline void BlockReduceMax(const LocalTensor<T>& dstLocal, const LocalTensor<T>& srcLocal,const int32_t repeat, const int32_t maskCount, const int32_t dstRepStride, const int32_t srcBlkStride,const int32_t srcRepStride)
参数说明
参数名 |
描述 |
|---|---|
T |
操作数数据类型。 |
isSetMask |
是否在接口内部设置mask。
|
参数名称 |
输入/输出 |
含义 |
|---|---|---|
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表示迭代的次数。 关于该参数的具体描述请参考重复迭代次数-Repeat times。 |
mask[2]/ maskCount |
输入 |
mask用于控制每次迭代内参与计算的元素。
|
dstRepStride |
输入 |
表示相邻迭代间,矢量目的操作数相邻迭代间相同block的地址步长。 输入类型为half类型时,RepStride单位为16Byte,输入类型为float类型时,RepStride单位为32Byte。 注意,此参数值Atlas 训练系列产品不支持配置0。 |
srcBlkStride |
输入 |
表示单次迭代内,矢量源操作数单次迭代内不同block间地址步长。 |
srcRepStride |
输入 |
表示相邻迭代间,矢量源操作数相邻迭代间相同block的地址步长。 |
返回值
无
支持的型号
Atlas 训练系列产品
Atlas推理系列产品AI Core
Atlas A2训练系列产品/Atlas 800I A2推理产品
Atlas 200I/500 A2推理产品
注意事项
- 为了节省地址空间,您可以定义一个Tensor,供源操作数与目的操作数同时使用(即地址重叠),需要注意计算后的目的操作数数据不能覆盖未参与计算的源操作数,需要谨慎使用。
- 对于Atlas 200I/500 A2推理产品,若配置mask[2]/maskCount参数后,存在某个block里的任何一个元素都不参与计算,则该block内所有元素的最大值会填充为-inf返回。比如float场景下,当maskCount配置为32,即只计算前4个block,则后四个block内的最大值会返回-inf。half场景下会返回-65504。
- 操作数地址偏移对齐要求请参见通用约束。
调用示例
本样例中只展示Compute流程中的部分代码。如果您需要运行样例代码,请将该代码段拷贝并替换样例模板中Compute函数的部分代码即可。
- BlockReduceMax-tensor高维切分计算样例-mask连续模式
uint64_t mask = 256/sizeof(half); int repeat = 1; // repeat = 1, 128 elements one repeat, 128 elements total // srcBlkStride = 1, no gap between blocks in one repeat // dstRepStride = 1, srcRepStride = 8, no gap between repeats BlockReduceMax<half>(dstLocal, srcLocal, repeat, mask, 1, 1, 8);
- BlockReduceMax-tensor高维切分计算样例-mask逐bit模式
uint64_t mask[2] = { UINT64_MAX, UINT64_MAX }; int repeat = 1; // repeat = 1, 128 elements one repeat, 128 elements total // srcBlkStride = 1, no gap between blocks in one repeat // dstRepStride = 1, srcRepStride = 8, no gap between repeats BlockReduceMax<half>(dstLocal, srcLocal, repeat, mask, 1, 1, 8);
输入数据(src_gm): [-8.781, 4.688, -0.09607, -5.445, 4.957, -4.832, 9.555, 8.391, 6.273, -2.412, 7.969, 3.9, -0.4238, 2.988, -6.855, -1.335, ... 9.68, -6.672, -6.488, -7.398, 8.562, 3.508, 3.135, -5.512, -7.883, -8.594, -5.895, -8.938, -7.676, -7.867, -9.188, -5.715] 输出数据(dst_gm): [9.555, ..., 9.68, 0, ... 0]
样例模板
#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();
}