SetAippFunctions

Product Support

Function

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 (original data in the middle green region and padding data in other regions)
    
  • 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 (original data in the middle green region and padding data in other regions)
    
  • 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 (original data in the middle green region and padding data in other regions)
    
  • 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 (original data in the middle green region and padding data in other regions)
    

• 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 transfer API is called, the configured destination image height parameter becomes invalid after single-row read is enabled, for example, loadImageToLocalParams.vertSize of the LoadImageToLocal API.

• Data type conversion: Converts the pixel data type from uint8_t to int8_t or half. Conversion from uint8_t to int8_t limits output values to the range [–128, 127].

  1 2 3 4

  // Example 1: Implement uint8_t-to-int8_t type conversion and mean subtraction. Set the mean value of each channel to the average value of all data within the channel. (The min and var values are invalid and require no configuration.) output[i][j][k] = input[i][j][k] - mean[k] // Example 2: Implement uint8_t-to-fp16 type conversion and normalization. Set the mean value of each channel to the average value of all data within the channel, the min value to the minimum value after mean subtraction of all data within the channel, and the var value to the reciprocal of the difference between the maximum and minimum values of all data within 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 configured by the template parameter U. When U is uint8_t, the data type conversion function is disabled.

  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. For uint8_t or int8_t outputs, the channels are padded to a multiple of 32; for fp16 outputs, the channels are padded to a multiple of 16.

  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 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 input format, the chip discards the data of the fourth channel by default and outputs 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 into RGBX. If X is in channel 3, this function must be disabled to output 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(); }