Device侧代码与说明

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

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

extern "C" __global__ __aicore__ void add_kernel(__gm__ uint8_t * __restrict tensor_a, __gm__ uint8_t * __restrict tensor_b, __gm__ uint8_t * __restrict tensor_c) {

    // Vector初始化配置
    set_vector_mask(-1, -1);
    set_mask_norm();
    set_atomic_none();

    // 指针类型转换，将通用的字节指针（uint8_t*）转换为特定类型的浮点指针（float*）
    __gm__ float *aGm = (__gm__ float *)tensor_a;
    __gm__ float *bGm = (__gm__ float *)tensor_b;
    __gm__ float *cGm = (__gm__ float *)tensor_c;
    
    // 初始化三个需要用到的tensor首地址，均在UB中
    __ubuf__ float *aUB = (__ubuf__ float *)get_imm(0);
    __ubuf__ float *bUB = (__ubuf__ float *)get_imm(256);
    __ubuf__ float *cUB = (__ubuf__ float *)get_imm(512);
    
    // 数据搬入
    copy_gm_to_ubuf(aUB, tensor_a, 0, 1, 8, 0, 0);
    copy_gm_to_ubuf(bUB, tensor_b, 0, 1, 8, 0, 0);

    // 同步
    set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
    wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);

    // 计算
    vadd(cUB, aUB, bUB, 1, 1, 1, 1, 0, 0, 0);
    
    // 同步
    set_flag(PIPE_V, PIPE_MTE3, EVENT_ID1);
    wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID1);

    // 数据搬出
    copy_ubuf_to_gm(cGm, cUB, 0, 1, 8, 0, 0, bm_t::BM_DISABLE);

    // 同步
    pipe_barrier(PIPE_ALL);
}

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：目标地址，aUB/bUB。
   • src：源地址，aGm/bGM。
   • sid：预留参数，此处默认为0即可。
   • nBurst：搬运数据块个数，此处不分块，所以nBurst设置为1。
   • lenBurst：一次连续搬运数据长度，单位32Bytes，元素个数为64，64*sizeof(float) = 256B，计算搬运长度256/32=8。
   • srcGap：两次搬运源地址间隙，源地址处连续搬运，此处设置0。
   • dstGap：两次搬运目的地址间隙，目标地址处连续搬运，此处设置0。
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：目标地址，cUB。
   • src0：源地址0，aUB。
   • src1：源地址1，bUB。
   • repeat：指令迭代次数，Vector 计算单元每个迭代会从UB中取出8个datablock（每个datablock 数据块内部地址连续，长度 32Byte，合计 256Byte）进行计算，本次是256B，共计8个block，1次执行完。
   • dstBlockStride：目的操作数不同block间地址步长。单位为32B，因为连续存储，所以dstBlockStride为1。
   • src0BlockStride：源操作数0不同block间地址步长。单位为32B，因为src0 block连续读取参与运算，所以src0BlockStride为1。
   • src1BlockStride：源操作数1不同block间地址步长。单位为32B，因为src1 block连续读取参与运算，所以src1BlockStride为1。
   • dstRepeatStride：相邻两次执行间，目的操作数相同block间地址步长。单位为32B，因为repeat设为1，所以此处不涉及，设为0。
   • src0RepeatStride：相邻两次执行间，源操作数0相同block间地址步长。单位为32B，因为repeat设为1，所以此处不涉及，设为0。
   • src1RepeatStride：相邻两次执行间，源操作数1相同block间地址步长。单位为32B，因为repeat设为1，所以此处不涉及，设为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, bm_t byteMode);

   • dst：目标地址，cGm。
   • src：源地址，cUB。
   • sid：预留参数，此处默认为0即可。
   • nBurst：搬运数据块个数，此处不分块，所以nBurst设置为1。
   • lenBurst：一次连续搬运数据长度，单位32Bytes，元素个数为64，64 * sizeof(float) = 256B，计算搬运长度256/32=8。
   • srcGap：两次搬运源地址间隙，源地址处连续搬运，此处设置0。
   • dstGap：两次搬运目的地址间隙，目标地址处连续搬运，此处设置0。
   • byteMode：通过设置byteMode枚举参数可改变block大小。设为bm_t::BM_DISABLE，即block大小为32Bytes。