CANN/asc-devkit SetFixPipeAddr寄存器配置

# SetFixPipeAddr【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit产品支持情况Ascend 950PR/Ascend 950DT不支持Atlas A3 训练系列产品/Atlas A3 推理系列产品不支持Atlas A2 训练系列产品/Atlas A2 推理系列产品不支持Atlas 200I/500 A2 推理产品支持Atlas 推理系列产品AI Core不支持Atlas 推理系列产品Vector Core不支持Atlas 训练系列产品不支持功能说明DataCopy数据搬运L0C Buffer-GM过程中进行随路量化后通过调用该接口设置Elementwise操作时LocalTensor的地址。函数原型template typename T __aicore__ inline void SetFixPipeAddr(const LocalTensorT eleWiseData, uint16_t c0ChStride)参数说明表1参数说明参数名称输入/输出含义eleWiseData输入L1 Buffer上的源操作数。类型为LocalTensor。支持的物理位置为L1 BufferTPosition为A1/B1/C1。起始地址需要保证32字节对齐仅支持half数据类型。c0ChStride输入在L1 Buffer上的C0 channel stride单位是C0_SIZE32字节。eleWiseData沿N方向以C0为单位切分得到的数据块称为C0 channel两块C0 channel的间隔称为C0 channel stride。数据类型LocalTensor支持的数据类型为half。返回值说明无约束说明无调用示例DataCopy随路量化搬运后可以逐个元素加/减一个大小为mSize * nSize的LocalTensor具体LocalTensor地址相关参数需要调用SetFixPipeAddr来设置。__aicore__ inline void SetEleSrcPara(const LocalTensorhalf eleWiseData, uint16_t c0ChStride) { AscendC::SetFixPipeAddr(eleWiseData, c0ChStride); }针对Atlas 200I/500 A2 推理产品完整示例如下。随路格式转换数据搬运通路L0C Buffer-GM。示例Mmad含有矩阵乘偏置左矩阵和右矩阵的数据类型为int8_t结果矩阵的数据类型为int32_t。量化模式DEQF16scalar量化参数为0.5将Mmad计算出的结果由int32_t量化成half并搬出。#ifdef ASCENDC_CPU_DEBUG #include tikicpulib.h #endif #include kernel_operator.h #include ../../instrs/common_utils/register_utils.h template typename dst_T, typename fmap_T, typename weight_T, typename dstCO1_T class KernelCubeDataCopy{ public: __aicore__ inline KernelCubeDataCopy(uint16_t CoutIn, uint8_t dilationHIn, uint8_t dilationWIn, QuantMode_t deqModeIn) { // ceiling of 16 Cout CoutIn; dilationH dilationHIn; dilationW dilationWIn; C0 32 / sizeof(fmap_T); C1 channelSize / C0; coutBlocks (Cout 16 - 1) / 16; ho H - dilationH * (Kh - 1); wo W - dilationW * (Kw - 1); howo ho * wo; howoRound ((howo 16 - 1) / 16) * 16; featureMapA1Size C1 * H * W * C0; // shape: [C1, H, W, C0] weightA1Size C1 * Kh * Kw * Cout * C0; // shape: [C1, Kh, Kw, Cout, C0] featureMapA2Size howoRound * (C1 * Kh * Kw * C0); weightB2Size (C1 * Kh * Kw * C0) * coutBlocks * 16; m howo; k C1 * Kh * Kw * C0; n Cout; biasSize Cout; // shape: [Cout] dstSize coutBlocks * howo * 16; // shape: [coutBlocks, howo, 16] dstCO1Size coutBlocks * howoRound * 16; fmRepeat featureMapA2Size / (16 * C0); weRepeat weightB2Size / (16 * C0); deqMode deqModeIn; } __aicore__ inline void Init(__gm__ uint8_t* fmGm, __gm__ uint8_t* weGm, __gm__ uint8_t* biasGm, __gm__ uint8_t* deqGm, __gm__ uint8_t* eleWiseGm, __gm__ uint8_t* dstGm) { fmGlobal.SetGlobalBuffer((__gm__ fmap_T*)fmGm); weGlobal.SetGlobalBuffer((__gm__ weight_T*)weGm); biasGlobal.SetGlobalBuffer((__gm__ dstCO1_T*)biasGm); deqGlobal.SetGlobalBuffer((__gm__ uint64_t*)deqGm); dstGlobal.SetGlobalBuffer((__gm__ dst_T*)dstGm); eleWiseGlobal.SetGlobalBuffer((__gm__ half*)eleWiseGm); pipe.InitBuffer(inQueueFmA1, 1, featureMapA1Size * sizeof(fmap_T)); pipe.InitBuffer(inQueueFmA2, 1, featureMapA2Size * sizeof(fmap_T)); pipe.InitBuffer(inQueueWeB1, 1, weightA1Size * sizeof(weight_T)); pipe.InitBuffer(inQueueWeB2, 1, weightB2Size * sizeof(weight_T)); pipe.InitBuffer(inQueueBiasA1, 1, biasSize * sizeof(dstCO1_T)); pipe.InitBuffer(inQueueDeqA1, 1, dstCO1Size * sizeof(uint64_t)); pipe.InitBuffer(inQueueDeqFB, 1, dstCO1Size * sizeof(uint64_t)); pipe.InitBuffer(outQueueCO1, 1, dstCO1Size * sizeof(dstCO1_T)); pipe.InitBuffer(inQueueC1, 1, dstSize * sizeof(half)); } __aicore__ inline void Process() { CopyIn(); Split(); Compute(); CopyOut(); } private: __aicore__ inline void CopyIn() { AscendC::LocalTensorfmap_T featureMapA1 inQueueFmA1.AllocTensorfmap_T(); AscendC::LocalTensorweight_T weightB1 inQueueWeB1.AllocTensorweight_T(); AscendC::LocalTensordstCO1_T biasA1 inQueueBiasA1.AllocTensordstCO1_T(); AscendC::DataCopy(featureMapA1, fmGlobal, { 1, static_castuint16_t(featureMapA1Size * sizeof(fmap_T) / 32), 0, 0 }); AscendC::DataCopy(weightB1, weGlobal, { 1, static_castuint16_t(weightA1Size * sizeof(weight_T) / 32), 0, 0 }); AscendC::DataCopy(biasA1, biasGlobal, { 1, static_castuint16_t(biasSize * sizeof(dstCO1_T) / 32), 0, 0 }); inQueueFmA1.EnQue(featureMapA1); inQueueWeB1.EnQue(weightB1); inQueueBiasA1.EnQue(biasA1); } __aicore__ inline void Split() { AscendC::LocalTensorfmap_T featureMapA1 inQueueFmA1.DeQuefmap_T(); AscendC::LocalTensorweight_T weightB1 inQueueWeB1.DeQueweight_T(); AscendC::LocalTensorfmap_T featureMapA2 inQueueFmA2.AllocTensorfmap_T(); AscendC::LocalTensorweight_T weightB2 inQueueWeB2.AllocTensorweight_T(); uint8_t padList[] {0, 0, 0, 0}; // load3dv2 AscendC::LoadData(featureMapA2, featureMapA1, { padList, H, W, channelSize, k, howoRound, 0, 0, 1, 1, Kw, Kh, dilationW, dilationH, false, false, 0 }); // load2d AscendC::LoadData(weightB2, weightB1, { 0, weRepeat, 1, 0, 0, false, 0 }); inQueueFmA2.EnQuefmap_T(featureMapA2); inQueueWeB2.EnQueweight_T(weightB2); inQueueFmA1.FreeTensor(featureMapA1); inQueueWeB1.FreeTensor(weightB1); } __aicore__ inline void Compute() { AscendC::LocalTensorfmap_T featureMapA2 inQueueFmA2.DeQuefmap_T(); AscendC::LocalTensorweight_T weightB2 inQueueWeB2.DeQueweight_T(); AscendC::LocalTensordstCO1_T dstCO1 outQueueCO1.AllocTensordstCO1_T(); AscendC::LocalTensordstCO1_T biasA1 inQueueBiasA1.DeQuedstCO1_T(); // C A * B bias // m左矩阵Heightk左矩阵Widthn右矩阵Width AscendC::Mmad(dstCO1, featureMapA2, weightB2, biasA1, { m, n, k, true, 0, false, false, false }); outQueueCO1.EnQuedstCO1_T(dstCO1); inQueueFmA2.FreeTensor(featureMapA2); inQueueWeB2.FreeTensor(weightB2); } __aicore__ inline void CopyOut() { AscendC::LocalTensordstCO1_T dstCO1 outQueueCO1.DeQuedstCO1_T(); // 开启DEQF16量化量化参数设置为0.5 float tmp (float)0.5; // 将float的tmp转换成uint64_t的deqScalar uint64_t deqScalar static_castuint64_t(*reinterpret_castint32_t*(tmp)); bool nz2ndEn false; // nz2nd不开启时nSize必须为16的倍数 uint16_t nSize coutBlocks * 16; uint16_t mSize m; // srcStride必须为16的倍数 uint16_t srcStride (m 16 - 1) / 16 * 16; // nz2nd不开启时dstStride为burst头到头的距离且为32字节对齐 uint32_t dstStride m * sizeof(dst_T) * 16 / 32; if (nz2ndEn) { // nd矩阵的数量为1src_nd_stride与dst_nd_stride填1 AscendC::SetFixpipeNz2ndFlag(1, 1, 1); // nz2nd开启时nSize可以不为16的倍数与Mmad的n保持一致 nSize n; // nz2nd开启时dstStride表示同一nd矩阵的相邻连续行的间隔与n保持一致 dstStride nSize; }; // 不开启ReLU与channelSplit AscendC::DataCopyCO12DstParams intriParams(nSize, mSize, dstStride, srcStride, deqMode, 0, false, nz2ndEn); // mov l0c to gm, deq scalar quant AscendC::SetFixpipePreQuantFlag(deqScalar); // 设置量化参数 AscendC::PipeBarrierPIPE_FIX(); AscendC::DataCopy(dstGlobal, dstCO1, intriParams); // // mov l0c to gm, deq tensor quant // // 需要额外申请deq tensor的gm空间将值搬运到workA1 // AscendC::LocalTensoruint64_t workA1 inQueueDeqA1.AllocTensoruint64_t(); // // deq tensor的size // uint16_t deqSize 128; // AscendC::DataCopy(workA1, deqGlobal, deqSize); // // deq tensor在fix上的地址 // AscendC::LocalTensoruint64_t deqFB inQueueDeqFB.AllocTensoruint64_t(); // // l1-fix, burst_len unit is 128Bytes // uint16_t fbufBurstLen deqSize / 128; // AscendC::DataCopyParams dataCopyParams(1, fbufBurstLen, 0, 0); // AscendC::DataCopy(deqFB, workA1, dataCopyParams); // // 设置量化tensor // AscendC::SetFixPipeConfig(deqFB); // AscendC::PipeBarrierPIPE_FIX(); // // mov l0c to gm量化操作后开启ClipReLU操作 // intriParams.clipReluPre 1; // // 设置clip ReLU的值到寄存器 // uint64_t clipReluVal 0x3c00; // value 1, half // SetFixPipeClipRelu(clipReluVal); // //mov l0c to gm量化操作后设置element-wise操作Add // intriParams.eltWiseOp 1; // // 需要额外申请element-wise tensor的gm空间将值搬到eleWiseTensor // AscendC::LocalTensorhalf eleWiseTensor inQueueC1.AllocTensorhalf(); // DataCopy(eleWiseTensor, eleWiseGlobal, { 1, static_castuint16_t(sizeof(half) * dst_size / 32), 0, 0 }); // AscendC::PipeBarrierPIPE_ALL(); // // 将存放element-wise tensor的地址设置到寄存器里 // SetFixPipeAddr(eleWiseTensor, 1); // AscendC::DataCopy(dstGlobal, dstCO1, intriParams); // inQueueDeqA1.FreeTensor(workA1); // inQueueDeqFB.FreeTensor(deqFB); // outQueueCO1.FreeTensor(dstCO1); // inQueueC1.FreeTensor(eleWiseTensor); } private: AscendC::TPipe pipe; // feature map queue AscendC::TQueAscendC::TPosition::A1, 1 inQueueFmA1; AscendC::TQueAscendC::TPosition::A2, 1 inQueueFmA2; // weight queue AscendC::TQueAscendC::TPosition::B1, 1 inQueueWeB1; AscendC::TQueAscendC::TPosition::B2, 1 inQueueWeB2; // bias queue AscendC::TQueAscendC::TPosition::A1, 1 inQueueBiasA1; // deq tensor queue AscendC::TQueAscendC::TPosition::A1, 1 inQueueDeqA1; // fb dst of deq tensor AscendC::TQueAscendC::TPosition::C2PIPE2GM, 1 inQueueDeqFB; // dst queue AscendC::TQueAscendC::TPosition::CO1, 1 outQueueCO1; // element-wise tensor AscendC::TQueAscendC::TPosition::C1, 1 inQueueC1; AscendC::GlobalTensorfmap_T fmGlobal; AscendC::GlobalTensorweight_T weGlobal; AscendC::GlobalTensordst_T dstGlobal; AscendC::GlobalTensoruint64_t deqGlobal; AscendC::GlobalTensordstCO1_T biasGlobal; AscendC::GlobalTensorhalf eleWiseGlobal; uint16_t channelSize 32; uint16_t H 4, W 4; uint8_t Kh 2, Kw 2; uint16_t Cout; uint16_t C0, C1; uint8_t dilationH, dilationW; uint16_t coutBlocks, ho, wo, howo, howoRound; uint32_t featureMapA1Size, weightA1Size, featureMapA2Size, weightB2Size, biasSize, dstSize, dstCO1Size; uint16_t m, k, n; uint8_t fmRepeat, weRepeat; QuantMode_t deqMode QuantMode_t::NoQuant; }; #define KERNEL_CUBE_DATACOPY(dst_type, fmap_type, weight_type, dstCO1_type, CoutIn, dilationHIn, dilationWIn, deqModeIn) \ extern C __global__ __aicore__ void cube_datacopy_kernel_##fmap_type(__gm__ uint8_t* fmGm, __gm__ uint8_t* weGm, \ __gm__ uint8_t* biasGm, __gm__ uint8_t* deqGm, __gm__ uint8_t* eleWiseGm, __gm__ uint8_t* dstGm) \ { \ if (g_coreType AscendC::AIV) { \ return; \ } \ KernelCubeDataCopydst_type, fmap_type, weight_type, dstCO1_type op(CoutIn, dilationHIn, dilationWIn, \ deqModeIn); \ op.Init(fmGm, weGm, biasGm, deqGm, eleWiseGm, dstGm); \ op.Process(); \ } KERNEL_CUBE_DATACOPY(half, int8_t, int8_t, int32_t, 128, 1, 1, QuantMode_t::DEQF16);【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考