SetAippFunctions

Supported Products

Function Usage

Sets the Artificial Intelligence Preprocessing (AIPP) parameters for images. This API is used together with the LoadImageToLocal API. After setting the parameters, you can call LoadImageToLocal to preprocess images during data movement, including data padding, channel swapping, single-line read, data type conversion, channel padding, and CSC. When SetAippFunctions is called, the matrix and format of the source image in the global memory need to be passed.

• Data padding: performs padding in the HW direction of the image. The following modes are available:
  • Mode 0: constant padding mode. Each position in the padding area is padded with a constant. The constant padded in each channel can be set. In this mode, only left and right padding is supported.
    Figure 1 Constant padding mode (the green area in the middle of the image indicates the original data, and the other areas indicate the padding data)
    
  • Mode 1: row/column padding mode. Each position in the padding area is padded with the data closest to the source image position in the row or column.
    Figure 2 Row/column padding mode (the green area in the middle of the image indicates the original data, and the other areas indicate the padding data)
    
  • Mode 2: block padding mode. Data blocks are copied from the source image for padding based on the padding width and height.
    Figure 3 Block padding mode (the green area in the middle of the image indicates the original data, and the other areas indicate the padding data)
    
  • Mode 3: mirrored block padding mode. The mirrored data blocks are copied from the source image for padding based on the padding width and height.
    Figure 4 Mirrored block padding mode (the green area in the middle of the image indicates the original data, and the other areas indicate the padding data)
    

• Channel swapping: Swaps image channels.

  For the RGB888 format, the R and B channels can be swapped.

  For the YUV420SP format, the U and V channels can be swapped.

  For the XRGB8888 format, the X channel can be moved backward (XRGB -> RGBX), and the R and B channels can be swapped.

• Single-line read: reads only one line from the source image.
  

  When the data movement API is called, the height parameter of the destination image is invalid after the single-line read function is enabled, for example, LoadImageToLocalParams.verSize of LoadImageToLocal.

• Data type conversion: The pixel data type can be converted from uint8_t to int8_t or half. When uint8_t is converted to int8_t, the output data range is limited to [–128, 127].

  1 2 3 4

  // Example 1: Convert the data type from uint8_t to int8_t and normalize the data. Set the mean value of each channel to the average value of all data in the channel. (The min and var values are invalid and do not need to be set.) output[i][j][k] = input[i][j][k] - mean[k] // Example 2: Convert the type from uint8 to fp16 and implement normalization. Set the mean value of each channel to the mean value of all data in the channel. The min value is set to the minimum value after zero mean of all data in the channel, while the var value is set to the reciprocal of the difference between the maximum value and the minimum value of all data in the channel. uint8_t -> fp16: output[i][j][k] = (input[i][j][k] - mean[k] - min[k]) * var[k]

  

  The data type after conversion is specified by the template parameter U. If U is of the uint8 type, the data type conversion function does not take effect.

  When the data movement API is called, the data type of the destination tensor must be the same as that of the output of this API. For example, the data type of dstLocal in LoadImageToLocal.

• Channel padding: performs padding on the image channel. The default value is 0.

  Mode 0: Pad the channel to the 32-byte channel. That is, if the output data type is uint8_t or int8_t, the channels are padded to 32 channels. If the output data type is fp16, the channels are padded to 16 channels.

  Mode 1: Pad channels to the 4-byte channel.

• CSC: converts the RGB format to the YUV format, or vice versa.

  

  

Prototype

• The input image format is YUV400, RGB888, or XRGB8888.

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

• The input image format is YUV420 semi-planar.

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

Parameters

enum class AippInputFormat : uint8_t {
    YUV420SP_U8 = 0,
    XRGB8888_U8 = 1,
    RGB888_U8 = 4,
    YUV400_U8 = 9,
};

1
2
3
4
5
6
7
8
9

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

Restrictions

• The address alignment requirements of src0 and src1 in the global memory are as follows.

  Image Format	src0	src1
  YUV420SP	Must be 2-byte aligned.	Must be 2-byte aligned.
  XRGB8888	Must be 4-byte aligned.	-
  RGB888	No alignment requirement.	-
  YUV400	No alignment requirement.	-

• For data in XRGB format, the chip discards the data of the fourth channel by default and outputs the data in RGB format. Therefore, if X is in channel 0, the function of moving the X channel backward must be enabled to convert the input channel to RGBX. Otherwise, if X is in channel 3, the function of moving the X channel backward must be disabled to output the data in RGB format.

Returns

None

Example

• This calling example supports the Atlas inference product 's AI Core platform. The example image format is 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() { // Size of the Y dimension of the image in YUV420SP format. gmSrc0Size = srcHorizSize * srcVertSize; // Size of the UV dimension of the image in YUV420SP format. 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 is the entire input image. Set src1 to the start address of the UV dimension of the image. 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::TPosition::A1, 1> inQueueA1; AscendC::TQue<AscendC::TPosition::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(); }