Brcb

功能说明

给定一个输入张量，每一次取输入张量中的8个数填充到结果张量的8个datablock（32Bytes）中去，每个数对应一个datablock。

定义原型

template <typename T>
__aicore__ inline void Brcb(const LocalTensor<T>& dstLocal, const LocalTensor<T>& src0Local, const uint8_t repeatTimes, const BrcbRepeatParams& repeatParams)

参数说明

表1 参数说明
参数名称	输入/输出	含义
dstLocal	输出	目的操作数。类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。支持数据类型(uint16_t/uint32_t/half/float), 地址需要32bytes对齐。
srcLocal	输入	源操作数，连续存储score的elements。类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。数据类型和dstLocal保持一致。
repeatTimes	输入	指令迭代次数，每次迭代完成8个block的数据收集，数据范围：repeatTimes∈[0,255]。
repeatParams	输入	指令迭代参数，类型为BrcbRepeatParams，参数说明参考表2。

表2 BrcbRepeatParams结构体内参数说明
参数名称	输入/输出	含义
dstBlkStride	输出	单次迭代内，矢量目的操作数不同block间地址步长。注意事项：当dstBlkStride值为0时，默认按照1来处理。
dstRepStride	输入	相邻迭代间，矢量目的操作数相同block地址步长。

支持的型号

Atlas A2训练系列产品/Atlas 800I A2推理产品。

Atlas推理系列产品（Ascend 310P处理器）AI Core

注意事项

不支持srcLocal与dstLocal为同一块内存地址

针对Atlas推理系列产品（Ascend 310P处理器）AI Core，使用时需要预留8K的Unified Buffer空间，作为接口的临时数据存放区。

调用示例

uint16_t数据类型brcb示例

#include "kernel_operator.h"
namespace AscendC {
class VbrcbCase {
public:
    __aicore__ inline VbrcbCase()
    {}
    __aicore__ inline void Init(__gm__ uint8_t *x, __gm__ uint8_t *y)
    {
        x_gm.SetGlobalBuffer(reinterpret_cast<__gm__ uint16_t *>(x));
        y_gm.SetGlobalBuffer(reinterpret_cast<__gm__ uint16_t *>(y));
        tpipe.InitBuffer(vecIn, 1, 16 * sizeof(uint16_t));
        tpipe.InitBuffer(vecOut, 1, 256 * sizeof(uint16_t));
    }

    __aicore__ inline void Process()
    {
        CopyIn();
        Compute();
        CopyOut();
    }

    __aicore__ inline void CopyIn()
    {
        auto x_buf = vecIn.AllocTensor<uint16_t>();
        AscendC::DataCopy(x_buf, x_gm, 16);
        vecIn.EnQue(x_buf);
    }

    __aicore__ inline void Compute()
    {
        auto x_buf = vecIn.DeQue<uint16_t>();
        auto y_buf = vecOut.AllocTensor<uint16_t>();
        Brcb(y_buf, x_buf, 2, {1,8});
        vecOut.EnQue(y_buf);
        vecIn.FreeTensor(x_buf);
    }

    __aicore__ inline void CopyOut()
    {
        auto y_buf = vecOut.DeQue<uint16_t>();
        AscendC::DataCopy(y_gm, y_buf, 256);
        vecOut.FreeTensor(y_buf);
    }

private:
    GlobalTensor<uint16_t> x_gm;
    GlobalTensor<uint16_t> y_gm;

    TPipe tpipe;
    TQue<QuePosition::VECIN, 1> vecIn;
    TQue<QuePosition::VECOUT, 1> vecOut;
};
}  // namespace AscendC

extern "C" __global__ __aicore__ void vbrcb_uint16_t_16(__gm__ uint8_t *x, __gm__ uint8_t *y)
{
    AscendC::VbrcbCase op;
    op.Init(x, y);
    op.Process();
}

结果示例：

输入数据(srcGlobal): [1 2 3 ... 16]
输出数据(dstGlobal):[1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 ... 15 15 15 15 15 15 15 15 15 15 15 15 15 15 15 15 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16]

父主题： 数据填充