SetAippFunctions

功能说明

设置图片预处理（AIPP，AI core pre-process）相关参数。和LoadImageToLocal接口配合使用。设置后，调用LoadImageToLocal接口可在搬运过程中完成图像预处理操作：包括数据填充，通道交换，单行读取、数据类型转换、通道填充、色域转换。调用SetAippFunctions接口时需传入源图片在Global Memory上的矩阵、源图片的图片格式。

• 数据填充：在图片HW方向上padding。分为如下几种模式：
  • 模式0：常量填充模式，padding区域各位置填充为常数，支持设置每个通道填充的常数。该模式下仅支持左右padding，不支持上下padding。
    图1 常量填充模式
    
  • 模式1：行列填充模式，padding区域各位置填充行/列上最邻近源图片位置的数据。
    图2 行列填充模式
    
  • 模式2：块填充模式，按照padding的宽高，从源图片拷贝数据块进行padding区域填充。
    图3 块填充模式
    
  • 模式3：镜像块填充模式，按照padding的宽高，从源图片拷贝数据块的镜像进行padding区域填充。
    图4 镜像块填充模式
    

• 通道交换：将图片通道进行交换。

  对于RGB888格式，支持交换R和B通道。

  对于YUV420SP格式，支持交换U和V通道。

  对于XRGB8888格式，支持X通道后移（XRGB->RGBX）、支持交换R和B通道。

  例：对于XRGB格式的数据，芯片在处理的时候会默认丢弃掉第四个通道的数据留下XRG。因此如果要保留RGB通道的数据，对于XRGB输入的需要后移X通道，将输入的通道转换为RGBX即可。

• 单行读取：源图片中仅读取一行。
  

  调用数据搬运接口时，开启单行读取后设置的目的图片高度参数无效，如LoadImageToLocal接口的loadImageToLocalParams.verSize。

• 数据类型转换：转换像素的数据类型，支持uint8转换为int8或fp16。

  1 2 3 4

  // 例1：实现uint8 ->int8 的类型转换，同时实现零均值化：设置每个通道mean值为该通道所有数据的平均值（min和var值无效，不用设置）。 output[i][j][k] = input[i][j][k] – mean[k] // 例2：实现uint8 -> fp16 的类型转换，同时实现归一化：设置每个通道mean值为该通道所有数据的平均值，min值为该通道所有数据零均值化后的最小值，var值为该通道所有数据的最大值减最小值的倒数。 uint8 -> fp16: output[i][j][k] = (input[i][j][k] – mean[k] – min[k]) * var[k]

  

  转换后的数据类型是由模板参数U配置，U为uint8时数据类型转换功能不生效。

  调用数据搬运接口时，目的Tensor的数据类型需要与本接口输出数据类型保持一致，如LoadImageToLocal的dstLocal参数的数据类型。

• 通道填充：在图片通道方向上padding。默认为模式0。

  模式0：将通道padding至32Bytes。即输出数据类型为uint8/int8时，padding至32通道；输出数据类型为fp16时，padding至16通道。

  模式1：将通道padding至4通道。

• 色域转换：RGB格式转换为YUV格式，或YUV模式转换为RGB格式。

  

  

函数原型

• 输入图片格式为YUV400、RGB888、XRGB8888

  1 2	template<typename T, typename U> void SetAippFunctions(const GlobalTensor<T>& src0, AippInputFormat format, AippParams<U> config)

• 输入图片格式为YUV420 Semi-Planar

  1 2	template<typename T, typename U> void SetAippFunctions(const GlobalTensor<T>& src0, const GlobalTensor<T>& src1, AippInputFormat format, AippParams<U> config)

参数说明

1
2
3
4

AippInputFormat::YUV420SP_U8  // YUV420 Semi-Planar，数据类型为uint8_t
AippInputFormat::XRGB8888_U8  // XRGB8888，数据类型为uint8_t
AippInputFormat::RGB888_U8  // RGB888，数据类型为uint8_t
AippInputFormat::YUV400_U8  // YUV400，数据类型为uint8_t

1
2
3
4
5
6
7
8
9

template <typename U>
struct AippParams {
    AippPaddingParams<U> paddingParams;
    AippSwapParams swapParams;
    AippSingleLineParams singleLineParams;
    AippDataTypeConvParams dtcParams;
    AippChannelPaddingParams<U> cPaddingParams;
    AippColorSpaceConvParams cscParams;
};

支持的型号

Atlas推理系列产品AI Core

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

Atlas 200/500 A2推理产品

注意事项

src0、src1在Global Memory上的地址对齐要求如下：

返回值

无

调用示例

• 该调用示例支持的运行平台为Atlas推理系列产品AI Core，示例图片格式为YUV420SP。

  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

  #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(); }