Device侧代码与说明

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

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

extern "C" __global__ __aicore__ void mat_mul_relu_mix_kernel(__gm__ uint16_t* __restrict__ ffts_addr, __gm__ uint8_t *tensor_a, __gm__ uint8_t *tensor_b, __gm__ uint8_t *tensor_c, int m, int k, int n) {

    // Vector与Cube核间同步，用于初始化FFTS_BASE_ADDR寄存器
    set_ffts_base_addr((uint64_t)ffts_addr);

    // 指针类型转换，将通用的字节指针（uint8_t*）转换为特定类型的浮点指针（float*）
    __gm__ float *aGm = (__gm__ float *)tensor_a;
    __gm__ float *bGm = (__gm__ float *)tensor_b;
    __gm__ float *cGm = (__gm__ float *)tensor_c;

    // Cube核
    #if defined(__DAV_C220_CUBE__)

    // 初始化三个需要用到的tensor首地址，分别在L0A、L0B和L0C中
    __ca__ float *aL0a = (__ca__ float *)get_imm(0);
    __cb__ float *bL0b = (__cb__ float *)get_imm(0);
    __cc__ float *cL0c = (__cc__ float *)get_imm(0);
    
    // 数据搬入
    load_gm_to_ca(aL0a, aGm, 0, 8, 1, 0, 0, inc); 
    load_gm_to_cb(bL0b, bGm, 0, 8, 1, 0, 0, inc); 
    
    // 同步
    set_flag(PIPE_MTE2, PIPE_M, EVENT_ID0);
    wait_flag(PIPE_MTE2, PIPE_M, EVENT_ID0);

    // 计算
    mad(cL0c, aL0a, bL0b, m, k, n, 0, 0, 0, 0, 0, 0, true);

    // 同步
    set_flag(PIPE_M, PIPE_FIX, EVENT_ID1);
    wait_flag(PIPE_M, PIPE_FIX, EVENT_ID1);

    // 数据搬出
    set_nd_para(4398046576641);
    copy_matrix_cc_to_gm(cGm, cL0c, 0, n, m, n, m, 0, 0, 0, false, true);

    // 同步
    pipe_barrier(PIPE_ALL);

    // Vector与Cube核间同步，发送同步信息数据到FFTS（地址由FFTS_BASE_ADDR寄存器确定），设置同步点
    ffts_cross_core_sync(PIPE_FIX, 0x0001 | ((2 & 0x0f) << 4) | ((static_cast<uint64_t>(1) & 0x0f) << 8));

    // Vector核
    #elif defined(__DAV_C220_VEC__)

    // Vector初始化配置
    set_vector_mask(-1, -1);
    set_mask_norm();
    set_atomic_none();
    
    // 初始化三个需要用到的tensor首地址，均在UB中
    __ubuf__ float* tensor_c_in_UB = (__ubuf__ float*)get_imm(0);
    __ubuf__ float* tensor_c_out_UB = (__ubuf__ float*)get_imm(2048);
    
    // Vec与Cube核间同步，和ffts_cross_core_sync配套使用（通过flagID关联），用于等待所有同步对象到达flagID对应的同步点。
    wait_flag_dev(1);

    // 数据搬入
    copy_gm_to_ubuf(tensor_c_in_UB, cGm, 0, 1, m * n * sizeof(float)/32, 0, 0);

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

    // 计算
    vrelu(tensor_c_out_UB, tensor_c_in_UB, 16, 1, 1, 8, 8);

    // 同步
    set_flag(PIPE_V, PIPE_MTE3, EVENT_ID1);
    wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID1);

    // 数据搬出
    copy_ubuf_to_gm(cGm, tensor_c_out_UB, 0, 1, m * n * sizeof(float)/32, 0, 0, bm_t::BM_DISABLE);
    
    // 同步
    pipe_barrier(PIPE_ALL);
    #endif
}

void set_ffts_base_addr(uint64_t config)

void ffts_cross_core_sync(pipe_t pipe, uint64_t config)

在将Cube的计算结果（GM上）搬运到UB之前，增加了wait_flag_dev，其接口原型为：

void wait_flag_dev(int64_t flagID)

flagID：与Cube中ffts_cross_core_sync接口设置的flagID相同，也为1。

当Cube侧操作完成后，wait_flag_dev等待结束，Vector侧开始执行。

Vector相关搬运计算接口介绍：

1. 数据搬入。
   将需要进行ReLU计算的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：目标地址，tensor_c_in_UB 。
   • src：源地址，cGm 。
   • sid：预留参数，此处默认为0即可。
   • nBurst：搬运数据块个数，此处不分块，所以nBurst设置为1。
   • lenBurst：一次连续搬运数据长度，单位32Bytes。元素个数为32*32，32*32*sizeof(float) = 4096B，计算搬运长度m*n*sizeof(float)/32。
   • srcGap：两次搬运源地址间隙，源地址处连续搬运，此处设置0。
   • dstGap：两次搬运目的地址间隙，目标地址处连续搬运，此处设置0。
2. 计算。
   将在UB上的tensor进行ReLU计算，这里使用vrelu接口，其接口原型为：

   void vrelu(__ubuf__ float *dst, __ubuf__ float *src, uint8_t repeat, uint16_t dstBlockStride, uint16_t srcBlockStride, uint16_t dstRepeatStride, uint16_t srcRepeatStride);

   • dst：目标地址，tensor_c_out_UB。
   • src：源地址，tensor_c_in_UB。
   • repeat：指令迭代次数，Vector 计算单元每个迭代会从UB中取出8个datablock（每个datablock 数据块内部地址连续，长度 32Byte，合计 256Byte）进行计算，本次是32*32*sizeof(float) = 4096B，4096/256=16次执行完。
   • dstBlockStride：目的操作数不同block间地址步长。单位为32B，因为连续存储，所以dstBlockStride为1。
   • srcBlockStride：源操作数不同block间地址步长。单位为32B，因为src1 block连续读取参与运算，所以srcBlockStride为1。
   • dstRepeatStride：相邻两次执行间，目的操作数相同block间地址步长。单位为32B，因为stride是头到头的距离，所以跨度距离是256B，设为256/32=8。
   • srcRepeatStride：相邻两次执行间，源操作数相同block间地址步长。单位为32B，因为stride是头到头的距离，所以跨度距离是256B，设为256/32=8。
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：源地址，tensor_c_out_UB。
   • sid：预留参数，此处默认为0即可。
   • nBurst：搬运数据块个数，此处不分块，所以nBurst设置为1。
   • lenBurst：一次连续搬运数据长度，单位32Bytes，元素个数为32*32，32*32* sizeof(float)=4096B，计算搬运长度m*n*sizeof(float)/32。
   • srcGap：两次搬运源地址间隙，源地址处连续搬运，此处设置0。
   • dstGap：两次搬运目的地址间隙，目标地址处连续搬运，此处设置0。
   • byteMode：通过设置byteMode枚举参数可改变block大小。设为bm_t::BM_DISABLE，即block大小为32Bytes。

至此，先在Cube完成矩阵乘法计算，设置核间同步，Vector等待其完成后进行ReLU计算的一套Cube+Vector的计算流程结束。