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：实现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

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

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

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

参数说明

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

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推理产品

注意事项

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

返回值

无

调用示例

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

  #include "kernel_operator.h"
  namespace AscendC {
  class KernelLoadImage {
  public:
      __aicore__ inline KernelLoadImage()
      {
          // YUV420SP 图片中，Y 维度的 size
          gmSrc0Size = srcHorizSize * srcVertSize;
          // YUV420SP 图片中，UV 维度的 size
          gmSrc1Size = (srcHorizeSize / 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()
      {
          LocalTensor<int8_t> featureMapA1 = inQueueA1.AllocTensor<int8_t>();
          uint64_t fm_addr = static_cast<uint64_t>(reinterpret_cast<uintptr_t>(fmGlobal.GetPhyAddr()));
          // aipp config
          AippParams<int8_t> aippConfig;
          aippConfig.cPaddingParams.cPaddingMode = cPadMode;
          aippConfig.cPaddingParams.cPaddingValue = cPaddingValue;
          // fmGlobal为整张输入图片，src1参数处填入图片UV维度的起始地址
          SetAippFunctions(fmGlobal, fmGlobal[gmSrc0Size], inputFormat, aippConfig);
          LoadImageToLocal(featureMapA1, { horizSize, vertSize, horizStartPos, vertStartPos, srcHorizSize, topPadSize, botPadSize, leftPadSize, rightPadSize });
          inQueueA1.EnQue(featureMapA1);
      }
      __aicore__ inline void CopyToUB()
      {
          LocalTensor<int8_t> featureMapA1 = inQueueA1.DeQue<int8_t>();
          LocalTensor<int8_t> featureMapUB = outQueueUB.AllocTensor<int8_t>();
          DataCopy(featureMapUB, featureMapA1, dstSize);
          event_t eventIdMTE1ToMTE3 = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE1_MTE3));
          SetFlag<HardEvent::MTE1_MTE3>(eventIdMTE1ToMTE3);
          WaitFlag<HardEvent::MTE1_MTE3>(eventIdMTE1ToMTE3);
          outQueueUB.EnQue<int8_t>(featureMapUB);
          inQueueA1.FreeTensor(featureMapA1);
      }
      __aicore__ inline void CopyOut()
      {
          LocalTensor<int8_t> featureMapUB = outQueueUB.DeQue<int8_t>();
          DataCopy(dstGlobal, featureMapUB, dstSize);
          outQueueUB.FreeTensor(featureMapUB);
      }
  
  private:
      TPipe pipe;
      TQue<QuePosition::A1, 1> inQueueA1;
      TQue<QuePosition::VECOUT, 1> outQueueUB;
  
      GlobalTensor<uint8_t> fmGlobal;
      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, dstSize = 0;
      AippInputFormat inputFormat = AippInputFormat::YUV420SP_U8;
      uint32_t cPadMode = 0;
      int8_t cPaddingValue = 0;
  };
  } // namespace AscendC
  extern "C" __global__ __aicore__ void load_image_simple_kernel(__gm__ uint8_t *fmGm, __gm__ uint8_t *dstGm) 
  {
      AscendC::KernelLoadImage op;
      op.Init(fmGm, dstGm);
      op.Process();
  }