LoadImageToLocal
功能说明
将图像数据从GM搬运到A1/B1。 搬运过程中可以完成图像预处理操作:包括图像翻转,改变图像尺寸(抠图,裁边,缩放,伸展),以及色域转换,类型转换等。图像预处理的相关参数通过SetAippFunctions进行配置。
函数原型
1 2 |
template <typename T> __aicore__ inline void LoadImageToLocal(const LocalTensor<T>& dstLocal, const LoadImageToLocalParams& loadDataParams) |
参数说明
参数名称 |
输入/输出 |
含义 |
---|---|---|
dstLocal |
输出 |
目的操作数,类型为LocalTensor,支持的TPosition为A1/B1。 LocalTensor的起始地址需要保证32字节对齐。 Atlas推理系列产品AI Core,支持的数据类型为:uint8_t/int8_t/half Atlas A2训练系列产品/Atlas 800I A2推理产品,支持数据类型:int8_t/half Atlas 200/500 A2推理产品 支持数据类型为: uint8_t/int8_t/half |
loadDataParams |
输入 |
LoadData参数结构体,类型为LoadImageToLocalParams,具体参考表2。 |
参数名称 |
输入/输出 |
含义 |
---|---|---|
horizSize |
输入 |
从源图中加载图片的水平宽度,单位为像素,取值范围:horSize∈[2, 4095] 。 |
vertSize |
输入 |
从源图中加载图片的垂直高度,单位为像素,取值范围:verSize∈[2, 4095]。 |
horizStartPos |
输入 |
加载图片在源图片上的水平起始地址,单位为像素,取值范围:horizStartPos∈[0, 4095] 。默认为0。 注意:当输入图片为YUV420SP、XRGB8888, RGB888和YUV400格式时,该参数需要是偶数。 |
vertStartPos |
输入 |
加载图片在源图片上的垂直起始地址,单位为像素,取值范围:vertStartPos∈[0, 4095] 。默认为0。 注意:当输入图片为YUV420SP格式时,该参数需要是偶数。 |
srcHorizSize |
输入 |
源图像水平宽度 ,单位为像素,取值范围:srcHorizSize∈[2, 4095] 。 注意:当输入图片为YUV420SP格式时,该参数需要是偶数。 |
topPadSize |
输入 |
目的图像顶部填充的像素数 ,取值范围:topPadSize∈[0, 32] ,默认为0。进行数据填充时使用,需要先调用SetAippFunctions通过AippPaddingParams配置填充的数值,再通过topPadSize、botPadSize、leftPadSize、rightPadSize配置填充的大小范围。 |
botPadSize |
输入 |
目的图像底部填充的像素数,取值范围:botPadSize∈[0, 32] ,默认为0。 |
leftPadSize |
输入 |
目的图像左边填充的像素数,取值范围:lPadSize∈[0, 32] ,默认为0。 |
rightPadSize |
输入 |
目的图像右边填充的像素数,取值范围:rPadSize∈[0, 32] ,默认为0。 |
支持的型号
Atlas推理系列产品AI Core
Atlas A2训练系列产品/Atlas 800I A2推理产品
Atlas 200/500 A2推理产品
注意事项
- 操作数地址偏移对齐要求请参见通用约束。
- 加载到dstLocal的图片的大小加padding的大小必须小于等于L1的大小。
返回值
无
调用示例
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 |
#include "kernel_operator.h" class KernelLoadImage { public: __aicore__ inline KernelLoadImage() { // YUV420SP 图片中,Y 维度的 size gmSrc0Size = srcHorizSize * srcVertSize; // YUV420SP 图片中,UV 维度的 size gmSrc1Size = (srcHorizSize / 2) * (srcVertSize / 2) * 2; dstSize = dstHorizSize * dstVertSize * cSize; } __aicore__ inline void Init(__gm__ uint8_t *fmGm, __gm__ uint8_t *dstGm) { fmGlobal.SetGlobalBuffer((__gm__ uint8_t *)fmGm); dstGlobal.SetGlobalBuffer((__gm__ int8_t *)dstGm); pipe.InitBuffer(inQueueA1, 1, (gmSrc0Size + gmSrc1Size) * sizeof(int8_t)); pipe.InitBuffer(outQueueUB, 1, dstSize * sizeof(int8_t)); } __aicore__ inline void Process() { CopyIn(); CopyToUB(); CopyOut(); } private: __aicore__ inline void CopyIn() { AscendC::LocalTensor<int8_t> featureMapA1 = inQueueA1.AllocTensor<int8_t>(); uint64_t fm_addr = static_cast<uint64_t>(reinterpret_cast<uintptr_t>(fmGlobal.GetPhyAddr())); // aipp config AscendC::AippParams<int8_t> aippConfig; aippConfig.cPaddingParams.cPaddingMode = cPadMode; aippConfig.cPaddingParams.cPaddingValue = cPaddingValue; // fmGlobal为整张输入图片,src1参数处填入图片UV维度的起始地址 AscendC::SetAippFunctions(fmGlobal, fmGlobal[gmSrc0Size], inputFormat, aippConfig); AscendC::LoadImageToLocal(featureMapA1, { horizSize, vertSize, horizStartPos, vertStartPos, srcHorizSize, topPadSize, botPadSize, leftPadSize, rightPadSize }); inQueueA1.EnQue(featureMapA1); } __aicore__ inline void CopyToUB() { AscendC::LocalTensor<int8_t> featureMapA1 = inQueueA1.DeQue<int8_t>(); AscendC::LocalTensor<int8_t> featureMapUB = outQueueUB.AllocTensor<int8_t>(); AscendC::DataCopy(featureMapUB, featureMapA1, dstSize); event_t eventIdMTE1ToMTE3 = static_cast<event_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::MTE1_MTE3)); AscendC::SetFlag<AscendC::HardEvent::MTE1_MTE3>(eventIdMTE1ToMTE3); AscendC::WaitFlag<AscendC::HardEvent::MTE1_MTE3>(eventIdMTE1ToMTE3); outQueueUB.EnQue<int8_t>(featureMapUB); inQueueA1.FreeTensor(featureMapA1); } __aicore__ inline void CopyOut() { AscendC::LocalTensor<int8_t> featureMapUB = outQueueUB.DeQue<int8_t>(); AscendC::DataCopy(dstGlobal, featureMapUB, dstSize); outQueueUB.FreeTensor(featureMapUB); } private: AscendC::TPipe pipe; AscendC::TQue<AscendC::QuePosition::A1, 1> inQueueA1; AscendC::TQue<AscendC::QuePosition::VECOUT, 1> outQueueUB; AscendC::GlobalTensor<uint8_t> fmGlobal; AscendC::GlobalTensor<int8_t> dstGlobal; uint16_t horizSize = 32, vertSize = 32, horizStartPos = 0, vertStartPos = 0, srcHorizSize = 32, srcVertSize = 32, leftPadSize = 0, rightPadSize = 0; uint32_t dstHorizSize = 32, dstVertSize = 32, cSize = 32; uint8_t topPadSize = 0, botPadSize = 0; uint32_t gmSrc0Size = 0, gmSrc1Size = 0, dstSize = 0; AscendC::AippInputFormat inputFormat = AscendC::AippInputFormat::YUV420SP_U8; uint32_t cPadMode = 0; int8_t cPaddingValue = 0; }; extern "C" __global__ __aicore__ void load_image_simple_kernel(__gm__ uint8_t *fmGm, __gm__ uint8_t *dstGm) { KernelLoadImage op; op.Init(fmGm, dstGm); op.Process(); } |