Device侧代码与说明

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

// Device侧文件：mixkernel.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__ float* __restrict__ tensor_a, __gm__ float* __restrict__ tensor_b, __gm__ float* __restrict__ tensor_c_gm) {
    set_ffts_base_addr((uint64_t)ffts_addr);
    #if defined(__DAV_C220_CUBE__)
    set_l1_3d_size(0);  // 设置l1循环缓冲区大小为整个l1
    set_padding(0);     // 设置pad大小为0
    if (((int32_t)block_idx) < 2) { // tiling相关分核策略，这里只使用两个核
        __cbuf__ float* tensor_a_fractal = (__cbuf__ float*)get_imm(0);       // 在L1上分配矩阵A的起始地址
        __cbuf__ float* tensor_b_fractal = (__cbuf__ float*)get_imm(2048);    // 在L1上分配矩阵B的起始地址
        __cc__ float* tensor_mmad = (__cc__ float*)get_imm(0);                // 在L0C上分配结果矩阵的起始地址
        __ca__ float* tensor_a_zz = (__ca__ float*)get_imm(0);                // 在L0A上分配矩阵A的起始地址
        __cb__ float* tensor_b_zn = (__cb__ float*)get_imm(0);                // 在L0B上分配矩阵B的起始地址
        copy_gm_to_cbuf_multi_nd2nz_b32s(tensor_a_fractal, tensor_a + block_idx * 256, 0, 1, 16, 16, 1, 16, 16, 1, 1);
        set_flag(PIPE_MTE2, PIPE_MTE1, EVENT_ID0);
        copy_gm_to_cbuf_multi_nd2nz_b32s(tensor_b_fractal, tensor_b, 0, 1, 16, 32, 1, 32, 16, 1, 1);
        set_flag(PIPE_MTE2, PIPE_MTE1, EVENT_ID1);
        wait_flag(PIPE_MTE2, PIPE_MTE1, EVENT_ID0);
        load_cbuf_to_ca(tensor_a_zz, tensor_a_fractal, 0, 2, 1, 0, 0, 0, inc);
        wait_flag(PIPE_MTE2, PIPE_MTE1, EVENT_ID1);
        load_cbuf_to_cb(tensor_b_zn, tensor_b_fractal, 0, 4, 1, 0, 0, 0, inc);
        set_flag(PIPE_MTE1, PIPE_M, EVENT_ID0);
        wait_flag(PIPE_MTE1, PIPE_M, EVENT_ID0);
        mad(tensor_mmad, tensor_a_zz, tensor_b_zn, 16, 16, 32, 0, 0, 0, 1);
        set_flag(PIPE_M, PIPE_FIX, EVENT_ID0);
        set_nd_para(4295032833);
        wait_flag(PIPE_M, PIPE_FIX, EVENT_ID0);
        copy_matrix_cc_to_gm(tensor_c_gm + block_idx * 512, tensor_mmad, 0, 32, 16, 32, 16, 0, NoQuant, 0, 0, 1);
    }
    pipe_barrier(PIPE_ALL);
    ffts_cross_core_sync(PIPE_FIX, 0x0001 | ((2 & 0x0f) << 4) | ((static_cast<uint64_t>(1) & 0x0f) << 8));
    
    #elif defined(__DAV_C220_VEC__)
    set_vector_mask(-1, -1);
    set_mask_norm();
    set_atomic_none();
    if (block_idx < 2) {
        __ubuf__ float* tensor_c_in_UB = (__ubuf__ float*)get_imm(0);
        __ubuf__ float* tensor_c_out_UB = (__ubuf__ float*)get_imm(2048);
        wait_flag_dev(1);
        copy_gm_to_ubuf_align_b32(tensor_c_in_UB, tensor_c_gm + block_idx * 512, 0, 1, 512 * sizeof(float), 0, 0, 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, 8, 1, 1, 8, 8);
        set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);
        wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);
        copy_ubuf_to_gm_align_b32(tensor_c_gm + block_idx * 512, tensor_c_out_UB, 0, 1, 512 * 4, 0, 0, 0, 0);
    }
    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：与AIC中ffts_cross_core_sync接口设置的flagID相同，也为1。

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

后续操作与Vector算子示例类似。

1. 数据搬入

   先将需要计算的数据从GM上搬入到UB，使用接口原型为：

   void copy_gm_to_ubuf_align_b32(__ubuf__ void *dst, __gm__ void *src, uint8_t sid, uint16_t nBurst, uint32_t lenBurst, uint8_t leftPaddingNum, uint8_t rightPaddingNum, uint32_t srcGap, uint32_t dstGap)

   • dst：UB上的起始地址。
   • src：GM上的起始地址，每个Vec核读取对应Cube核产出的数据，所以偏移block_idx * 512。
   • sid：预留参数，此处默认为0即可。
   • nBurst：只需循环一次。
   • lenBurst：512个float类型元素所占Bytes。
   • leftPaddingNUm：无pad。
   • rightPaddingNum：无pad。
   • srcGap：只循环一次，无关参数。
   • dstGap：只循环一次，无关参数。
2. 计算

   由于后续计算依赖搬运的数据，所以需要一对同步指令控制，当搬运完到达同步点后，开始进行ReLU计算，其使用的接口原型为：

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

   • dst：UB上计算输出结果的首地址。
   • src：源数据搬入到UB上的首地址。
   • repeat：重复次数，512个float类型数据每次计算8*32B（block），需要循环8次。
   • dstBlockStride：数据连续设置为1。
   • srcBlockStride：数据连续设置为1。
   • dstRepeatStride：数据连续，一次循环计算8个block，间隔为8。
   • srcRepeatStride：数据连续，一次循环计算8个block，间隔为8。

待数据计算完成后使用copy_ubuf_to_gm_align_b32接口将数据搬出到GM上，接口使用与搬入时类似，可参考使用。

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