【2023CANN训练营第二季】——Ascend C算子开发进阶

【2023CANN训练营第二季】——Ascend C算子开发进阶—Ascend C Tiling计算

了解Tiling基本概念

在这一小节中接触到了一个新的概念，叫Tiling计算，指的是在Ascend C 算子开发过程中，矢量的算子流程分为3个基本任务：CopyIn，Compute，CopyOut。CopyIn任务负责将Global Memory上的输入Tensor xGm和yGm搬运至Local Memory，但是Local Memory不能容纳所有算子的输入和输出，所以需要每次搬入一部分数据进行计算，然后再搬出，再搬入另一部分数据，重复上述过程，直到最终得到完整结果，这个把全部数据进行切分、分块的计算就叫做Tiling计算。

Tiling两种实现方式

有两种场景的Tiling实现，分别为固定shape场景与动态shape场景。

固定shape场景：输入大小都是固定的，实现难度低，只要考虑shape的逻辑处理，优化难度低。
动态shape场景：可以将形状通过核函数的入参传入核函数，满足shape变动的场景，实现难度高，要考虑不同逻辑分支处理，优化难度也高。

两种场景的核函数add_custom对比

固定shape核函数实现

#include "add_custom_unalign_tiling.h"
#include "register/op_def_registry.h"

namespace optiling {
constexpr uint32_t BLOCK_DIM = 8;
constexpr uint32_t SIZE_OF_HALF = 2;
constexpr uint32_t BLOCK_SIZE = 32;
// shape需要对齐到的最小单位
constexpr uint32_t ALIGN_NUM = BLOCK_SIZE / SIZE_OF_HALF;

这段代码的目的是定义一些常量并计算一个需要对齐到的最小单位的值。

动态shape核函数实现：

#include "kernel_operator.h"
using namespace AscendC;
constexpr int32_t BUFFER_NUM = 2;
extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling)
{
    GET_TILING_DATA(tilingData, tiling);
    KernelAdd op;
    op.Init(x, y, z, tilingData.totalLength, tilingData.tileNum);
    if (TILING_KEY_IS(1)) {
        op.Process();
    }
}

动态shape场景样例演示

固定shape 下Ascend C矢量加法实现代码在samples/cplusplus/level1_single_api/4_op_dev/6_ascendc_custom_op/kernel_invocation/Add/add_custom.cpp文件中，动态shape 对应的实现在samples/cplusplus/level1_single_api/4_op_dev/6_ascendc_custom_op/kernel_invocation/Add_tile/add_custom.cpp文件中

下面分别介绍两种场景下的样例演示
一、核函数
对于两种场景下，核函数的区别在于动态场景下会多了两个参数，workspace和tiling。而固定shape的核函数只有x,y,z。除此之外在核函数中，动态shape还多了GET_TILING_DATA函数以及op.Init函数中多了两个入参。