在所有的输入数据中找出最大值及最大值对应的索引位置。
接口级别 |
原型定义 |
|---|---|
0级接口 |
|
2级接口 |
template <typename T> __aicore__ inline void ReduceMax(const LocalTensor<T>& dstLocal, const LocalTensor<T>& srcLocal, const LocalTensor<T>& workLocal, const int32_t count, bool calIndex = 0) |
参数名称 |
输入/输出 |
含义 |
|---|---|---|
dstLocal |
输出 |
目的操作数。 类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。 Atlas 训练系列产品，支持的数据类型为：half Atlas推理系列产品AI Core，支持的数据类型为：half/float Atlas A2训练系列产品，支持的数据类型为：half/float |
srcLocal |
输入 |
源操作数。 类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。 源操作数的数据类型需要与目的操作数保持一致。 Atlas 训练系列产品，支持的数据类型为：half Atlas推理系列产品AI Core，支持的数据类型为：half/float Atlas A2训练系列产品，支持的数据类型为：half/float |
workLocal |
输入 |
本指令执行期间，需要一块空间用于存储中间结果，需特别注意空间大小，具体请参考指令约束说明中空间大小的要求说明。 类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。 源操作数的数据类型需要与目的操作数保持一致。 Atlas 训练系列产品，支持的数据类型为：half Atlas推理系列产品AI Core，支持的数据类型为：half/float Atlas A2训练系列产品，支持的数据类型为：half/float |
mask |
输入 |
mask用于控制每次迭代内参与计算的元素。
|
repeatTimes |
输入 |
重复迭代次数。 矢量计算单元，每次读取连续的256 Bytes数据进行计算，为完成对输入数据的处理，必须通过多次迭代（repeat）才能完成所有数据的读取与计算。repeatTimes表示迭代的次数。关于该参数的具体描述请参考重复迭代次数-Repeat times。 |
srcRepStride |
输入 |
相邻迭代间，源操作数相同block地址步长，即源操作数每次迭代跳过的block数目。详细说明请参考相邻迭代间相同block的地址步长。 |
calIndex |
输入 |
指定是否获取最大值的索引，bool类型，默认值为false，取值：
|
参数名称 |
输入/输出 |
含义 |
|---|---|---|
dstLocal |
输出 |
目的操作数。 类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。 Atlas 训练系列产品，支持的数据类型为：half Atlas推理系列产品AI Core，支持的数据类型为：half/float Atlas A2训练系列产品，支持的数据类型为：half/float |
srcLocal |
输入 |
源操作数。 类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。 源操作数的数据类型需要与目的操作数保持一致。 Atlas 训练系列产品，支持的数据类型为：half Atlas推理系列产品AI Core，支持的数据类型为：half/float Atlas A2训练系列产品，支持的数据类型为：half/float |
workLocal |
输入 |
源操作数。本指令执行期间，需要一块空间用于存储中间结果，需特别注意空间大小，具体请参考指令约束说明中空间大小的要求说明。 类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。 源操作数的数据类型需要与目的操作数保持一致。 Atlas 训练系列产品，支持的数据类型为：half Atlas推理系列产品AI Core，支持的数据类型为：half/float Atlas A2训练系列产品，支持的数据类型为：half/float |
calIndex |
输入 |
指定是否获取最大值的索引，仅支持bool类型，默认值为false，取值：
|
count |
输入 |
输入数据元素个数。 参数取值范围和操作数的数据类型有关，数据类型不同，能够处理的元素个数最大值不同，但是最大处理的数据量不会超过UB大小限制。 |
无
Atlas 训练系列产品
Atlas推理系列产品AI Core
Atlas A2训练系列产品
// 先定义一个向上取整函数
int RoundUp(int a, int b)
{
return (a + b - 1) / b;
}
// 然后定义参与计算的数据类型
int typeSize = 2; // half类型为2Bytes，float类型为4Bytes，按需填入
// 再根据数据类型定义两个单位
int elementsPerBlock = 32 / typeSize; // 1个block存放的元素个数
int elementsPerRepeat = 256 / typeSize; // 1次repeat可以处理的元素个数
// 最后确定首次最大repeat值
int firstMaxRepeat = repeatTimes; // 此处需要注意：对于0级接口，firstMaxRepeat就是repeatTimes；对于2级接口，firstMaxRepeat为count/elementsPerRepeat，比如在half类型下firstMaxRepeat就是count/128，在float类型下为count/64，按需填入
对于calIndex = false的情况，只需一轮操作就可以计算出最小所需空间：
int iter1OutputCount = firstMaxRepeat * 2; // 第一轮操作产生的元素个数
int iter1AlignEnd = RoundUp(iter1OutputCount, elementsPerBlock) * elementsPerBlock; // 第一轮产生的元素个数做向上取整
int finalWorkLocalNeedSize = iter1AlignEnd; // 最终workLocal所需的空间大小
对于calIndex = true的情况，分为三轮操作计算出最小所需空间：
int iter1OutputCount = firstMaxRepeat * 2; // 第一轮操作产生的元素个数
int iter2AlignStart = RoundUp(iter1OutputCount, elementsPerBlock) * elementsPerBlock; // 第二轮操作起始位置偏移，即第一轮产生的元素个数做向上取整后的结果
int iter2OutputCount = RoundUp(iter1OutputCount, elementsPerRepeat) * 2; // 第二轮操作产生的元素个数
int iter3AlignStart = RoundUp(iter2OutputCount, elementsPerBlock) * elementsPerBlock; // 第三轮操作起始位置偏移，即第二轮产生的元素个数做向上取整后的结果
int iter3OutputCount = RoundUp(iter2OutputCount, elementsPerRepeat) * 2; // 第三轮操作产生的元素个数
int iter3AlignEnd = RoundUp(iter3OutputCount, elementsPerBlock) * elementsPerBlock; // 第三轮产生的元素个数做向上取整
int finalWorkLocalNeedSize = iter2AlignStart + iter3AlignStart + iter3AlignEnd; // 最终workLocal所需的空间大小
以上计算出来的最终的空间大小单位是元素个数,若转成Bytes数表示就是finalWorkLocalNeedSize * typeSize (Bytes)
具体计算示例请参考调用示例中workLocal空间计算示例
// dstLocal,srcLocal和workLocal均为half类型,srcLocal的计算数据量为8320,并且连续排布，需要索引值，使用0级接口，设定repeatTimes为65，mask为全部元素参与计算 uint64_t mask = 128; ReduceMax<half>(dstLocal, srcLocal, workLocal, mask, 65, 8, true);
// dstLocal,srcLocal和workLocal均为half类型,srcLocal的计算数据量为8320,并且连续排布，需要索引值，使用0级接口，设定repeatTimes为65,mask为全部元素参与计算
uint64_t mask[2] = { 0xFFFFFFFFFFFFFFFF, 0xFFFFFFFFFFFFFFFF };
ReduceMax<half>(dstLocal, srcLocal, workLocal, mask, 65, 8, true);
// dstLocal,srcLocal和workLocal均为half类型,srcLocal的计算数据量为8320,并且连续排布，需要索引值，使用2级接口 ReduceMax<half>(dstLocal, srcLocal, workLocal, 8320, true);
ReduceMax接口workLocal计算示例一: dstLocal,srcLocal和workLocal均为half类型,srcLocal的计算数据量为8320, 使用0级接口, repeatTimes为65, mask为128,需要索引值 0级接口调用示例为: ReduceMax<half>(dstLocal, srcLocal, workLocal, 128, 65, 8, true); 此时workLocal所需的最小空间计算过程为: int typeSize = 2; int elementsPerBlock = 32 / typeSize = 16; int elementsPerRepeat = 256 / typeSize = 128; int firstMaxRepeat = repeatTimes; int iter1OutputCount = firstMaxRepeat * 2 = 130; // 第一轮操作产生的元素个数 int iter2AlignStart = RoundUp(iter1OutputCount, elementsPerBlock)*elementsPerBlock = 144; // 对第一轮操作输出个数向上取整 int iter2OutputCount = RoundUp(iter1OutputCount, elementsPerRepeat)*2 = 4; // 第二轮操作产生的元素个数 int iter3AlignStart = RoundUp(iter2OutputCount, elementsPerBlock)*elementsPerBlock = 16; // 对第二轮操作输出个数向上取整 int iter3OutputCount = RoundUp(iter2OutputCount, elementsPerRepeat)*2 = 2; // 第三轮操作产生的元素个数 int iter3AlignEnd = RoundUp(iter3OutputCount, elementsPerBlock) * elementsPerBlock = 16; // 第三轮产生的元素个数做向上取整 最终workLocal所需的最小空间就是 iter2AlignStart + iter3AlignStart + iter3AlignEnd = 144 + 16 + 16 = 176 ,也就是352Bytes ReduceMax接口workLocal计算示例二: dstLocal,srcLocal和workLocal均为half类型,srcLocal的计算数据量为32640, 使用0级接口,repeatTimes为255, mask为128,需要索引值 0级接口调用示例为: ReduceMax<half>(dstLocal, srcLocal, workLocal, 128, 255, 8, true); 此时workLocal所需的最小空间计算过程为: int typeSize = 2; int elementsPerBlock = 32 / typeSize = 16; int elementsPerRepeat = 256 / typeSize = 128; int firstMaxRepeat = repeatTimes; int iter1OutputCount = firstMaxRepeat * 2 = 510; // 第一轮操作产生的元素个数 int iter2AlignStart = RoundUp(iter1OutputCount, elementsPerBlock)*elementsPerBlock = 512; // 对第一轮操作输出个数向上取整 int iter2OutputCount = RoundUp(iter1OutputCount, elementsPerRepeat)*2 = 8; // 第二轮操作产生的元素个数 int iter3AlignStart = RoundUp(iter2OutputCount, elementsPerBlock)*elementsPerBlock = 16; // 对第二轮操作输出个数向上取整 int iter3OutputCount = RoundUp(iter2OutputCount, elementsPerRepeat)*2 = 2; // 第三轮操作产生的元素个数 int iter3AlignEnd = RoundUp(iter3OutputCount, elementsPerBlock) * elementsPerBlock = 16; // 第三轮产生的元素个数做向上取整 需要的空间就是 iter2AlignStart + iter3AlignStart + iter3OutputCount = 512 + 16 + 16 = 544 ,也就是1088Bytes ReduceMax接口workLocal计算示例三: dstLocal,srcLocal和workLocal均为half类型,srcLocal的计算数据量为65408,使用2级接口,count=65408,需要索引值 2级接口调用示例为: ReduceMax<half>(dstLocal, srcLocal, workLocal, 65408, true); 此时workLocal所需的最小空间计算过程为: int typeSize = 2; int elementsPerBlock = 32 / typeSize = 16; int elementsPerRepeat = 256 / typeSize = 128; int firstMaxRepeat = count / elementsPerRepeat = 511; int iter1OutputCount = firstMaxRepeat * 2 = 1022; // 第一轮操作产生的元素个数 int iter2AlignStart = RoundUp(iter1OutputCount, elementsPerBlock)*elementsPerBlock = 1024; // 对iter1OutputCount输出个数向上取整 int iter2OutputCount = RoundUp(iter1OutputCount, elementsPerRepeat)*2 = 16; // 第二轮操作产生的元素个数 int iter3AlignStart = RoundUp(iter2OutputCount, elementsPerBlock)*elementsPerBlock = 16; // 对iter2OutputCount输出个数向上取整 int iter3OutputCount = RoundUp(iter2OutputCount, elementsPerRepeat)*2 = 2; // 第三轮操作产生的元素个数 int iter3AlignEnd = RoundUp(iter3OutputCount, elementsPerBlock) * elementsPerBlock = 16; // 第三轮产生的元素个数做向上取整 需要的空间就是 iter2AlignStart + iter3AlignStart + iter3OutputCount = 1024 + 16 + 16 = 1056,也就是2112Bytes ReduceMax接口workLocal计算示例四: dstLocal,srcLocal和workLocal均为half类型,srcLocal的的计算数据量为512,使用0级接口,repeatTimes为4, mask为128,需要索引值 0级接口调用示例为: ReduceMax<half>(dstLocal, srcLocal, workLocal, 128, 4, 8, true); 此时workLocal所需的最小空间计算过程为: int typeSize = 2; int elementsPerBlock = 32 / typeSize = 16; int elementsPerRepeat = 256 / typeSize = 128; int firstMaxRepeat = repeatTimes; int iter1OutputCount = firstMaxRepeat * 2 = 8; // 第一轮操作产生的元素个数 int iter2AlignStart = RoundUp(iter1OutputCount, elementsPerBlock)*elementsPerBlock = 16; // 对iter1OutputCount输出个数向上取整 int iter2OutputCount = RoundUp(iter1OutputCount, elementsPerRepeat)*2 = 2; // 第二轮操作产生的元素个数 本用例中，由于第二轮操作产生的元素个数为2，即第二轮结束就可以拿到最大值和其索引值，因此需要的空间就是iter2AlignStart + RoundUp(iter2OutputCount, elementsPerBlock) * elementsPerBlock = 16 + 16 = 32，也就是64Bytes ReduceMax接口workLocal计算示例五: dstLocal,srcLocal和workLocal均为half类型,srcLocal的计算数据量count为65408,使用2级接口,count=65408，不需要索引值 2级接口调用示例为: ReduceMax<half>(dstLocal, srcLocal, workLocal, 65408, false); 此时workLocal所需的最小空间计算过程为: int typeSize = 2; int elementsPerBlock = 32 / typeSize = 16; int elementsPerRepeat = 256 / typeSize = 128; int firstMaxRepeat = count / elementsPerRepeat = 511; int iter1OutputCount = firstMaxRepeat * 2 = 1022; // 第一轮操作产生的元素个数 int iter1AlignEnd = RoundUp(iter1OutputCount, elementsPerBlock) * elementsPerBlock = 1024; // 第一轮产生的元素个数做向上取整 由于calIndex为false,因此最终workLocal所需的最小空间大小就是对第一轮产生元素做向上取整后的结果，此处就是1024，也就是2048Bytes ReduceMax接口workLocal计算示例六: dstLocal,srcLocal和workLocal均为float类型,srcLocal的计算数据量为8320, 使用0级接口, repeatTimes为130, mask为64,需要索引值 0级接口调用示例为: ReduceMax<float>(dstLocal, srcLocal, workLocal, 64, 130, 8, true); 此时workLocal所需的最小空间计算过程为: int typeSize = 4; int elementsPerBlock = 32 / typeSize = 8; int elementsPerRepeat = 256 / typeSize = 64; int firstMaxRepeat = repeatTimes; int iter1OutputCount = firstMaxRepeat * 2 = 260; // 第一轮操作产生的元素个数 int iter2AlignStart = RoundUp(iter1OutputCount, elementsPerBlock)*elementsPerBlock = 264; // 对第一轮操作输出个数向上取整 int iter2OutputCount = RoundUp(iter1OutputCount, elementsPerRepeat)*2 = 10; // 第二轮操作产生的元素个数 int iter3AlignStart = RoundUp(iter2OutputCount, elementsPerBlock)*elementsPerBlock = 16; // 对第二轮操作输出个数向上取整 int iter3OutputCount = RoundUp(iter2OutputCount, elementsPerRepeat)*2 = 2; // 第三轮操作产生的元素个数 int iter3AlignEnd = RoundUp(iter3OutputCount, elementsPerBlock) * elementsPerBlock = 16; // 第三轮产生的元素个数做向上取整 最终workLocal所需的最小空间就是 iter2AlignStart + iter3AlignStart + iter3AlignEnd = 264 + 16 + 16 = 296 ,也就是1184Bytes
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);
repeat = srcDataSize / mask;
pipe.InitBuffer(inQueueSrc, 1, srcDataSize * sizeof(half));
pipe.InitBuffer(workQueue, 1, 32 * sizeof(half)); // 此处按照公式计算所需的最小work空间为32，也就是64Bytes
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>();
LocalTensor<half> workLocal = workQueue.AllocTensor<half>();
// level0
ReduceMax<half>(dstLocal, srcLocal, workLocal, mask, repeat, repStride, true);
outQueueDst.EnQue<half>(dstLocal);
inQueueSrc.FreeTensor(srcLocal);
workQueue.FreeTensor(workLocal);
}
__aicore__ inline void CopyOut()
{
LocalTensor<half> dstLocal = outQueueDst.DeQue<half>();
DataCopy(dstGlobal, dstLocal, srcDataSize);
outQueueDst.FreeTensor(dstLocal);
}
private:
TPipe pipe;
TQue<QuePosition::VECIN, 1> inQueueSrc;
TQue<QuePosition::VECOUT, 1> workQueue;
TQue<QuePosition::VECOUT, 1> outQueueDst;
GlobalTensor<half> srcGlobal, dstGlobal;
int srcDataSize = 512;
int dstDataSize = 512;
int mask = 128;
int repStride = 8;
int repeat = 0;
};
} // namespace AscendC
extern "C" __global__ __aicore__ void kernel_ReduceMax_lv0_half_512(__gm__ uint8_t* src, __gm__ uint8_t* dstGm)
{
AscendC::KernelReduce op;
op.Init(src, dstGm);
op.Process();
}
示例结果 输入数据(src_gm):
[0.4795 0.951 0.866 0.008545 0.8037 0.551 0.754 0.73 0.6035 0.251 0.4841 0.05914 0.9414 0.379 0.664 0.6914 0.9307 0.3853 0.4048
0.7754 0.1265 0.709 0.7695 0.8057 0.9673 0.2566 0.8696 0.243 0.871 0.123 0.76 0.1844 0.7324 0.5757 0.0172 0.7188 0.556 0.3699
0.7334 0.655 0.919 0.4219 0.82 0.1046 0.5796 0.4773 0.1405 0.3777 0.4421 0.983 0.728 0.642 0.37 0.9473 0.52 0.7783 0.699
0.716 0.1791 0.1272 0.2471 0.3298 0.3518 0.9756 0.2268 0.6167 0.742 0.4185 0.8193 0.919 0.03827 0.02957 0.2598 0.798 0.3752 0.2109
0.1753 0.7227 0.829 0.6978 0.347 0.463 0.685 0.1992 0.847 0.941 0.835 0.03336 0.1359 0.04736 0.758 0.5347 0.616 0.869 0.582
0.694 0.2035 0.3613 0.8413 0.68 0.0896 0.3833 0.0768 0.292 0.11053 0.5586 0.578 0.3286 0.09314 0.5845 0.7124 0.2058 0.6523 0.784
0.9985 0.6626 0.8975 0.405 0.884 0.7744 0.0258 0.484 0.768 0.7197 0.577 0.03143 0.9185 0.3608 0.3352 0.9077 0.709 0.85 0.4607
0.61 0.4277 0.1004 0.1995 0.1608 0.2852 0.8887 0.813 0.3396 0.272 0.703 0.1312 0.734 0.2612 0.6895 0.8647 0.9165 0.1455 0.9233
0.3027 0.7163 0.927 0.1995 0.155 0.6953 0.66 0.04163 0.99 0.544 0.4243 0.804 0.4612 0.01912 0.5127 0.8755 0.6665 0.707 0.01018
0.874 0.8545 0.9375 0.9844 0.578 0.934 0.683 0.4668 0.63 0.2032 0.3188 0.9478 0.9375 0.03357 0.9927 0.996 0.451 0.1105 0.762
0.82 0.8047 0.911 0.926 0.1973 0.9175 0.4521 0.4487 0.1273 0.718 0.737 0.305 0.922 0.1396 0.618 0.753 0.5913 0.874 0.08905
0.003582 0.05252 0.674 0.3923 0.527 0.4106 0.7812 0.113 0.965 0.6157 0.4368 0.6646 0.7944 0.7964 0.531 0.6665 0.517 0.04468 0.5737
0.752 0.4 0.4463 0.05496 0.939 0.6353 0.2036 0.667 0.3994 0.2573 0.118 0.973 0.5923 0.558 0.7114 0.785 0.714 0.7485 0.854
0.2585 0.274 0.9824 0.4158 0.283 0.2194 0.3074 0.2793 0.531 0.8965 0.01456 0.5264 0.992 0.856 0.5986 0.734 0.908 0.12317 0.8374
0.6665 0.1904 0.97 0.2546 0.364 0.6914 0.462 0.05353 0.02975 0.6235 0.4941 0.4714 0.788 0.06537 0.8423 0.2527 0.7734 0.591 0.443
0.3022 0.02116 0.01605 0.772 0.6924 0.01032 0.594 0.1865 0.7393 0.8887 0.916 0.9653 0.696 0.901 0.1255 0.5513 0.2742 0.5586 0.988
0.0954 0.4365 0.677 0.894 0.8413 0.05655 0.932 0.4426 0.336 0.848 0.9434 0.1976 0.813 0.773 0.2605 0.1543 0.8555 0.3596 0.997
0.10315 0.5796 0.5327 0.2283 0.7583 0.3674 0.513 0.9126 0.751 0.532 0.399 0.832 0.549 0.2358 0.6655 0.477 0.5864 0.3528 0.989
0.1412 0.748 0.3652 0.05292 0.3552 0.5767 0.826 0.4792 0.8477 0.03488 0.8267 0.2345 0.931 0.0884 0.6816 0.4685 0.618 0.09973 0.4385
0.782 0.6465 0.03882 0.4158 0.1422 0.822 0.8203 0.95 0.3274 0.724 0.929 0.8726 0.004307 0.815 0.67 0.4368 0.7793 0.593 0.4663
0.2207 0.01773 0.39 0.008896 0.4238 0.716 0.1155 0.601 0.9214 0.3708 0.4285 0.951 0.00431 0.726 0.977 0.1254 0.6484 0.4648 0.891
0.723 0.6333 0.9077 0.4849 0.3008 0.0495 0.4575 0.266 0.2014 0.1106 0.6914 0.2744 0.4956 0.532 0.1752 0.709 0.3464 0.6104 0.4067
0.1317 0.8647 0.8 0.4832 0.013855 0.6733 0.4524 0.6865 0.7017 0.9385 0.2957 0.2444 0.4167 0.55 0.8926 0.8364 0.506 0.9966 0.7207
0.51 0.8745 0.3188 0.847 0.86 0.64 0.08453 0.59 0.2062 0.1031 0.1459 0.3806 0.2096 0.469 0.1492 0.10065 0.536 0.572 0.353
0.068 0.07855 0.6177 0.3408 0.1538 0.2732 0.997 0.1158 0.4028 0.9536 0.7197 0.585 0.0899 0.3994 0.1835 0.737 0.4639 0.3071 0.47
0.993 0.3862 0.293 0.1813 0.8193 0.745 0.064 0.7407 0.329 0.198 0.596 0.3 0.6562 0.819 0.2803 0.04095 0.703 0.3425 0.9224
0.776 0.8057 0.734 0.2534 0.1824 0.793 0.3542 0.2595 0.2607 0.838 0.39 0.631 0.3542 0.1968 0.643 0.015366 0.4106 0.604 ]
输出数据(dst_gm):
[0.9985, 6.8e-06], 6.8e-06使用reinterpret_cast方法转换后就是索引值114
2级接口完整调用示例:
#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);
repeat = srcDataSize / mask;
pipe.InitBuffer(inQueueSrc, 1, srcDataSize * sizeof(half));
pipe.InitBuffer(workQueue, 1, 32 * sizeof(half)); // 此处按照公式计算所需的最小work空间为32，也就是64Bytes
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>();
LocalTensor<half> workLocal = workQueue.AllocTensor<half>();
// level2
ReduceMax<half>(dstLocal, srcLocal, workLocal, srcDataSize, true);
outQueueDst.EnQue<half>(dstLocal);
inQueueSrc.FreeTensor(srcLocal);
workQueue.FreeTensor(workLocal);
}
__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> workQueue;
TQue<QuePosition::VECOUT, 1> outQueueDst;
GlobalTensor<half> srcGlobal, dstGlobal;
int srcDataSize = 288;
int dstDataSize = 16;
int mask = 128;
int repStride = 8;
int repeat = 0;
};
} // namespace AscendC
extern "C" __global__ __aicore__ void kernel_ReduceMax_lv2_half_288(__gm__ uint8_t* src, __gm__ uint8_t* dstGm)
{
AscendC::KernelReduce op;
op.Init(src, dstGm);
op.Process();
}
示例结果 输入数据(src_gm):
[0.4778 0.5903 0.2433 0.698 0.1943 0.407 0.891 0.1766 0.5977 0.9473 0.6523 0.10913 0.0143 0.86 0.2366 0.625 0.3696 0.708 0.946
0.538 0.3826 0.08215 0.516 0.9116 0.1548 0.507 0.8145 0.89 0.5435 0.563 0.1125 0.543 0.3142 0.8086 0.6885 0.874 0.855 0.4019
0.1613 0.04462 0.945 0.6064 0.6904 0.00758 0.9463 0.528 0.9966 0.629 0.714 0.03134 0.4407 0.0322 0.5376 0.04443 0.03778 0.522 0.793
0.3086 0.4 0.3984 0.5693 0.8203 0.673 0.796 0.2747 0.2246 0.468 0.1146 0.4468 0.419 0.3816 0.1636 0.1414 0.4028 0.9785 0.8984
0.4355 0.874 0.864 0.7856 0.739 0.895 0.2487 0.5034 0.958 0.661 0.8755 0.302 0.802 0.563 0.9067 0.1562 0.1337 0.1844 0.3047
0.543 0.3855 0.9536 0.8633 0.5435 0.002748 0.8916 0.9614 0.3665 0.1588 0.51 0.77 0.552 0.84 0.2798 0.7217 0.8633 0.3794 0.5376
0.03 0.7783 0.9297 0.9556 0.609 0.1776 0.5957 0.2954 0.6675 0.7183 0.4182 0.8804 0.1837 0.3235 0.3486 0.43 0.8633 0.3972 0.1307
0.7915 0.43 0.2544 0.827 0.04843 0.1637 0.3376 0.4087 0.4993 0.5923 0.3057 0.04306 0.4905 0.693 0.7393 0.777 0.01379 0.2742 0.669
0.6826 0.04028 0.0423 0.281 0.12476 0.5366 0.2098 0.559 0.8833 0.82 0.0745 0.7485 0.04004 0.776 0.863 0.1909 0.7876 0.734 0.4727
0.3655 0.944 0.006794 0.01872 0.687 0.5664 0.9697 0.2437 0.2014 0.0269 0.3975 0.08405 0.36 0.0751 0.02632 0.135 0.531 0.554 0.378
0.9365 0.5254 0.8687 0.181 0.329 0.322 0.3076 0.508 0.638 0.3462 0.3882 0.7705 0.5933 0.994 0.1188 0.0782 0.94 0.00856 0.1396
0.2191 0.00648 0.8994 0.6714 0.6724 0.57 0.3127 0.4905 0.2119 0.3938 0.5957 0.1493 0.9424 0.716 0.3699 0.829 0.647 0.8286 0.04514
0.4028 0.5786 0.148 0.3425 0.999 0.869 0.04288 0.817 0.7075 0.03098 0.621 0.612 0.0774 0.532 0.4395 0.0711 0.4805 0.5835 0.5947
0.1768 0.52 0.3428 0.9146 0.7324 0.5054 0.7397 0.2737 0.6313 0.1704 0.5093 0.8105 0.1312 0.752 0.3647 0.781 0.4197 0.2329 0.787
0.762 0.63 0.9263 0.2673 0.1846 0.765 0.921 0.2913 0.3135 0.337 0.2598 0.1782 0.8013 0.641 0.6865 0.736 0.618 0.8755 0.2756
0.9854 0.8296 0.262 ]
输出数据(dst_gm):
[0.999, 1.38e-05], 1.38e-05使用reinterpret_cast方法转换后就是索引值232