Device侧代码与说明
Device侧代码示例如下,并参考本节内容了解主要操作。
// Device侧文件:cube_kernel.cce
#ifdef __CCE_KT_TEST__
#define __aicore__
#else
#define __aicore__ [aicore]
#endif
extern "C" __global__ __aicore__ void mat_mul_kernel(__gm__ uint8_t *tensor_a, __gm__ uint8_t *tensor_b, __gm__ uint8_t *tensor_c, int m, int k, int n) {
// 指针类型转换,将通用的字节指针(uint8_t*)转换为特定类型的浮点指针(float*)
__gm__ half *aGm = (__gm__ half *)tensor_a;
__gm__ half *bGm = (__gm__ half *)tensor_b;
__gm__ float *cGm = (__gm__ float *)tensor_c;
// 初始化三个需要用到的tensor首地址,分别在L0A、L0B和L0C中
__cbuf__ half *aL1 = (__cbuf__ half *)get_imm(0);
__cbuf__ half *bL1 = (__cbuf__ half *)get_imm(2048);
__ca__ half *aL0a = (__ca__ half *)get_imm(0);
__cb__ half *bL0b = (__cb__ half *)get_imm(0);
__cc__ float *cL0c = (__cc__ float *)get_imm(0);
// 数据搬入L1,ND格式转为NZ格式
copy_gm_to_cbuf_multi_nd2nz_b16(aL1, aGm, 0, 1, m, n, m*n, n, m, 1, 16*16);
copy_gm_to_cbuf_multi_nd2nz_b16(bL1, bGm, 0, 1, m, n, m*n, n, m, 1, 16*16);
// 同步
set_flag(PIPE_MTE2, PIPE_MTE1, EVENT_ID0);
wait_flag(PIPE_MTE2, PIPE_MTE1, EVENT_ID0);
// 数据搬入L0A和L0B,NZ格式分别转为ZZ格式和ZN格式
for (int i = 0; i < 2; i++) {
load_cbuf_to_ca(aL0a+512*i*2/2, aL1, i, 2, 2, 0, 0, false, inc);
}
for (int i = 0; i < 2; i++) {
load_cbuf_to_cb(bL0b+512*i*2/2, bL1, i, 2, 2, 0, 0, true, inc);
}
// 同步
set_flag(PIPE_MTE1, PIPE_M, EVENT_ID1);
wait_flag(PIPE_MTE1, PIPE_M, EVENT_ID1);
// 计算
mad(cL0c, aL0a, bL0b, m, k, n, 0, 0, 0, 0, 0, 0, true);
// 同步
set_flag(PIPE_M, PIPE_FIX, EVENT_ID2);
wait_flag(PIPE_M, PIPE_FIX, EVENT_ID2);
// 数据搬出
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);
}
- 数据搬入L1,该接口具备分形转换能力,可将ND转为NZ,所以GM中的数据仅需ND格式即可。
将需要进行矩阵计算的两个tensor,从GM搬到L1上,并实现ND2NZ的分形转换,这里使用了copy_gm_to_cbuf_multi_nd2nz_b16接口,其接口原型为:
void copy_gm_to_cbuf_multi_nd2nz_b16(__cbuf__ half *dst, __gm__ half *src, uint8_t sid, uint16_t ndNum, uint16_t nValue, uint16_t dValue, uint16_t srcNdMatrixStride, uint16_t srcDValue, uint16_t dstNzC0Stride, uint16_t dstNzNStride, uint16_t dstNzMatrixStride);
- dst:目的地址,aL1/bL1。
- src:源地址,aGm/bGM。
- sid:预留参数,此处默认为0即可。
- ndNum:待搬运数据的nd块数量。整块搬运,设为1。
- nValue:数据块n方向长度。源矩阵的行数,设为m(=32)。
- dValue:数据块d方向长度。源矩阵的列数,设为n(=32)。
- srcNdMatrixStride:源数据块nd块之间的距离,单位为元素。此处即为源矩阵的大小,m*n。
- srcDValue:用于指示源nd矩阵的d维大小。连续读取,此处即为n(=32)。
- dstNzC0Stride:L1中nz上2个c0块之间的距离,单位为C0_size。连续存储,此处即为m(=32)。
- dstNzNStride:L1中n维上2个c0块之间的距离,单位为C0_size。连续存储,此处即为1。
- dstNzMatrixStride:2个nd矩阵中2个NZ矩阵的距离。连续存储,且只有1个nd块,此处即为分形的大小,16*16。
- 数据搬入L0A/L0B,该接口仅具备将单个分形转置的能力。经过步骤1后L1中的数据为NZ,而L0A和L0B所需的数据为ZZ和ZN,所以在传入L0A和L0B时均需要间隔取分形的操作将分形间的格式由N转为Z。该操作后L0A所需的数据已满足要求,故接口中转置的参数设为false即可,而L0B所需数据的分形内部格式还需由Z转为N,所以接口中转置的参数需要设为true。这里使用了load_cbuf_to_ca/load_cbuf_to_cb接口,其接口原型为:
void load_cbuf_to_ca(__ca__ half *dst, __cbuf__ half *src, uint16_t baseIdx, uint8_t repeat, uint16_t srcStride, uint16_t dstGap, uint8_t sid, bool transpose, __cce_scalar::addr_cal_mode_t addr_cal_mode); void load_cbuf_to_cb(__cb__ half *dst, __cbuf__ half *src, uint16_t baseIdx, uint8_t repeat, uint16_t srcStride, uint16_t dstGap, uint8_t sid, bool transpose, __cce_scalar::addr_cal_mode_t addr_cal_mode);
- dst:目的地址,分别为aL0a+512*i*2/2和bL0b+512*i*2/2。因为使用for循环间隔取分形来改变分形间的格式,所以目标地址需要加偏移量,偏移量为元素个数。每次搬运512B*2(repeat次数)/2(half大小)的元素个数,所以偏移量为512*i*2/2。
- src:源地址,aGm/bGM。源地址无需偏移量,可使用baseIdx实现偏移。
- baseIdx:表示src matrix的fractal matrix的index ID。每次循环从分形i开始读取,实现源地址的偏移,所以设为i。
- repeat:重复搬运的次数。每次搬512字节的分形,32*32的half类型矩阵包含4个分形,通过for循环间隔取分形来改变分形间的格式,每次循环中需要搬运2次,故设为2。
- srcStride:每个repeat迭代的源地址stride,单位是16*16的fractal matrix,每个repeat迭代的源地址分形头到第二个分形头的距离。此处需间隔取分形以修改分形间的格式,设为2。
- dstGap:每个repeat迭代的目的地址gap,单位是16*16的fractal matrix,每个repeat迭代的目的地址分形尾到第二个分形头的距离。连续存储时取0。
- sid:预留参数,此处默认为0即可。
- transpose:该参数仅在src={L1}, dst={L0A, L0B},且分形的数据类型为b16时可以设为true,可以实现单个分形的转置。由于L1的格式为NZ,L0A需要的格式为ZZ,所以分形内部的格式是一致的,设为false。而L0B需要的格式为ZN,所以分形内部的格式是不同的,设为true。
- addr_cal_mode:设为inc,表示src地址是增加,数据连续往后顺序读取;设为dec,表示从src往前读取,地址减小。需要递增读取,设为inc。
- 计算。
将在L0A/L0B上的两个矩阵进行矩阵相乘,这里使用mad接口,其接口原型为:
void mad(__cc__ float *dst, __ca__ float *src0, __cb__ float *src1, uint16_t m, uint16_t k, uint16_t n, uint8_t featOffset, uint8_t smaskOffset, uint8_t unitFlag, bool kDirectionAlign, bool isWeightOffset, bool cmatrixSource, bool cmatrixInitVal);
- c:目的地址,cL0c。
- a:矩阵A地址,aL0a。
- b:矩阵B地址,bL0ab。
- m:A矩阵的高,传入参数为m(=32)。
- k:A矩阵的宽或B矩阵的高,传入参数为k(=32)。
- n:B矩阵的宽,传入参数为n(=32)。
- featOffset:特征图矩阵偏移。该参数在当前版本未启用,设为0即可。
- smaskOffset:SMASK缓冲区地址。该参数在当前版本未启用,设为0即可。
- unitFlag:设置为0即可。
- kDirectionAlign:此位仅用于(源矩阵类型为f32, 并且目的矩阵类型也是f32,即f32*f32+f32),其他类型忽略此位。如果=1,则L0AL0B中的矩阵在K方向上对齐到16。否则它将与8对齐。设为0即可。
- isWeightOffset:权重矩阵偏移使能位。该参数在当前版本未启用,设为0即可。
- cmatrixcSource:当cmatrixInitVal设为False时,该参数才有意义。如果cmatrixSource=0,则C矩阵在L0C中,其在L0C中的地址与c[31:0]中的地址相同。如果cmatrixSource=1,则C矩阵位于偏置表中。设为0即可。
- cmatrixInitVal:表示矩阵c初始值控制位,True表示c矩阵初始值为0,False使用c矩阵初始值使用其中的具体数据。设置为true,使得C矩阵初始值为0。
- 数据搬出。
将计算完的结果tensor,从L0C上搬出结果到GM指定位置,并实现NZ2ND的分形转换,这里使用了copy_matrix_cc_to_gm接口,同时为了实现NZ2ND,还需要使用set_nd_para对进行设置。
copy_matrix_cc_to_gm接口的原型为:
void copy_matrix_cc_to_gm(__gm__ float *dst, __cc__ float *src, uint8_t sid, uint16_t NSize, uint16_t MSize, uint32_t dstStride_dst_D, uint16_t srcStride, uint8_t UnitFlagMode, uint64_t QuantPRE, uint8_t ReLUPRE, bool channelSplit, bool NZ2ND_EN);
- dst:目的地址,cGm。
- src:源地址,cL0c。
- sid:预留参数,此处默认为0即可。
- NSize:矩阵C的宽,传入参数为n(=32)。
- MSize:矩阵C的高,传入参数为m(=32)。
- dstStride_dst_D:因为使能了NZ2ND,且目的ND矩阵为连续存储,所以该参数为矩阵C的宽,传入参数为n(=32)。
- srcStride:L0C源矩阵中不同数据块间距离。连续读取时为矩阵C的高,传入参数为m(=32)。
- UnitFlagMode:填0即可。
- QuantPRE:预量化模式,预留参数,设为0即可,代表不做量化。
- ReLUPRE:ReLU模式。不使用ReLU,设为0。
- channelSpilt:是否使能通道拆分。不进行通道切分,设为0。
- NZ2ND_EN:是否使能NZ2ND_EN格式转换。设为1,实现NZ2ND的转换。
set_nd_para接口的原型为:
void set_nd_para(uint64_t config);
- config[0:15]位:表示nd块数量。nd块数量为1。
- config[16:31]位:表示源数据nd块步长,其单位为分形大小。连续读取为1。
- config[32:47]位:表示目的数据nd块步长,其单位为元素。连续存储为32*32=1024。
- config[0:47]为:0100 0000 0000 0000 0000 0000 0001 0000 0000 0000 0001,其十进制为4398046576641。
父主题: 示例2