PhiloxRandom

产品支持情况

功能说明

基于Philox随机数生成算法，给定随机数种子，生成若干的随机数。

Philox随机数生成的核心算法是一个基于记数的伪随机数生成算法，输入为一个128bit的记数器C，两个32bit的key（k0和k1），输出为4个32bit的整数。

函数原型

• 连续模式

  template <uint16_t Rounds = 7, typename T>
  __aicore__ inline void PhiloxRandom(const LocalTensor<T>& dstLocal, const PhiloxKey& philoxKey, const PhiloxCounter& philoxCounter, uint16_t count)
  

• stride模式

  template <uint16_t Rounds = 7, typename T>
  __aicore__ inline void PhiloxRandom(const LocalTensor<T>& dstLocal, const PhiloxKey& philoxKey, const PhiloxCounter& philoxCounter, const PhiloxRandomParams& params)
  

参数说明

表 1 模板参数说明

表 2 参数说明

using PhiloxKey = uint32_t[2];

using PhiloxCounter = uint32_t[4];

struct PhiloxRandomParams {
   uint32_t stride;  // 两行元素之间的间隔
   uint32_t row;     // 表示生成的行数
   uint32_t column;  // 表示生成的每一行的元素个数
}

图 1 PhiloxRandom示意图

上图是一个生成随机数的示意图。

• 连续模式下使用philoxCounter={0, 0, 0, 0}，count=32来生成32个随机数。
• stride模式下可按列分两次生成，调用两次接口。第一次调用参数为philoxCounter={0, 0, 0, 0}，stride=8，row=4，column=4；第二次调用参数为philoxCounter={1, 0, 0, 0}（每次记数器C自增会生成128bit的随机数），stride=8，row=4，column=4。

返回值说明

无

约束说明

无

调用示例

• 接口使用样例

  // philoxKey={0,0}, philoxCounter={0,0,0,0}, params={1024, 8, 1024}
  LocalTensor<uint32_t> dstLocal = outQueue.AllocTensor<uint32_t>();  
  PhiloxRandom<7>(dstLocal, {0, 0}, {0, 0, 0, 0},{1024, 8, 1024});
  

• 完整样例

  #include "kernel_operator.h"
  
  template <uint16_t Rounds, typename srcType>
  class KernelPhiloxStride {
  public:
      __aicore__ inline KernelPhiloxStride() {}
      __aicore__ inline void Init(GM_ADDR dstGm, uint32_t paramStride, uint32_t paramRow, uint32_t paramColumn)
      {
          stride = paramStride;
          row = paramRow;
          column = paramColumn;
          count = row * column;
          const int alginSize = AscendC::GetDataBlockSizeInBytes() / sizeof(srcType);
          dstSize = (count + alginSize - 1) / alginSize * alginSize;
          dstGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ srcType *>(dstGm), dstSize);
          pipe.InitBuffer(outQueue, 1, dstSize * sizeof(srcType));
      }
      __aicore__ inline void Process(uint32_t seed0, uint32_t seed1, uint32_t seed2, uint32_t seed3, uint32_t seed4,
          uint32_t seed5)
      {
          Compute(seed0, seed1, seed2, seed3, seed4, seed5);
          CopyOut();
      }
  private:
      __aicore__ inline void Compute(uint32_t seed0, uint32_t seed1, uint32_t seed2, uint32_t seed3, uint32_t seed4,
          uint32_t seed5)
      {
          AscendC::LocalTensor<srcType> dstLocal = outQueue.AllocTensor<srcType>();
          AscendC::PhiloxRandom<Rounds>(dstLocal, { seed0, seed1 }, { seed2, seed3, seed4, seed5 },
              { stride, row, column });
          outQueue.EnQue<srcType>(dstLocal);
      }
      __aicore__ inline void CopyOut()
      {
          AscendC::LocalTensor<srcType> dstLocal = outQueue.DeQue<srcType>();
          AscendC::DataCopy(dstGlobal, dstLocal, dstSize);
          outQueue.FreeTensor(dstLocal);
      }
  private:
      AscendC::GlobalTensor<srcType> dstGlobal;
      AscendC::TPipe pipe;
      AscendC::TQue<AscendC::TPosition::VECOUT, 1> outQueue;
      uint32_t count;
      uint32_t stride;
      uint32_t row;
      uint32_t column;
      uint32_t dstSize;
  };
  
  extern "C" __global__ __aicore__ void philox_kernel_stride(GM_ADDR dstGm)
  {
      KernelPhiloxStride<7, uint32_t> op;
      op.Init(dstGm, 1024, 8, 1024);
      op.Process(0, 0, 0, 0, 0, 0);
  }