更多样例

样例模板

为了方便您快速运行具体指令中的参考样例，本章节提供双目指令的样例模板。

您可以将以下样例模板作为代码框架，只需将具体指令中的样例片段拷贝替换下文代码段中的加粗内容即可。

#include "kernel_operator.h"
 
class KernelAdd {
public:
    __aicore__ inline KernelAdd() {}
    __aicore__ inline void Init(__gm__ uint8_t* src0Gm, __gm__ uint8_t* src1Gm, __gm__ uint8_t* dstGm)
    {
        src0Global.SetGlobalBuffer((__gm__ int16_t*)src0Gm);
        src1Global.SetGlobalBuffer((__gm__ int16_t*)src1Gm);
        dstGlobal.SetGlobalBuffer((__gm__ int16_t*)dstGm);
        pipe.InitBuffer(inQueueSrc0, 1, 512 * sizeof(int16_t));
        pipe.InitBuffer(inQueueSrc1, 1, 512 * sizeof(int16_t));
        pipe.InitBuffer(outQueueDst, 1, 512 * sizeof(int16_t));
    }
    __aicore__ inline void Process()
    {
        CopyIn();
        Compute();
        CopyOut();
    }
private:
    __aicore__ inline void CopyIn()
    {
        AscendC::LocalTensor<int16_t> src0Local = inQueueSrc0.AllocTensor<int16_t>();
        AscendC::LocalTensor<int16_t> src1Local = inQueueSrc1.AllocTensor<int16_t>();
        AscendC::DataCopy(src0Local, src0Global, 512);
        AscendC::DataCopy(src1Local, src1Global, 512);
        inQueueSrc0.EnQue(src0Local);
        inQueueSrc1.EnQue(src1Local);
    }
    __aicore__ inline void Compute()
    {
        AscendC::LocalTensor<int16_t> src0Local = inQueueSrc0.DeQue<int16_t>();
        AscendC::LocalTensor<int16_t> src1Local = inQueueSrc1.DeQue<int16_t>();
        AscendC::LocalTensor<int16_t> dstLocal = outQueueDst.AllocTensor<int16_t>();
 
        AscendC::Add(dstLocal, src0Local, src1Local, 512);

        outQueueDst.EnQue<int16_t>(dstLocal);
        inQueueSrc0.FreeTensor(src0Local);
        inQueueSrc1.FreeTensor(src1Local);
    }
    __aicore__ inline void CopyOut()
    {
        AscendC::LocalTensor<int16_t> dstLocal = outQueueDst.DeQue<int16_t>();
        AscendC::DataCopy(dstGlobal, dstLocal, 512);
        outQueueDst.FreeTensor(dstLocal);
    }
private:
    AscendC::TPipe pipe;
    AscendC::TQue<AscendC::QuePosition::VECIN, 1> inQueueSrc0, inQueueSrc1;
    AscendC::TQue<AscendC::QuePosition::VECOUT, 1> outQueueDst;
    AscendC::GlobalTensor<int16_t> src0Global, src1Global, dstGlobal;
};
 
extern "C" __global__ __aicore__ void add_simple_kernel(__gm__ uint8_t* src0Gm, __gm__ uint8_t* src1Gm, __gm__ uint8_t* dstGm)
{
    KernelAdd op;
    op.Init(src0Gm, src1Gm, dstGm);
    op.Process();
}

更多样例

您可以参考以下样例，了解如何使用双目指令的tensor高维切分计算接口，进行更灵活的操作、实现更高级的功能。

如果您需要运行样例代码，请将代码段拷贝并替换双目指令样例模板中Compute函数的加粗部分代码即可。

• 通过tensor高维切分计算接口中的mask连续模式，实现数据非连续计算。

  uint64_t mask = 64;  // 每个迭代内只计算前64个数
  AscendC::Add(dstLocal, src0Local, src1Local, mask, 4, { 1, 1, 1, 8, 8, 8 });

  结果示例如下：

  输入数据(src0Local): [1 2 3 ... 512]
  输入数据(src1Local): [513 514 515 ... 1024]
  输出数据(dstLocal): 
  [514 516 518 ... 640 undefined ... undefined
   770 772 774 ... 896 undefined ... undefined
   1026 1028 1030 ... 1152 undefined ... undefined
   1282 1284 1286 ... 1408 undefined ... undefined]

• 通过tensor高维切分计算接口中的mask逐比特模式，实现数据非连续计算。

  uint64_t mask[2] = { UINT64_MAX, 0 };  // mask[0]满，mask[1]空，每次只计算前64个数
  AscendC::Add(dstLocal, src0Local, src1Local, mask, 4, { 1, 1, 1, 8, 8, 8 });

  结果示例如下：

  输入数据(src0Local): [1 2 3 ... 512]
  输入数据(src1Local): [513 514 515 ... 1024]
  输出数据(dstLocal): 
  [514 516 518 ... 640 undefined ... undefined
   770 772 774 ... 896 undefined ... undefined
   1026 1028 1030 ... 1152 undefined ... undefined
   1282 1284 1286 ... 1408 undefined ... undefined]

• 通过控制tensor高维切分计算接口的repeatStride参数，实现数据非连续计算。

  uint64_t mask = 128;
  // repeatTimes设置为2，表示一共需要进行2次迭代
  // src0BlkStride, src1BlkStride设置为1，表示每个迭代内src0参与计算的数据地址间隔为1个datablock
  // src0RepStride设置为16, 表示相邻迭代之间src0起始地址间隔为16个datablock
  AscendC::Add(dstLocal, src0Local, src1Local, mask, 2, { 1, 1, 1, 8, 16, 8 });

  结果示例如下：

  输入数据(src0Local): [1 2 3 ... 512]
  输入数据(src1Local): [513 514 515 ... 1024]
  输出数据(dstLocal):
  [514 516 518 ...768 898 900 902 ... 1150 1152 undefined ... undefined]

• 通过控制tensor高维切分计算接口的dataBlockStride和repeatStride参数，实现数据非连续计算。

  uint64_t mask = 128;
  // repeatTimes设置为2，表示一共需要进行2次迭代
  // src0BlkStride设置为2，表示每个迭代内src0参与计算的数据地址间隔为2个datablock
  // src0RepStride设置为16, 表示相邻迭代之间src0起始地址间隔为16个datablock
  AscendC::Add(dstLocal, src0Local, src1Local, mask, 2, { 1, 2, 1, 8, 16, 8 });

  结果示例如下：

  输入数据(src0Local): [1 2 3 ... 512]
  输入数据(src1Local): [513 514 515 ... 1024]
  输出数据(dstLocal):
  [514 516 518 ... 544  562 564 566 ... 592  610 612 614 ... 640  658 660 662 ... 688
   706 708 710 ... 736  754 756 758 ... 784  802 804 806 ... 832  850 852 854 ... 880 
   898 900 902 ... 928  946 948 950 ... 976  994 996 998 ... 1024  1042 1044 1046 ... 1072
  1090 1092 1094 ... 1120  1138 1140 1142 ... 1168  1186 1188 1190 ... 1216 1234 1236 1238 … 1264
  undefined ... undefined]