copy_cbuf_to_gm

功能说明

从L1读取数据写入到GM中。

接口原型

void copy_cbuf_to_gm(__gm__ void *dst, __cbuf__ void *src, uint8_t sid, uint16_t nBurst, uint16_t lenBurst, uint16_t srcGap, uint16_t dstGap);

参数说明

参数含义见表1 通用搬运指令参数说明。

流水类型

PIPE_MTE3

用例

// input 为 64 维的 float 向量
extern "C" __global__ __aicore__ void move(
    __gm__ float* __restrict__ input,    // 输入数据指针
    __gm__ float* __restrict__ output    // 输出数据指针
) {
    // 分配 CBuffer 内存区域
    __cbuf__ float* input_cb = (__cbuf__ float*)get_imm(0);    
    
    // 将全局内存的数据搬入 L1
    copy_gm_to_cbuf(input_cb, input, 0, 1, 8, 0, 0, pad_t::PAD_NONE);
         
    // 等待输入数据搬运完成
    set_flag(PIPE_MTE2, PIPE_MTE3, EVENT_ID0);
    wait_flag(PIPE_MTE2, PIPE_MTE3, EVENT_ID0);
    
    // 将 L1 的结果写回全局内存
    copy_cbuf_to_gm(output, input_cb, 0, 1, 8, 0, 0);
        
    // 等待所有操作完成
    pipe_barrier(PIPE_ALL);
}

copy_cbuf_to_gm 参数解释：

1、dst: output；

2、src: input_cb ；

3、sid: 无需关注，设置为 0 即可；

4、nBurst: 所有数据作为整块搬运，故设为1；

5、lenBurst: 搬运总字节数为 64 * 4 = 256，以 32B 为单位，故为 8。

6、srcGap: 由于是单个数据块搬运，所以间隙为0。

7、dstGap: 由于是单个数据块搬运，所以间隙为0。