Sort
函数功能
排序函数，按照数值大小进行降序排序。一次迭代可以完成32个数的排序，排序后的数据按照如下排布方式进行保存：
Atlas A2训练系列产品/Atlas 800I A2推理产品采用方式一
Atlas推理系列产品AI Core采用方式二
- 排布方式一：
- 排布方式二：Region Proposal排布
输入输出数据均为Region Proposal。每个Region Proposal占用连续8个half/float类型的元素，约定其格式：
[x1, y1, x2, y2, score, label, reserved_0, reserved_1]
对于数据类型half，每一个Region Proposal占16Bytes，Byte[15:12]是无效数据，Byte[11:0]包含6个half类型的元素，其中Byte[11:10]定义为label，Byte[9:8]定义为score，Byte[7:6]定义为y2，Byte[5:4]定义为x2，Byte[3:2]定义为y1，Byte[1:0]定义为x1。
如下图所示，总共包含16个Region Proposals。
对于数据类型float，每一个Region Proposal占32Bytes，Byte[31:24]是无效数据，Byte[23:0]包含6个float类型的元素，其中Byte[23:20]定义为label，Byte[19:16]定义为score，Byte[15:12]定义为y2，Byte[11:8]定义为x2，Byte[7:4]定义为y1，Byte[3:0]定义为x1。
如下图所示，总共包含16个Region Proposals。
函数原型
1 2 |
template <typename T, bool isFullSort> __aicore__ inline void Sort(const LocalTensor<T> &dstLocal, const LocalTensor<T> &concatLocal, const LocalTensor<uint32_t> &indexLocal, LocalTensor<T> &tmpLocal, const int32_t repeatTimes) |
参数说明
接口 |
功能 |
|---|---|
T |
操作数的数据类型。 |
isFullSort |
是否开启全排序模式。全排序模式指将全部输入降序排序，非全排序模式下，排序方式请参考表2中的repeatTimes说明。 |
参数名称 |
输入/输出 |
含义 |
|---|---|---|
dstLocal |
输出 |
目的操作数。 类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。 Atlas A2训练系列产品/Atlas 800I A2推理产品，支持的数据类型为：half/float Atlas推理系列产品AI Core，支持的数据类型为：half/float |
concatLocal |
输入 |
源操作数。 类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。 此源操作数的数据类型需要与目的操作数保持一致。 Atlas A2训练系列产品/Atlas 800I A2推理产品，支持的数据类型为：half/float Atlas推理系列产品AI Core，支持的数据类型为：half/float |
indexLocal |
输入 |
源操作数。 类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。 此源操作数固定为uint32_t数据类型。 |
tmpLocal |
输入 |
临时空间。接口内部复杂计算时用于存储中间变量，由开发者提供。数据类型与源操作数保持一致。 类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。 Atlas A2训练系列产品/Atlas 800I A2推理产品，支持的数据类型为：half/float Atlas推理系列产品AI Core，支持的数据类型为：half/float |
repeatTimes |
输入 |
重复迭代次数，int32_t类型。
|
返回值
无
支持的型号
Atlas A2训练系列产品/Atlas 800I A2推理产品
Atlas推理系列产品AI Core
约束说明
- 当存在score[i]与score[j]相同时，如果i>j，则score[j]将首先被选出来，排在前面，即index的顺序与输入顺序一致。
- 非全排序模式下，每次迭代内的数据会进行排序，不同迭代间的数据不会进行排序。
- 操作数地址偏移对齐要求请参见通用约束。
调用示例
- 处理128个half类型数据。
Atlas A2训练系列产品/Atlas 800I A2推理产品
#include "kernel_operator.h" template <typename T> class FullSort { public: __aicore__ inline FullSort() {} __aicore__ inline void Init(__gm__ uint8_t *srcValueGm, __gm__ uint8_t *srcIndexGm, __gm__ uint8_t *dstValueGm, __gm__ uint8_t *dstIndexGm) { concatRepeatTimes = elementCount / 16; inBufferSize = elementCount * sizeof(uint32_t); outBufferSize = elementCount * sizeof(uint32_t); calcBufferSize = elementCount * 8; tmpBufferSize = elementCount * 8; sortedLocalSize = elementCount * 4; sortRepeatTimes = elementCount / 32; extractRepeatTimes = elementCount / 32; sortTmpLocalSize = elementCount * 4; valueGlobal.SetGlobalBuffer((__gm__ T *)srcValueGm); indexGlobal.SetGlobalBuffer((__gm__ uint32_t *)srcIndexGm); m_dstValueGlobal.SetGlobalBuffer((__gm__ T *)dstValueGm); dstIndexGlobal.SetGlobalBuffer((__gm__ uint32_t *)dstIndexGm); m_pipe.InitBuffer(queIn, 2, inBufferSize); m_pipe.InitBuffer(queOut, 2, outBufferSize); m_pipe.InitBuffer(queCalc, 1, calcBufferSize * sizeof(T)); m_pipe.InitBuffer(queTmp, 2, tmpBufferSize * sizeof(T)); } __aicore__ inline void Process() { CopyIn(); Compute(); CopyOut(); } private: __aicore__ inline void CopyIn() { AscendC::LocalTensor<T> valueLocal = queIn.AllocTensor<T>(); AscendC::DataCopy(valueLocal, valueGlobal, elementCount); queIn.EnQue(valueLocal); AscendC::LocalTensor<uint32_t> indexLocal = queIn.AllocTensor<uint32_t>(); AscendC::DataCopy(indexLocal, indexGlobal, elementCount); queIn.EnQue(indexLocal); } __aicore__ inline void Compute() { AscendC::LocalTensor<T> valueLocal = queIn.DeQue<T>(); AscendC::LocalTensor<uint32_t> indexLocal = queIn.DeQue<uint32_t>(); AscendC::LocalTensor<T> sortedLocal = queCalc.AllocTensor<T>(); AscendC::LocalTensor<T> concatTmpLocal = queTmp.AllocTensor<T>(); AscendC::LocalTensor<T> sortTmpLocal = queTmp.AllocTensor<T>(); AscendC::LocalTensor<T> dstValueLocal = queOut.AllocTensor<T>(); AscendC::LocalTensor<uint32_t> dstIndexLocal = queOut.AllocTensor<uint32_t>(); AscendC::LocalTensor<T> concatLocal; AscendC::Concat(concatLocal, valueLocal, concatTmpLocal, concatRepeatTimes); AscendC::Sort<T, true>(sortedLocal, concatLocal, indexLocal, sortTmpLocal, sortRepeatTimes); AscendC::Extract(dstValueLocal, dstIndexLocal, sortedLocal, extractRepeatTimes); queTmp.FreeTensor(concatTmpLocal); queTmp.FreeTensor(sortTmpLocal); queIn.FreeTensor(valueLocal); queIn.FreeTensor(indexLocal); queCalc.FreeTensor(sortedLocal); queOut.EnQue(dstValueLocal); queOut.EnQue(dstIndexLocal); } __aicore__ inline void CopyOut() { AscendC::LocalTensor<T> dstValueLocal = queOut.DeQue<T>(); AscendC::LocalTensor<uint32_t> dstIndexLocal = queOut.DeQue<uint32_t>(); AscendC::DataCopy(m_dstValueGlobal, dstValueLocal, elementCount); AscendC::DataCopy(dstIndexGlobal, dstIndexLocal, elementCount); queOut.FreeTensor(dstValueLocal); queOut.FreeTensor(dstIndexLocal); } private: AscendC::TPipe m_pipe; AscendC::TQue<AscendC::QuePosition::VECIN, 2> queIn; AscendC::TQue<AscendC::QuePosition::VECOUT, 2> queOut; AscendC::TQue<AscendC::QuePosition::VECIN, 2> queTmp; AscendC::TQue<AscendC::QuePosition::VECIN, 1> queCalc; AscendC::GlobalTensor<T> valueGlobal; AscendC::GlobalTensor<uint32_t> indexGlobal; AscendC::GlobalTensor<T> m_dstValueGlobal; AscendC::GlobalTensor<uint32_t> dstIndexGlobal; uint32_t elementCount = 128; uint32_t concatRepeatTimes; uint32_t inBufferSize; uint32_t outBufferSize; uint32_t calcBufferSize; uint32_t tmpBufferSize; uint32_t sortedLocalSize; uint32_t sortTmpLocalSize; uint32_t sortRepeatTimes; uint32_t extractRepeatTimes; }; extern "C" __global__ __aicore__ void sort_operator(__gm__ uint8_t *src0Gm, __gm__ uint8_t *src1Gm, __gm__ uint8_t *dst0Gm, __gm__ uint8_t *dst1Gm) { FullSort<half> op; op.Init(src0Gm, src1Gm, dst0Gm, dst1Gm); op.Process(); }示例结果 输入数据(srcValueGm): 128个float类型数据 [31 30 29 ... 2 1 0 63 62 61 ... 34 33 32 95 94 93 ... 66 65 64 127 126 125 ... 98 97 96] 输入数据(srcIndexGm): [31 30 29 ... 2 1 0 63 62 61 ... 34 33 32 95 94 93 ... 66 65 64 127 126 125 ... 98 97 96] 输出数据(dstValueGm): [127 126 125 ... 2 1 0] 输出数据(dstIndexGm): [127 126 125 ... 2 1 0]
- 处理64个half类型数据。
Atlas推理系列产品AI Core
#include "kernel_operator.h" template <typename T> class FullSort { public: __aicore__ inline FullSort() {} __aicore__ inline void Init(__gm__ uint8_t *srcValueGm, __gm__ uint8_t *srcIndexGm, __gm__ uint8_t *dstValueGm, __gm__ uint8_t *dstIndexGm) { concatRepeatTimes = elementCount / 16; inBufferSize = elementCount * sizeof(uint32_t); outBufferSize = elementCount * sizeof(uint32_t); calcBufferSize = elementCount * 8; tmpBufferSize = elementCount * 8; sortedLocalSize = elementCount * 8 * sizeof(T); sortRepeatTimes = elementCount / 16; extractRepeatTimes = elementCount / 16; sortTmpLocalSize = elementCount * 8 * sizeof(T); m_valueGlobal.SetGlobalBuffer((__gm__ T *)srcValueGm); indexGlobal.SetGlobalBuffer((__gm__ uint32_t *)srcIndexGm); m_dstValueGlobal.SetGlobalBuffer((__gm__ T *)dstValueGm); dstIndexGlobal.SetGlobalBuffer((__gm__ uint32_t *)dstIndexGm); m_pipe.InitBuffer(queIn, 2, inBufferSize); m_pipe.InitBuffer(queOut, 2, outBufferSize); m_pipe.InitBuffer(queCalc, 1, calcBufferSize * sizeof(T)); m_pipe.InitBuffer(queTmp, 2, tmpBufferSize * sizeof(T)); } __aicore__ inline void Process() { CopyIn(); Compute(); CopyOut(); } private: __aicore__ inline void CopyIn() { AscendC::LocalTensor<T> valueLocal = queIn.AllocTensor<T>(); AscendC::DataCopy(valueLocal, m_valueGlobal, elementCount); queIn.EnQue(valueLocal); AscendC::LocalTensor<uint32_t> indexLocal = queIn.AllocTensor<uint32_t>(); AscendC::DataCopy(indexLocal, indexGlobal, elementCount); queIn.EnQue(indexLocal); } __aicore__ inline void Compute() { AscendC::LocalTensor<T> valueLocal = queIn.DeQue<T>(); AscendC::LocalTensor<uint32_t> indexLocal = queIn.DeQue<uint32_t>(); AscendC::LocalTensor<T> sortedLocal = queCalc.AllocTensor<T>(); AscendC::LocalTensor<T> concatTmpLocal = queTmp.AllocTensor<T>(); AscendC::LocalTensor<T> sortTmpLocal = queTmp.AllocTensor<T>(); AscendC::LocalTensor<T> dstValueLocal = queOut.AllocTensor<T>(); AscendC::LocalTensor<uint32_t> dstIndexLocal = queOut.AllocTensor<uint32_t>(); AscendC::LocalTensor<T> concatLocal; AscendC::Concat(concatLocal, valueLocal, concatTmpLocal, concatRepeatTimes); AscendC::Sort<T, true>(sortedLocal, concatLocal, indexLocal, sortTmpLocal, sortRepeatTimes); AscendC::Extract(dstValueLocal, dstIndexLocal, sortedLocal, extractRepeatTimes); queTmp.FreeTensor(concatTmpLocal); queTmp.FreeTensor(sortTmpLocal); queIn.FreeTensor(valueLocal); queIn.FreeTensor(indexLocal); queCalc.FreeTensor(sortedLocal); queOut.EnQue(dstValueLocal); queOut.EnQue(dstIndexLocal); } __aicore__ inline void CopyOut() { AscendC::LocalTensor<T> dstValueLocal = queOut.DeQue<T>(); AscendC::LocalTensor<uint32_t> dstIndexLocal = queOut.DeQue<uint32_t>(); AscendC::DataCopy(m_dstValueGlobal, dstValueLocal, elementCount); AscendC::DataCopy(dstIndexGlobal, dstIndexLocal, elementCount); queOut.FreeTensor(dstValueLocal); queOut.FreeTensor(dstIndexLocal); } private: AscendC::TPipe m_pipe; AscendC::TQue<AscendC::QuePosition::VECIN, 2> queIn; AscendC::TQue<AscendC::QuePosition::VECOUT, 2> queOut; AscendC::TQue<AscendC::QuePosition::VECIN, 2> queTmp; AscendC::TQue<AscendC::QuePosition::VECIN, 1> queCalc; AscendC::GlobalTensor<T> m_valueGlobal; AscendC::GlobalTensor<uint32_t> indexGlobal; AscendC::GlobalTensor<T> m_dstValueGlobal; AscendC::GlobalTensor<uint32_t> dstIndexGlobal; uint32_t elementCount = 64; uint32_t concatRepeatTimes; uint32_t inBufferSize; uint32_t outBufferSize; uint32_t calcBufferSize; uint32_t tmpBufferSize; uint32_t sortedLocalSize; uint32_t sortTmpLocalSize; uint32_t sortRepeatTimes; uint32_t extractRepeatTimes; }; extern "C" __global__ __aicore__ void sort_operator(__gm__ uint8_t *src0Gm, __gm__ uint8_t *src1Gm, __gm__ uint8_t *dst0Gm, __gm__ uint8_t *dst1Gm) { FullSort<half> op; op.Init(src0Gm, src1Gm, dst0Gm, dst1Gm); op.Process(); }示例结果 输入数据(srcValueGm): 128个float类型数据 [15 14 13 ... 2 1 0 31 30 29 ... 18 17 16 47 46 45 ... 34 33 32 63 62 61 ... 50 49 48] 输入数据(srcIndexGm): [15 14 13 ... 2 1 0 31 30 29 ... 18 17 16 47 46 45 ... 34 33 32 63 62 61 ... 50 49 48] 输出数据(dstValueGm): [63 62 61 ... 2 1 0] 输出数据(dstIndexGm): [63 62 61 ... 2 1 0]
父主题： Sort
