Device侧代码与说明

Device侧代码示例如下，并参考本节内容了解主要操作。

// Device侧文件：kernel_aiv.cce
#ifdef __CCE_KT_TEST__
#define __aicore__ 
#else
#define __aicore__ [aicore]
#endif

extern "C" __global__ __aicore__ void add_kernel(__gm__ float* __restrict__ input_1, __gm__ float* __restrict__ input_2, __gm__ float* __restrict__ output) {
    set_vector_mask(-1, -1);
    set_mask_norm();
    set_atomic_none();
    __ubuf__ float* input_1_local_UB = (__ubuf__ float*)get_imm(0);
    __ubuf__ float* input_2_local_UB = (__ubuf__ float*)get_imm(1024);
    __ubuf__ float* output_local_UB = (__ubuf__ float*)get_imm(2048);
    copy_gm_to_ubuf(input_1_local_UB, input_1, 0, 1, 8, 0, 0);
    copy_gm_to_ubuf(input_2_local_UB, input_2, 0, 1, 8, 0, 0);
    set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
    wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
    vadd(output_local_UB, input_1_local_UB, input_2_local_UB, 1, 1, 1, 1, 0, 0, 0);
    set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);
    wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);
    copy_ubuf_to_gm(output, output_local_UB, 0, 1, 8, 0, 0);
    pipe_barrier(PIPE_ALL);
}

函数内首先通过get_imm()接口传入立即数偏移量，初始化三个需要在UB用到tensor首地址。

1. 数据搬入。
   将需要进行加法计算的两个加数tensor，从GM搬到UB上，这里使用了copy_gm_to_ubuf指令，其接口原型为：

   void copy_gm_to_ubuf(__ubuf__ void *dst, __gm__ void *src, uint8_t sid, uint16_t nBurst, uint16_t lenBurst, uint16_t srcGap, uint16_t dstGap)

   • dst：UB上首地址空间首地址。
   • src：GM起始地址。
   • sid：预留参数，此处默认为0即可。
   • nBurst：示例从GM搬运256B（64个float类型），一次搬完，所以nBurst设置为1。
   • lenBurst：一次连续搬运数据长度，单位32Bytes，元素个数为64，64 * sizeof(float) = 256B，计算搬运长度256/32=8。
   • srcStride：两次搬运源地址间隙，单位32Bytes，示例中只搬运一块数据，所以不涉及stride，此处设置0。
   • dstStride：两次搬运目的地址间隙，单位32Bytes，示例中只搬运一块数据，所以不涉及stride，此处设置0。

   由于下一步计算依赖搬运数据的完整性，所以需要插入同步对，前置指令的流水类型改为PIPE_MTE2，等待其完成过后，才能开始PIPE_V的计算。

2. 计算。
   将在ub上的两个加数进行逐元素相加，这里使用vadd指令，其接口原型为：

   void vadd(__ubuf__ float *dst, __ubuf__ float *src0, __ubuf__ float *src1, uint8_t repeat, uint8_t dstBlockStride, uint8_t src0BlockStride, uint8_t src1BlockStride, uint8_t dstRepeatStride, uint8_t src0RepeatStride, uint8_t src1RepeatStride)

   • dst：计算结果存放UB上的首地址。
   • src0：第一个加数在UB上的首地址。
   • src1：第二个加数在UB上的首地址。
   • repeat：指令迭代次数，本次是256B，共计8个block，1次执行完。
   • dstBlockStride：同一次执行，目的操作数不同block间地址步长。单位为32B，示例dst block连续存储，所以dstBlockStride为1，每个dst block的起始地址间隔为1个block(32B)。
   • src0BlockStride：同一次执行，源操作数0不同block间地址步长。单位为32B，示例src0 block连续读取参与运算，所以dstBlockStride为1，每个dst block的起始地址间隔为1个block(32B)。
   • src1BlockStride：同一次执行，源操作数1不同block间地址步长。单位为32B，示例src1 block连续读取参与运算，所以dstBlockStride为1，每个dst block的起始地址间隔为1个block(32B)。
   • dstRepeatStride：相邻两次执行，目的操作数相同block地址步长。单位为32B，此处不涉及，所以设为0。
   • src0RepeatStride：相邻两次执行，源操作数0相同block地址步长。单位为32B，此处不涉及，所以设为0。
   • src1RepeatStride：相邻两次执行，源操作数1相同block地址步长。单位为32B，此处不涉及，所以设为0
3. 数据搬出。
   将计算完的结果tensor，从UB上搬出到结果GM指定位置，这里使用了copy_ubuf_to_gm指令，接口原型为：

   void copy_ubuf_to_gm(__gm__ void *dst, __ubuf__ void *src, uint8_t sid, uint16_t nBurst, uint16_t lenBurst, uint16_t srcGap, uint16_t dstGap)

   • dst：GM上输出结果的首地址。
   • src：计算结果在UB上的首地址。
   • sid：预留参数，此处默认为0即可。
   • nBurst：示例从UB搬运256B（128个half类型），一次搬完，所以nBurst设置为1。
   • lenBurst：一次连续搬运数据长度，单位32Bytes；元素个数为64，64 * sizeof(float) = 256B，计算搬运长度256/32=8。
   • srcStride：两次搬运源地址间隙，单位32Bytes；示例中只搬运一块数据，所以不涉及stride，此处设置0。
   • dstStride：两次搬运目的地址间隙，单位32Bytes；示例中只搬运一块数据，所以不涉及stride，此处设置0。