Brcb

功能说明

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

定义原型

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

参数说明

支持的型号

Atlas A2训练系列产品

注意事项

无

调用示例

• 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]