寒武纪显卡实现softmax的pingpong流水并行

在上一篇文章添加链接描述中我们介绍了寒武纪显卡实现基本的softmax代码,这里我们借助于寒武纪的流水并行来编写进一步的策略。

pingpongGDRAM2NRAM流水

仅仅计算max和sum使用流水

我们先考虑不使用SRAM的流水,我们设置两个NRAM上的长度为maxNum上的数组src0和src1以及两个NRAM上的指针read和write,一开始设置write=src0,并且使用memcpy把sourc1的开头maxNum数据加载到write上,然后开始进入循环,其中循环次数为repeat-1。在循环内部不断交换read和write指针指向的数组,内部使用memcpy_async把后面的数据加载到write里面,发起这段命令以后马上针对read的数据进行计算,参考表格如下:
在这里插入图片描述

循环结束以后,发现最后还有maxNum的元素存储在write里面没有计算,为此需要针对这部分数据特殊处理,最后针对remain这部分不能整除的数据继续特殊处理即可。这里一定要注意的就是在循环内部一定要在计算结束以后加入同步机制__sync_all_ipu();
但是在计算出全局最大值和数值和以后,我们还要重新从GDRAM中读取数据到NRAM实现指数变换,在这个过程我们可以类似的使用这种方法来做流水,此时如果仅仅针对数据从GDRAM到NRAM和计算这两部分做流水,那么这个循环就没必要做__sync_all_ipu(),因为计算结束以后把数据写回GDRAM的这个过程可以保证数据结束,参考下面这个代码。

#include <bang.h>
#include <bang_device_functions.h>
#define EPS 1e-7
const int NRAM_MAX_SIZE = 1024 * 128;//后续树状求和必须保证NRAM_MAX_SIZE为2的幂次
const int maxNum = NRAM_MAX_SIZE/sizeof(float); //NRAM上最多存储maxNum个float元素
const int warpSize = 32;//__bang_reduce_sum每次从src取128字节数据相加,对应32个float元素,并且0-31的结果保存在索引0,32-63的结果保存在索引1

__nram__ float src1[maxNum];//每次搬运maxNum数据到NRAM
__nram__ float src0[maxNum];//每次搬运maxNum数据到NRAM
__nram__ float destSum[maxNum];//后面数值求和
__nram__ float destSumFinal[warpSize];//将destSum规约到destFinal[0]
__nram__ float srcMax[2];

__mlu_entry__ void softmaxKernel(float* dst, float* source1, float* globalMax, float* globalSum, int num) {
  
  int remain = num%taskDim;//如果不能整除,则让前部分taskId多处理一个元素
  int stepEasy = (num - remain)/taskDim;
  int stepHard = stepEasy + 1;
  int step = (taskId < remain ? stepHard : stepEasy);//前部分taskId多处理一个元素
  int indStart = (taskId < remain ? taskId * stepHard : remain * stepHard + (taskId - remain) * stepEasy);
  int remainNram = step%maxNum;
  int repeat = step/maxNum;//如果一个task处理元素个数超出NRAM最大内存,则需要for循环
  //maxNum尽量取大一些,免得repeat过大导致求和过程累加过于严重,使得大数吃小数
  source1 = source1 + indStart;//设定起始偏移量

  //------------------------------------下面开始计算max
  __nram__ float destOldMax;
  __nram__ float destNewMax;
  __bang_write_zero(destSum, maxNum);
  destNewMax = -INFINITY;//初始化为负无穷
  __nram__ float *read;
  __nram__ float *write;
  write = src0;
  __memcpy(write, source1, NRAM_MAX_SIZE, GDRAM2NRAM);
  for(int i = 0; i < repeat - 1; i++){
    if(i%2 == 0){
      read = src0;
      write = src1;
    }
    else{
      read = src1;
      write = src0;
    }
    __memcpy_async(write, source1 + (i + 1) * maxNum, NRAM_MAX_SIZE, GDRAM2NRAM);
    __bang_argmax(srcMax, read, maxNum);//针对taskId处理的这step数据,借助于for循环把信息集中到长度为maxNum的向量src1中
    if(destNewMax < srcMax[0]){
      destNewMax = srcMax[0];//更新最大值
    }
    __bang_sub_scalar(read, read, destNewMax, maxNum);//src1 = src1 - 最大值
    __bang_active_exp_less_0(read, read, maxNum);//src1 = exp(src1 - 最大值)
    if(i > 0){
      __bang_mul_scalar(destSum, destSum, exp(destOldMax - destNewMax), maxNum);//destSum = destSum * exp(destOldMax - destNewMax)
    }
    __bang_add(destSum, destSum, read, maxNum);//destSum = destSum + exp(src1 - destNewMax)
    destOldMax = destNewMax;
    __sync_all_ipu();//必须同步
  }
  //------------特殊处理最后一部分
  __bang_argmax(srcMax, write, maxNum);//针对taskId处理的这step数据,借助于for循环把信息集中到长度为maxNum的向量src1中
  if(destNewMax < srcMax[0]){
    destNewMax = srcMax[0];//更新最大值
  }
  __bang_sub_scalar(write, write, destNewMax, maxNum);//src1 = src1 - 最大值
  __bang_active_exp_less_0(write, write, maxNum);//src1 = exp(src1 - 最大值)
  __bang_mul_scalar(destSum, destSum, exp(destOldMax - destNewMax), maxNum);
  __bang_add(destSum, destSum, write, maxNum);//destSum = destSum + exp(src1 - destNewMax)
  destOldMax = destNewMax;
  //-------------特殊处理结束
  if(remainNram){
    __bang_write_value(src1, maxNum, -INFINITY);//必须要初始化src1全部元素为负无穷
    __memcpy(src1, source1 + repeat * maxNum, remainNram * sizeof(float), GDRAM2NRAM);
    __bang_argmax(srcMax, src1, maxNum);//针对taskId处理的这step数据,借助于for循环把信息集中到长度为maxNum的向量src1中
    if(destNewMax < srcMax[0]){
      destNewMax = srcMax[0];
    }
    __bang_write_value(src1, maxNum, destNewMax);//必须重新初始化为destNewMax
    __memcpy(src1, source1 + repeat * maxNum, remainNram * sizeof(float), GDRAM2NRAM);//必须再次读取
    __bang_sub_scalar(src1, src1, destNewMax, maxNum);//后面maxNum-remainNram部分为0
    __bang_active_exp_less_0(src1, src1, maxNum);//相当于多加了maxNum-remainNram
    __bang_mul_scalar(destSum, destSum, exp(destOldMax - destNewMax), maxNum);
    __bang_add(destSum, destSum, src1, maxNum);
    destOldMax = destNewMax;
  }//结束以后向量destNewMax保存了source1[indSart:indStart+step]这部分数据的全局最大值,destSum保存数值和
  //----------
  __bang_write_zero(destSumFinal, warpSize);//初始化destSumFinal全部元素为0
  int segNum = maxNum / warpSize;//将destSum分成segNum段,每段向量长度为warpSize,分段进行树状求和,segNum要求是2的幂次
  for(int strip = segNum/2; strip > 0; strip = strip / 2){//segNum要求是2的幂次即maxNum必须选取2的幂次
    for(int i = 0; i < strip ; i++){
      __bang_add(destSum + i * warpSize, destSum + i * warpSize, destSum + (i + strip) * warpSize, warpSize);
    } 
  }
  __bang_reduce_sum(destSumFinal, destSum, warpSize);
  destSumFinal[0] = destSumFinal[0] - (maxNum - remainNram);//把上面多加的(maxNum - remainNram)减掉
  
  //----------
  globalMax[0] = -INFINITY;
  globalSum[0] = 0.0;
  __sync_all();
  __bang_atomic_max(&destNewMax, globalMax, &destNewMax, 1);//globalMax[0]必须初始化为负无穷
  destSumFinal[0] = destSumFinal[0] * exp(destOldMax - globalMax[0]);
  //__bang_printf("taskId:%d, step:%d, sum:%.6f\n", taskId, step, destSumFinal[0]);
  __sync_all();
  __bang_atomic_add(destSumFinal, globalSum, destSumFinal, 1);//globalSum[0]必须初始化为0
  
  
  dst = dst + indStart;//设定起始偏移量
  float globalSumInv = 1.0/globalSum[0];
  
  write = src0;
  __memcpy(write, source1, NRAM_MAX_SIZE, GDRAM2NRAM);
  for(int i = 0; i < repeat - 1; i++){
    if(i%2 == 0){
      read = src0;
      write = src1;
    }
    else{
      read = src1;
      write = src0;
    }
    __memcpy_async(write, source1 + (i + 1) * maxNum, NRAM_MAX_SIZE, GDRAM2NRAM);
    __bang_sub_scalar(read, read, globalMax[0], maxNum);//src1 = src1 - globalMax[0] 
    __bang_active_exp_less_0(read, read, maxNum);//src1 = exp(src1 - globalMax[0])
    __bang_mul_scalar(read, read, globalSumInv, maxNum);//倒数和另一个向量逐元素相乘得到除法结果
    __memcpy(dst + i * maxNum, read, NRAM_MAX_SIZE, NRAM2GDRAM);
  }
  //-----------特殊处理最后一部分
  __bang_sub_scalar(write, write, globalMax[0], maxNum);//src1 = src1 - globalMax[0] 
  __bang_active_exp_less_0(write, write, maxNum);//src1 = exp(src1 - globalMax[0])
  __bang_mul_scalar(write, write, globalSumInv, maxNum);//倒数和另一个向量逐元素相乘得到除法结果
  __memcpy(dst + (repeat - 1) * maxNum, write, NRAM_MAX_SIZE, NRAM2GDRAM);
  if(remainNram){ 
    __bang_write_value(src1, maxNum, globalMax[0]);
    __memcpy(src1, source1 + repeat * maxNum, remainNram * sizeof(float), GDRAM2NRAM);
    __bang_sub_scalar(src1, src1, globalMax[0], maxNum);//src1 = src1 - globalMax[0] 
    __bang_active_exp_less_0(src1, src1, maxNum);//src1 = exp(src1 - globalMax[0])
    __bang_mul_scalar(src1, src1, globalSumInv, maxNum);//倒数和另一个向量逐元素相乘得到除法结果
    __memcpy(dst + repeat * maxNum, src1, remainNram * sizeof(float), NRAM2GDRAM);
  }
  
  __bang_printf("taskId:%d,repeat:%d,max:%.6f, sum:%.6f\n",taskId, repeat, globalMax[0], globalSum[0]);
  

}


int main(void)
{
  int num = 1024 * 1024 * 1024;
  //int num = 11;
  cnrtQueue_t queue;
  CNRT_CHECK(cnrtSetDevice(0));
  CNRT_CHECK(cnrtQueueCreate(&queue));

  cnrtDim3_t dim = {4, 1, 1};
  int taskNum = dim.x * dim.y * dim.z;
  cnrtFunctionType_t ktype = CNRT_FUNC_TYPE_UNION1;

  cnrtNotifier_t start, end;
  CNRT_CHECK(cnrtNotifierCreate(&start));
  CNRT_CHECK(cnrtNotifierCreate(&end));

  float* host_dst = (float*)malloc(num * sizeof(float));
  float* host_src1 = (float*)malloc(num * sizeof(float));
  

  for (int i = 0; i < num; i++) {
    host_src1[i] = i%4;
  }

  float* mlu_dst;
  float* mlu_src1;
  float* globalMax;
  float* globalSum;
  CNRT_CHECK(cnrtMalloc((void**)&mlu_dst, num * sizeof(float)));
  CNRT_CHECK(cnrtMalloc((void**)&mlu_src1, num * sizeof(float)));
  CNRT_CHECK(cnrtMalloc((void**)&globalMax, sizeof(float)));
  CNRT_CHECK(cnrtMalloc((void**)&globalSum, sizeof(float)));

  CNRT_CHECK(cnrtMemcpy(mlu_src1, host_src1, num * sizeof(float), cnrtMemcpyHostToDev));
  
  //----------------------------
  CNRT_CHECK(cnrtPlaceNotifier(start, queue));
  softmaxKernel<<<dim, ktype, queue>>>(mlu_dst, mlu_src1, globalMax, globalSum, num);
  CNRT_CHECK(cnrtPlaceNotifier(end, queue));
  cnrtQueueSync(queue);

  //---------------------------
  CNRT_CHECK(cnrtMemcpy(host_dst, mlu_dst, num * sizeof(float), cnrtMemcpyDevToHost));
  for(int i = 0; i < 10; i++){
    printf("softmax[%d]:%.6e,origin:%.6f\n", i, host_dst[i], host_src1[i]);
  }
  float timeTotal;
  CNRT_CHECK(cnrtNotifierDuration(start, end, &timeTotal));
  printf("Total Time: %.3f ms\n", timeTotal / 1000.0);

  CNRT_CHECK(cnrtQueueDestroy(queue));

  cnrtFree(mlu_dst);
  cnrtFree(mlu_src1);
  cnrtFree(globalMax);
  cnrtFree(globalSum);
  
  free(host_dst);
  free(host_src1);
  

  return 0;
}




指数变换也使用流水

但是如果我们把数据从NRAM写回GDRAM也加入流水,那么就需要同步,我们以两个表格来展示流水模型
计算全局最大值和数值和,对于一个长度为4maxNum的向量来说,循环总数是4+1,其中数据拷贝只存在于i<4的循环,数据计算只存在于i>0的循环,我们以i=1为例,当i=1时,此时数据拷贝的是[maxNum:2maxNum]这部分数据,但是计算的却是[0:maxNum]这部分数据,因此这两个完美避开。

在这里插入图片描述

指数变换,这个过程和上面类似,不过循环总数变成了4+2,GDRAM2NRAM这个过程只存在于i<4,Compute这个过程只存在于0<i<5,NRAM2GDRAM这个过程只存在于i>1.
在这里插入图片描述

因此,我们定义一个长度为3maxNum的NRAM向量src,一开始的计算全局最大值和数值和过程只使用前面2maxNum的空间,后面指数变换的时候才使用全部空间。

#include <bang.h>
#include <bang_device_functions.h>
#define EPS 1e-7
const int NRAM_MAX_SIZE = 1024 * 128;//后续树状求和必须保证NRAM_MAX_SIZE为2的幂次
const int maxNum = NRAM_MAX_SIZE/sizeof(float); //NRAM上最多存储maxNum个float元素
const int warpSize = 32;//__bang_reduce_sum每次从src取128字节数据相加,对应32个float元素,并且0-31的结果保存在索引0,32-63的结果保存在索引1

__nram__ float src[3 * maxNum];//后面GDRAM2NRAM,计算,NRAM2GDRAM三份数据
__nram__ float destSum[maxNum];//后面数值求和
__nram__ float destSumFinal[warpSize];//将destSum规约到destFinal[0]
__nram__ float srcMax[2];

__mlu_entry__ void softmaxKernel(float* dst, float* source, float* globalMax, float* globalSum, int num) {
  
  int remain = num%taskDim;//如果不能整除,则让前部分taskId多处理一个元素
  int stepEasy = (num - remain)/taskDim;
  int stepHard = stepEasy + 1;
  int step = (taskId < remain ? stepHard : stepEasy);//前部分taskId多处理一个元素
  int indStart = (taskId < remain ? taskId * stepHard : remain * stepHard + (taskId - remain) * stepEasy);
  int remainNram = step%maxNum;
  int repeat = step/maxNum;//如果一个task处理元素个数超出NRAM最大内存,则需要for循环
  //maxNum尽量取大一些,免得repeat过大导致求和过程累加过于严重,使得大数吃小数
  source = source + indStart;//设定起始偏移量

  //------------------------------------下面开始计算max
  __nram__ float destOldMax;
  __nram__ float destNewMax;
  __bang_write_zero(destSum, maxNum);
  destNewMax = -INFINITY;//初始化为负无穷
  for(int i = 0; i < repeat + 1; i++){
    if(i < repeat){
      __memcpy_async(src + i%2 * maxNum, source + i * maxNum, NRAM_MAX_SIZE, GDRAM2NRAM);
    }
    if(i > 0){
      __bang_argmax(srcMax, src + (i - 1)%2 * maxNum, maxNum);
      if(destNewMax < srcMax[0]){
        destNewMax = srcMax[0];//更新最大值
      }
      __bang_sub_scalar(src + (i - 1)%2 * maxNum, src + (i - 1)%2 * maxNum, destNewMax, maxNum);//src = src - 最大值
      __bang_active_exp_less_0(src + (i - 1)%2 * maxNum, src + (i - 1)%2 * maxNum, maxNum);//src = exp(src - 最大值)
      if(i > 1){
        __bang_mul_scalar(destSum, destSum, exp(destOldMax - destNewMax), maxNum);
      }
      __bang_add(destSum, destSum, src + (i - 1)%2 * maxNum, maxNum);//destSum = destSum + exp(src - destNewMax)
      destOldMax = destNewMax;
    }
    __sync_all_ipu();
  }
  if(remainNram){
    __bang_write_value(src, 3 * maxNum, -INFINITY);//必须要初始化src全部元素为负无穷
    __memcpy(src, source + repeat * maxNum, remainNram * sizeof(float), GDRAM2NRAM);
    __bang_argmax(srcMax, src, maxNum);//针对taskId处理的这step数据,借助于for循环把信息集中到长度为maxNum的向量src中
    if(destNewMax < srcMax[0]){
      destNewMax = srcMax[0];
    }
    __bang_write_value(src, 3 * maxNum, destNewMax);//必须重新初始化为destNewMax
    __memcpy(src, source + repeat * maxNum, remainNram * sizeof(float), GDRAM2NRAM);//必须再次读取
    __bang_sub_scalar(src, src, destNewMax, maxNum);//后面maxNum-remainNram部分为0
    __bang_active_exp_less_0(src, src, maxNum);//相当于多加了maxNum-remainNram
    if(repeat > 0){
      __bang_mul_scalar(destSum, destSum, exp(destOldMax - destNewMax), maxNum);
    }
    __bang_add(destSum, destSum, src, maxNum);
    destOldMax = destNewMax;
  }//结束以后向量destNewMax保存了source[indSart:indStart+step]这部分数据的全局最大值,destSum保存数值和
  //----------
  __bang_write_zero(destSumFinal, warpSize);//初始化destSumFinal全部元素为0
  int segNum = maxNum / warpSize;//将destSum分成segNum段,每段向量长度为warpSize,分段进行树状求和,segNum要求是2的幂次
  for(int strip = segNum/2; strip > 0; strip = strip / 2){//segNum要求是2的幂次即maxNum必须选取2的幂次
    for(int i = 0; i < strip ; i++){
      __bang_add(destSum + i * warpSize, destSum + i * warpSize, destSum + (i + strip) * warpSize, warpSize);
    } 
  }
  __bang_reduce_sum(destSumFinal, destSum, warpSize);
  destSumFinal[0] = destSumFinal[0] - (maxNum - remainNram);//把上面多加的(maxNum - remainNram)减掉
  
  //----------
  globalMax[0] = -INFINITY;
  globalSum[0] = 0.0;
  __sync_all();
  __bang_atomic_max(&destNewMax, globalMax, &destNewMax, 1);//globalMax[0]必须初始化为负无穷
  destSumFinal[0] = destSumFinal[0] * exp(destOldMax - globalMax[0]);
  //__bang_printf("taskId:%d, step:%d, sum:%.6f\n", taskId, step, destSumFinal[0]);
  __sync_all();
  __bang_atomic_add(destSumFinal, globalSum, destSumFinal, 1);//globalSum[0]必须初始化为0
  
  
  dst = dst + indStart;//设定起始偏移量
  float globalSumInv = 1.0/globalSum[0];
  
  for(int i = 0; i < repeat + 2; i++){
    if(i < repeat){
      __memcpy_async(src + i%3 * maxNum, source + i * maxNum, NRAM_MAX_SIZE, GDRAM2NRAM);
    }
    if(i > 0 && i < repeat + 1){
      __bang_sub_scalar(src + (i - 1)%3 * maxNum, src + (i - 1)%3 * maxNum, globalMax[0], maxNum);//src = src - globalMax[0] 
      __bang_active_exp_less_0(src + (i - 1)%3 * maxNum, src + (i - 1)%3 * maxNum, maxNum);//src = exp(src - globalMax[0])
      __bang_mul_scalar(src + (i - 1)%3 * maxNum, src + (i - 1)%3 * maxNum, globalSumInv, maxNum);
    }
    if(i > 1){
      __memcpy_async(dst + (i - 2) * maxNum, src + (i - 2)%3 * maxNum, NRAM_MAX_SIZE, NRAM2GDRAM);
    }
    __sync_all_ipu();
  }
  
  if(remainNram){ 
    __bang_write_value(src, 3 * maxNum, globalMax[0]);
    __memcpy(src, source + repeat * maxNum, remainNram * sizeof(float), GDRAM2NRAM);
    __bang_sub_scalar(src, src, globalMax[0], maxNum);//src = src - globalMax[0] 
    __bang_active_exp_less_0(src, src, maxNum);//src = exp(src - globalMax[0])
    __bang_mul_scalar(src, src, globalSumInv, maxNum);//倒数和另一个向量逐元素相乘得到除法结果
    __memcpy(dst + repeat * maxNum, src, remainNram * sizeof(float), NRAM2GDRAM);
  }
  
  __bang_printf("taskId:%d,repeat:%d,max:%.6f, sum:%.6f\n",taskId, repeat, globalMax[0], globalSum[0]);
  

}


int main(void)
{
  int num = 1024 * 1024 * 1024;
  //int num = 11;
  cnrtQueue_t queue;
  CNRT_CHECK(cnrtSetDevice(0));
  CNRT_CHECK(cnrtQueueCreate(&queue));

  cnrtDim3_t dim = {4, 1, 1};
  int taskNum = dim.x * dim.y * dim.z;
  cnrtFunctionType_t ktype = CNRT_FUNC_TYPE_UNION1;

  cnrtNotifier_t start, end;
  CNRT_CHECK(cnrtNotifierCreate(&start));
  CNRT_CHECK(cnrtNotifierCreate(&end));

  float* host_dst = (float*)malloc(num * sizeof(float));
  float* host_src = (float*)malloc(num * sizeof(float));
  

  for (int i = 0; i < num; i++) {
    host_src[i] = i%4;
  }

  float* mlu_dst;
  float* mlu_src;
  float* globalMax;
  float* globalSum;
  CNRT_CHECK(cnrtMalloc((void**)&mlu_dst, num * sizeof(float)));
  CNRT_CHECK(cnrtMalloc((void**)&mlu_src, num * sizeof(float)));
  CNRT_CHECK(cnrtMalloc((void**)&globalMax, sizeof(float)));
  CNRT_CHECK(cnrtMalloc((void**)&globalSum, sizeof(float)));

  CNRT_CHECK(cnrtMemcpy(mlu_src, host_src, num * sizeof(float), cnrtMemcpyHostToDev));
  
  //----------------------------
  CNRT_CHECK(cnrtPlaceNotifier(start, queue));
  softmaxKernel<<<dim, ktype, queue>>>(mlu_dst, mlu_src, globalMax, globalSum, num);
  CNRT_CHECK(cnrtPlaceNotifier(end, queue));
  cnrtQueueSync(queue);

  //---------------------------
  CNRT_CHECK(cnrtMemcpy(host_dst, mlu_dst, num * sizeof(float), cnrtMemcpyDevToHost));
  for(int i = 0; i < 10; i++){
    printf("softmax[%d]:%.6e,origin:%.6f\n", i, host_dst[i], host_src[i]);
  }
  float timeTotal;
  CNRT_CHECK(cnrtNotifierDuration(start, end, &timeTotal));
  printf("Total Time: %.3f ms\n", timeTotal / 1000.0);

  CNRT_CHECK(cnrtQueueDestroy(queue));

  cnrtFree(mlu_dst);
  cnrtFree(mlu_src);
  cnrtFree(globalMax);
  cnrtFree(globalSum);
  
  free(host_dst);
  free(host_src);
  

  return 0;
}




不能整除部分加入流水

我们在上面的基础上继续针对计算全局最大值,数值和以及指数变换建模,此时我们需要判断一下数组的长度是否整除,如果remain不等于0,那么计算模型的循环总数就设置为repeat+2

在这里插入图片描述

下面是指数变换的表格模型,如果remain不等于0,那么计算模型的循环总数就设置为repeat+3
在这里插入图片描述

这种做法代码实现相对复杂,这里就不提供代码了。

SRAM应用

上面这几个版本的softmax都有一个共同点,即不同taskId处理的那step数据是相互隔离的,类似于taskId=0的时候处理[0:step],taskId = j的时候处理[j×step:(j + 1)×step]这部分数据,对于taskId来说再对这部分数据不断分成maxNum段来处理。这里我们换一种思路,我们把数据分成以长度为taskDim×maxNum的多个小单元,然后针对小单元让不同taskId来处理对应的这部分数据,比如说对于第一个小单元,taskId=j的时候就处理[j×maxNum:(j + 1)×maxNum]这部分数据,对于第N个小单元,taskId=j的时候就处理[(N - 1)×taskDim×maxNum + j×maxNum: N×taskDim×maxNum + (j + 1)×maxNum]这部分数据,在一定程度上可以起到合并访存的效果,最重要的是,此时不同taskId处理的数据是连续的,我们可以使用共享内存SRAM来加速,参考链接添加链接描述
我们开辟一个长度为taskDim×maxNum的共享内存src2SRAM,每次先从source1当中读取一个小单元,把数据先存储到共享内存中,然后做一个cluster上的同步,之后再将这部分数据从共享内存中读取到NRAM上,经过代码测试,这样做可以有效提高速度,完整代码如下所示:

#include <bang.h>
#include <bang_device_functions.h>
#define EPS 1e-7
const int NRAM_MAX_SIZE = 1024 * 256;//后续树状求和必须保证NRAM_MAX_SIZE为2的幂次
const int maxNum = NRAM_MAX_SIZE/sizeof(float); //NRAM上最多存储maxNum个float元素
const int warpSize = 32;//__bang_reduce_sum每次从src取128字节数据相加,对应32个float元素,并且0-31的结果保存在索引0,32-63的结果保存在索引1

__nram__ float src1[maxNum];//每次搬运maxNum数据到NRAM
__nram__ float destSum[maxNum];//后面数值求和
__nram__ float destSumFinal[warpSize];//将destSum规约到destFinal[0]
__nram__ float srcMax[2];

template<int taskNum>
__mlu_entry__ void softmaxKernel(float* dst, float* source1, float* globalMax, float* globalSum, int num) {
  __mlu_shared__ float src2SRAM[taskNum * maxNum];
  int size = taskDim * maxNum;
  int remain = num%size;//如果不能整除,则让前部分taskId多处理一个元素
  int repeat = (num - remain)/size;

  int remainTask = remain%taskDim;
  int stepEasy = (remain - remainTask)/taskDim;
  int stepHard = stepEasy + 1;
  int step = (taskId < remainTask ? stepHard : stepEasy);//前部分taskId多处理一个元素
  int indStart = (taskId < remainTask ? taskId * stepHard : remainTask * stepHard + (taskId - remainTask) * stepEasy);

  __nram__ float destOldMax;
  __nram__ float destNewMax;
  __bang_write_zero(destSum, maxNum);
  destNewMax = -INFINITY;//初始化为负无穷
  for(int i = 0; i < repeat; i++){
    __memcpy(src2SRAM, source1 + i * size, taskDim * NRAM_MAX_SIZE, GDRAM2SRAM);
    __sync_cluster();   //设置sync barrier
    __memcpy(src1, src2SRAM + taskId * maxNum, NRAM_MAX_SIZE, SRAM2NRAM);
    __bang_argmax(srcMax, src1, maxNum);
    if(destNewMax < srcMax[0]){
      destNewMax = srcMax[0];//更新最大值
    }
    __bang_sub_scalar(src1, src1, destNewMax, maxNum);//src1 = src1 - 最大值
    __bang_active_exp_less_0(src1, src1, maxNum);//src1 = exp(src1 - 最大值)
    if(i > 0){
      __bang_mul_scalar(destSum, destSum, exp(destOldMax - destNewMax), maxNum);
    }
    __bang_add(destSum, destSum, src1, maxNum);//destSum = destSum + exp(src1 - destNewMax)
    destOldMax = destNewMax;
  }
  if(remain){

    __bang_write_value(src1, maxNum, -INFINITY);//必须要初始化src1全部元素为负无穷
    __memcpy(src1, source1 + repeat * size + indStart, step * sizeof(float), GDRAM2NRAM);
    __bang_argmax(srcMax, src1, maxNum);
    if(destNewMax < srcMax[0]){
      destNewMax = srcMax[0];
    }
    __bang_write_value(src1, maxNum, destNewMax);//必须重新初始化为destNewMax
    __memcpy(src1, source1 + repeat * size + indStart, step * sizeof(float), GDRAM2NRAM);//必须再次读取
    __bang_sub_scalar(src1, src1, destNewMax, maxNum);//后面maxNum-step部分为0
    __bang_active_exp_less_0(src1, src1, maxNum);//相当于多加了maxNum-step
    if(repeat > 0){
      __bang_mul_scalar(destSum, destSum, exp(destOldMax - destNewMax), maxNum);
    }
    __bang_add(destSum, destSum, src1, maxNum);
    destOldMax = destNewMax;
  }
  
  //----------
  __bang_write_zero(destSumFinal, warpSize);//初始化destSumFinal全部元素为0
  int segNum = maxNum / warpSize;//将destSum分成segNum段,每段向量长度为warpSize,分段进行树状求和,segNum要求是2的幂次
  for(int strip = segNum/2; strip > 0; strip = strip / 2){//segNum要求是2的幂次即maxNum必须选取2的幂次
    for(int i = 0; i < strip ; i++){
      __bang_add(destSum + i * warpSize, destSum + i * warpSize, destSum + (i + strip) * warpSize, warpSize);
    } 
  }
  __bang_reduce_sum(destSumFinal, destSum, warpSize);
  destSumFinal[0] = destSumFinal[0] - (maxNum - step);//把上面多加的(maxNum - step)减掉
  
  //----------
  globalMax[0] = -INFINITY;
  globalSum[0] = 0.0;
  __sync_all();
  __bang_atomic_max(&destNewMax, globalMax, &destNewMax, 1);//globalMax[0]必须初始化为负无穷
  destSumFinal[0] = destSumFinal[0] * exp(destOldMax - globalMax[0]);
  //__bang_printf("taskId:%d, step:%d, sum:%.6f\n", taskId, step, destSumFinal[0]);
  __sync_all();
  __bang_atomic_add(destSumFinal, globalSum, destSumFinal, 1);//globalSum[0]必须初始化为0
  
  
  float globalSumInv = 1.0/globalSum[0];
  
  for(int i = 0; i < repeat; i++){
    __memcpy(src2SRAM, source1 + i * size, taskDim * NRAM_MAX_SIZE, GDRAM2SRAM);
    __sync_cluster();   //设置sync barrier
    __memcpy(src1, src2SRAM + taskId * maxNum, NRAM_MAX_SIZE, SRAM2NRAM);
    __bang_sub_scalar(src1, src1, globalMax[0], maxNum);//src1 = src1 - globalMax[0] 
    __bang_active_exp_less_0(src1, src1, maxNum);//src1 = exp(src1 - globalMax[0])
    __bang_mul_scalar(src1, src1, globalSumInv, maxNum);//倒数和另一个向量逐元素相乘得到除法结果
    __memcpy(dst + i * size + taskId * maxNum, src1, NRAM_MAX_SIZE, NRAM2GDRAM);
  }
  
  if(remain){ 
    __bang_write_value(src1, maxNum, globalMax[0]);
    __memcpy(src1, source1 + repeat * size + indStart, step * sizeof(float), GDRAM2NRAM);
    __bang_sub_scalar(src1, src1, globalMax[0], maxNum);//src1 = src1 - globalMax[0] 
    __bang_active_exp_less_0(src1, src1, maxNum);//src1 = exp(src1 - globalMax[0])
    __bang_mul_scalar(src1, src1, globalSumInv, maxNum);//倒数和另一个向量逐元素相乘得到除法结果
    __memcpy(dst + repeat * size + indStart, src1, step * sizeof(float), NRAM2GDRAM);
  }
  
  __bang_printf("taskId:%d,repeat:%d,max:%.6f, sum:%.6f\n",taskId, repeat, globalMax[0], globalSum[0]);
  

}


int main(void)
{
  int num = 1024 * 1024 * 1024;
  //int num = 11;
  cnrtQueue_t queue;
  CNRT_CHECK(cnrtSetDevice(0));
  CNRT_CHECK(cnrtQueueCreate(&queue));

  cnrtDim3_t dim = {4, 1, 1};
  int taskNum = dim.x * dim.y * dim.z;
  cnrtFunctionType_t ktype = CNRT_FUNC_TYPE_UNION1;

  cnrtNotifier_t start, end;
  CNRT_CHECK(cnrtNotifierCreate(&start));
  CNRT_CHECK(cnrtNotifierCreate(&end));

  float* host_dst = (float*)malloc(num * sizeof(float));
  float* host_src1 = (float*)malloc(num * sizeof(float));
  

  for (int i = 0; i < num; i++) {
    host_src1[i] = i%4;
  }

  float* mlu_dst;
  float* mlu_src1;
  float* globalMax;
  float* globalSum;
  CNRT_CHECK(cnrtMalloc((void**)&mlu_dst, num * sizeof(float)));
  CNRT_CHECK(cnrtMalloc((void**)&mlu_src1, num * sizeof(float)));
  CNRT_CHECK(cnrtMalloc((void**)&globalMax, sizeof(float)));
  CNRT_CHECK(cnrtMalloc((void**)&globalSum, sizeof(float)));

  CNRT_CHECK(cnrtMemcpy(mlu_src1, host_src1, num * sizeof(float), cnrtMemcpyHostToDev));
  
  //----------------------------
  CNRT_CHECK(cnrtPlaceNotifier(start, queue));
  softmaxKernel<4><<<dim, ktype, queue>>>(mlu_dst, mlu_src1, globalMax, globalSum, num);
  CNRT_CHECK(cnrtPlaceNotifier(end, queue));
  cnrtQueueSync(queue);

  //---------------------------
  CNRT_CHECK(cnrtMemcpy(host_dst, mlu_dst, num * sizeof(float), cnrtMemcpyDevToHost));
  for(int i = 0; i < 10; i++){
    printf("softmax[%d]:%.6e,origin:%.6f\n", i, host_dst[i], host_src1[i]);
  }
  float timeTotal;
  CNRT_CHECK(cnrtNotifierDuration(start, end, &timeTotal));
  printf("Total Time: %.3f ms\n", timeTotal / 1000.0);

  CNRT_CHECK(cnrtQueueDestroy(queue));

  cnrtFree(mlu_dst);
  cnrtFree(mlu_src1);
  cnrtFree(globalMax);
  cnrtFree(globalSum);
  
  free(host_dst);
  free(host_src1);
  

  return 0;
}

四级流水

在这里插入图片描述

#include <bang.h>
#include <bang_device_functions.h>
#define EPS 1e-7
const int NRAM_MAX_SIZE = 1024 * 128;//后续树状求和必须保证NRAM_MAX_SIZE为2的幂次
const int maxNum = NRAM_MAX_SIZE/sizeof(float); //NRAM上最多存储maxNum个float元素
const int warpSize = 32;//__bang_reduce_sum每次从src取128字节数据相加,对应32个float元素,并且0-31的结果保存在索引0,32-63的结果保存在索引1

__nram__ float src1[4 * maxNum];//每次搬运maxNum数据到NRAM
__nram__ float destSum[maxNum];//后面数值求和
__nram__ float destSumFinal[warpSize];//将destSum规约到destFinal[0]
__nram__ float srcMax[2];

template<int taskNum>
__mlu_entry__ void softmaxKernel(float* dst, float* source1, float* globalMax, float* globalSum, int num) {
  __mlu_shared__ float src2SRAM[4 * taskNum * maxNum];
  int size = taskDim * maxNum;
  int remain = num%size;//如果不能整除,则让前部分taskId多处理一个元素
  int repeat = (num - remain)/size;

  int remainTask = remain%taskDim;
  int stepEasy = (remain - remainTask)/taskDim;
  int stepHard = stepEasy + 1;
  int step = (taskId < remainTask ? stepHard : stepEasy);//前部分taskId多处理一个元素
  int indStart = (taskId < remainTask ? taskId * stepHard : remainTask * stepHard + (taskId - remainTask) * stepEasy);

  __nram__ float destOldMax;
  __nram__ float destNewMax;
  __bang_write_zero(destSum, maxNum);
  destNewMax = -INFINITY;//初始化为负无穷
  for(int i = 0; i < repeat + 2; i++){
    if(i < repeat){
      __memcpy_async(src2SRAM + i%4 * size, source1 + i * size, taskDim * NRAM_MAX_SIZE, GDRAM2SRAM);
      __sync_cluster();   //i=0才需要设置sync barrier
    }
    if(i > 0 && i < repeat + 1){
      __memcpy_async(src1 + (i - 1)%4 * maxNum, src2SRAM + (i - 1)%4 * size + taskId * maxNum, NRAM_MAX_SIZE, SRAM2NRAM);
    }
    if(i > 1){
      __bang_argmax(srcMax, src1 + (i - 2)%4 * maxNum, maxNum);
      if(destNewMax < srcMax[0]){
        destNewMax = srcMax[0];//更新最大值
      }
      __bang_sub_scalar(src1 +  + (i - 2)%4 * maxNum, src1 +  + (i - 2)%4 * maxNum, destNewMax, maxNum);//src1 = src1 - 最大值
      __bang_active_exp_less_0(src1 +  + (i - 2)%4 * maxNum, src1 + (i - 2)%4 * maxNum, maxNum);//src1 = exp(src1 - 最大值)
      if(i > 2){
        __bang_mul_scalar(destSum, destSum, exp(destOldMax - destNewMax), maxNum);
      }
      __bang_add(destSum, destSum, src1 + (i - 2)%4 * maxNum, maxNum);//destSum = destSum + exp(src1 - destNewMax)
      destOldMax = destNewMax;
    }
    __sync_all_ipu();
  }
  if(remain){

    __bang_write_value(src1, 3 * maxNum, -INFINITY);//必须要初始化src1全部元素为负无穷
    __memcpy(src1, source1 + repeat * size + indStart, step * sizeof(float), GDRAM2NRAM);
    __bang_argmax(srcMax, src1, maxNum);
    if(destNewMax < srcMax[0]){
      destNewMax = srcMax[0];
    }
    __bang_write_value(src1, 3* maxNum, destNewMax);//必须重新初始化为destNewMax
    __memcpy(src1, source1 + repeat * size + indStart, step * sizeof(float), GDRAM2NRAM);//必须再次读取
    __bang_sub_scalar(src1, src1, destNewMax, maxNum);//后面maxNum-step部分为0
    __bang_active_exp_less_0(src1, src1, maxNum);//相当于多加了maxNum-step
    if(repeat > 0){
      __bang_mul_scalar(destSum, destSum, exp(destOldMax - destNewMax), maxNum);
    }
    __bang_add(destSum, destSum, src1, maxNum);
    destOldMax = destNewMax;
  }
  
  //----------
  __bang_write_zero(destSumFinal, warpSize);//初始化destSumFinal全部元素为0
  int segNum = maxNum / warpSize;//将destSum分成segNum段,每段向量长度为warpSize,分段进行树状求和,segNum要求是2的幂次
  for(int strip = segNum/2; strip > 0; strip = strip / 2){//segNum要求是2的幂次即maxNum必须选取2的幂次
    for(int i = 0; i < strip ; i++){
      __bang_add(destSum + i * warpSize, destSum + i * warpSize, destSum + (i + strip) * warpSize, warpSize);
    } 
  }
  __bang_reduce_sum(destSumFinal, destSum, warpSize);
  destSumFinal[0] = destSumFinal[0] - (maxNum - step);//把上面多加的(maxNum - step)减掉
  
  //----------
  globalMax[0] = -INFINITY;
  globalSum[0] = 0.0;
  __sync_all();
  __bang_atomic_max(&destNewMax, globalMax, &destNewMax, 1);//globalMax[0]必须初始化为负无穷
  destSumFinal[0] = destSumFinal[0] * exp(destOldMax - globalMax[0]);
  //__bang_printf("taskId:%d, step:%d, sum:%.6f\n", taskId, step, destSumFinal[0]);
  __sync_all();
  __bang_atomic_add(destSumFinal, globalSum, destSumFinal, 1);//globalSum[0]必须初始化为0
  
  
  float globalSumInv = 1.0/globalSum[0];
  
  for(int i = 0; i < repeat + 3; i++){
    if(i < repeat){
      __memcpy_async(src2SRAM + i%4 * size, source1 + i * size, taskDim * NRAM_MAX_SIZE, GDRAM2SRAM);
      __sync_cluster();   //i=0才需要设置sync barrier
    }
    if(i > 0 && i < repeat + 1){
      __memcpy_async(src1 + (i - 1)%4 * maxNum, src2SRAM + (i - 1)%4 * size + taskId * maxNum, NRAM_MAX_SIZE, SRAM2NRAM);
    }
    if(i > 1 && i < repeat + 2){
      __bang_sub_scalar(src1 + (i - 2)%4 * maxNum, src1 + (i - 2)%4 * maxNum, globalMax[0], maxNum);//src1 = src1 - globalMax[0] 
      __bang_active_exp_less_0(src1 + (i - 2)%4 * maxNum, src1 + (i - 2)%4 * maxNum, maxNum);//src1 = exp(src1 - globalMax[0])
      __bang_mul_scalar(src1 + (i - 2)%4 * maxNum, src1 + (i - 2)%4 * maxNum, globalSumInv, maxNum);//倒数和另一个向量逐元素相乘得到除法结果
    }
    if(i > 2){
      __memcpy_async(dst + (i - 3) * size + taskId * maxNum, src1 + (i - 3)%4 * maxNum, NRAM_MAX_SIZE, NRAM2GDRAM);
    }
    __sync_all_ipu();
  }
  
  if(remain){ 
    __bang_write_value(src1, maxNum, globalMax[0]);
    __memcpy(src1, source1 + repeat * size + indStart, step * sizeof(float), GDRAM2NRAM);
    __bang_sub_scalar(src1, src1, globalMax[0], maxNum);//src1 = src1 - globalMax[0] 
    __bang_active_exp_less_0(src1, src1, maxNum);//src1 = exp(src1 - globalMax[0])
    __bang_mul_scalar(src1, src1, globalSumInv, maxNum);//倒数和另一个向量逐元素相乘得到除法结果
    __memcpy(dst + repeat * size + indStart, src1, step * sizeof(float), NRAM2GDRAM);
  }
  
  __bang_printf("taskId:%d,repeat:%d,max:%.6f, sum:%.6f\n",taskId, repeat, globalMax[0], globalSum[0]);
  

}


int main(void)
{
  int num = 1024 * 1024 * 1024;
  //int num = 11;
  cnrtQueue_t queue;
  CNRT_CHECK(cnrtSetDevice(0));
  CNRT_CHECK(cnrtQueueCreate(&queue));

  cnrtDim3_t dim = {4, 1, 1};
  int taskNum = dim.x * dim.y * dim.z;
  cnrtFunctionType_t ktype = CNRT_FUNC_TYPE_UNION1;

  cnrtNotifier_t start, end;
  CNRT_CHECK(cnrtNotifierCreate(&start));
  CNRT_CHECK(cnrtNotifierCreate(&end));

  float* host_dst = (float*)malloc(num * sizeof(float));
  float* host_src1 = (float*)malloc(num * sizeof(float));
  

  for (int i = 0; i < num; i++) {
    host_src1[i] = i%4;
  }

  float* mlu_dst;
  float* mlu_src1;
  float* globalMax;
  float* globalSum;
  CNRT_CHECK(cnrtMalloc((void**)&mlu_dst, num * sizeof(float)));
  CNRT_CHECK(cnrtMalloc((void**)&mlu_src1, num * sizeof(float)));
  CNRT_CHECK(cnrtMalloc((void**)&globalMax, sizeof(float)));
  CNRT_CHECK(cnrtMalloc((void**)&globalSum, sizeof(float)));

  CNRT_CHECK(cnrtMemcpy(mlu_src1, host_src1, num * sizeof(float), cnrtMemcpyHostToDev));
  
  //----------------------------
  CNRT_CHECK(cnrtPlaceNotifier(start, queue));
  softmaxKernel<4><<<dim, ktype, queue>>>(mlu_dst, mlu_src1, globalMax, globalSum, num);
  CNRT_CHECK(cnrtPlaceNotifier(end, queue));
  cnrtQueueSync(queue);

  //---------------------------
  CNRT_CHECK(cnrtMemcpy(host_dst, mlu_dst, num * sizeof(float), cnrtMemcpyDevToHost));
  for(int i = 0; i < 10; i++){
    printf("softmax[%d]:%.6e,origin:%.6f\n", i, host_dst[i], host_src1[i]);
  }
  float timeTotal;
  CNRT_CHECK(cnrtNotifierDuration(start, end, &timeTotal));
  printf("Total Time: %.3f ms\n", timeTotal / 1000.0);

  CNRT_CHECK(cnrtQueueDestroy(queue));

  cnrtFree(mlu_dst);
  cnrtFree(mlu_src1);
  cnrtFree(globalMax);
  cnrtFree(globalSum);
  
  free(host_dst);
  free(host_src1);
  

  return 0;
}

五级流水

上面这个是四级流水,下面这个是五级流水,但是五级流水的运行时间相比于四级流水更长。
在这里插入图片描述

#include <bang.h>
#include <bang_device_functions.h>
#define EPS 1e-7
const int NRAM_MAX_SIZE = 1024 * 64;//后续树状求和必须保证NRAM_MAX_SIZE为2的幂次
const int maxNum = NRAM_MAX_SIZE/sizeof(float); //NRAM上最多存储maxNum个float元素
const int warpSize = 32;//__bang_reduce_sum每次从src取128字节数据相加,对应32个float元素,并且0-31的结果保存在索引0,32-63的结果保存在索引1

__nram__ float src1[5 * maxNum];//每次搬运maxNum数据到NRAM
__nram__ float destSum[maxNum];//后面数值求和
__nram__ float destSumFinal[warpSize];//将destSum规约到destFinal[0]
__nram__ float srcMax[2];

template<int taskNum>
__mlu_entry__ void softmaxKernel(float* dst, float* source1, float* globalMax, float* globalSum, int num) {
  __mlu_shared__ float src2SRAM[5 * taskNum * maxNum];
  int size = taskDim * maxNum;
  int remain = num%size;//如果不能整除,则让前部分taskId多处理一个元素
  int repeat = (num - remain)/size;

  int remainTask = remain%taskDim;
  int stepEasy = (remain - remainTask)/taskDim;
  int stepHard = stepEasy + 1;
  int step = (taskId < remainTask ? stepHard : stepEasy);//前部分taskId多处理一个元素
  int indStart = (taskId < remainTask ? taskId * stepHard : remainTask * stepHard + (taskId - remainTask) * stepEasy);

  __nram__ float destOldMax;
  __nram__ float destNewMax;
  __bang_write_zero(destSum, maxNum);
  destNewMax = -INFINITY;//初始化为负无穷
  for(int i = 0; i < repeat + 2; i++){
    if(i < repeat){
      __memcpy_async(src2SRAM + i%5 * size, source1 + i * size, taskDim * NRAM_MAX_SIZE, GDRAM2SRAM);
      __sync_cluster();   //设置sync barrier
    }
    if(i > 0 && i < repeat + 1){
      __memcpy_async(src1 + (i - 1)%5 * maxNum, src2SRAM + (i - 1)%5 * size + taskId * maxNum, NRAM_MAX_SIZE, SRAM2NRAM);
    }
    if(i > 1){
      __bang_argmax(srcMax, src1 + (i - 2)%5 * maxNum, maxNum);
      if(destNewMax < srcMax[0]){
        destNewMax = srcMax[0];//更新最大值
      }
      __bang_sub_scalar(src1 +  + (i - 2)%5 * maxNum, src1 +  + (i - 2)%5 * maxNum, destNewMax, maxNum);//src1 = src1 - 最大值
      __bang_active_exp_less_0(src1 +  + (i - 2)%5 * maxNum, src1 + (i - 2)%5 * maxNum, maxNum);//src1 = exp(src1 - 最大值)
      if(i > 2){
        __bang_mul_scalar(destSum, destSum, exp(destOldMax - destNewMax), maxNum);
      }
      __bang_add(destSum, destSum, src1 + (i - 2)%5 * maxNum, maxNum);//destSum = destSum + exp(src1 - destNewMax)
      destOldMax = destNewMax;
    }
    __sync_all_ipu();
  }
  if(remain){

    __bang_write_value(src1, 3 * maxNum, -INFINITY);//必须要初始化src1全部元素为负无穷
    __memcpy(src1, source1 + repeat * size + indStart, step * sizeof(float), GDRAM2NRAM);
    __bang_argmax(srcMax, src1, maxNum);
    if(destNewMax < srcMax[0]){
      destNewMax = srcMax[0];
    }
    __bang_write_value(src1, 3* maxNum, destNewMax);//必须重新初始化为destNewMax
    __memcpy(src1, source1 + repeat * size + indStart, step * sizeof(float), GDRAM2NRAM);//必须再次读取
    __bang_sub_scalar(src1, src1, destNewMax, maxNum);//后面maxNum-step部分为0
    __bang_active_exp_less_0(src1, src1, maxNum);//相当于多加了maxNum-step
    if(repeat > 0){
      __bang_mul_scalar(destSum, destSum, exp(destOldMax - destNewMax), maxNum);
    }
    __bang_add(destSum, destSum, src1, maxNum);
    destOldMax = destNewMax;
  }
  
  //----------
  __bang_write_zero(destSumFinal, warpSize);//初始化destSumFinal全部元素为0
  int segNum = maxNum / warpSize;//将destSum分成segNum段,每段向量长度为warpSize,分段进行树状求和,segNum要求是2的幂次
  for(int strip = segNum/2; strip > 0; strip = strip / 2){//segNum要求是2的幂次即maxNum必须选取2的幂次
    for(int i = 0; i < strip ; i++){
      __bang_add(destSum + i * warpSize, destSum + i * warpSize, destSum + (i + strip) * warpSize, warpSize);
    } 
  }
  __bang_reduce_sum(destSumFinal, destSum, warpSize);
  destSumFinal[0] = destSumFinal[0] - (maxNum - step);//把上面多加的(maxNum - step)减掉
  
  //----------
  globalMax[0] = -INFINITY;
  globalSum[0] = 0.0;
  __sync_all();
  __bang_atomic_max(&destNewMax, globalMax, &destNewMax, 1);//globalMax[0]必须初始化为负无穷
  destSumFinal[0] = destSumFinal[0] * exp(destOldMax - globalMax[0]);
  //__bang_printf("taskId:%d, step:%d, sum:%.6f\n", taskId, step, destSumFinal[0]);
  __sync_all();
  __bang_atomic_add(destSumFinal, globalSum, destSumFinal, 1);//globalSum[0]必须初始化为0
  
  
  float globalSumInv = 1.0/globalSum[0];
  
  for(int i = 0; i < repeat + 4; i++){
    if(i < repeat){
      __memcpy_async(src2SRAM + i%5 * size, source1 + i * size, taskDim * NRAM_MAX_SIZE, GDRAM2SRAM);
      __sync_cluster();   //设置sync barrier
    }
    if(i > 0 && i < repeat + 1){
      __memcpy_async(src1 + (i - 1)%5 * maxNum, src2SRAM + (i - 1)%5 * size + taskId * maxNum, NRAM_MAX_SIZE, SRAM2NRAM);
    }
    if(i > 1 && i < repeat + 2){
      __bang_sub_scalar(src1 + (i - 2)%5 * maxNum, src1 + (i - 2)%5 * maxNum, globalMax[0], maxNum);//src1 = src1 - globalMax[0] 
      __bang_active_exp_less_0(src1 + (i - 2)%5 * maxNum, src1 + (i - 2)%5 * maxNum, maxNum);//src1 = exp(src1 - globalMax[0])
      __bang_mul_scalar(src1 + (i - 2)%5 * maxNum, src1 + (i - 2)%5 * maxNum, globalSumInv, maxNum);//倒数和另一个向量逐元素相乘得到除法结果
    }
    if(i > 2 && i < repeat + 3){
      __memcpy_async(src2SRAM + (i - 3)%5 * size + taskId * maxNum, src1 + (i - 3)%5 * maxNum, NRAM_MAX_SIZE, NRAM2SRAM);
    }
    if(i > 3){
      __memcpy_async(dst + (i - 4) * size, src2SRAM + (i - 4)%5 * size, taskDim * NRAM_MAX_SIZE, SRAM2GDRAM);
    }
    __sync_all_ipu();
  }
  
  if(remain){ 
    __bang_write_value(src1, maxNum, globalMax[0]);
    __memcpy(src1, source1 + repeat * size + indStart, step * sizeof(float), GDRAM2NRAM);
    __bang_sub_scalar(src1, src1, globalMax[0], maxNum);//src1 = src1 - globalMax[0] 
    __bang_active_exp_less_0(src1, src1, maxNum);//src1 = exp(src1 - globalMax[0])
    __bang_mul_scalar(src1, src1, globalSumInv, maxNum);//倒数和另一个向量逐元素相乘得到除法结果
    __memcpy(dst + repeat * size + indStart, src1, step * sizeof(float), NRAM2GDRAM);
  }
  
  __bang_printf("taskId:%d,repeat:%d,max:%.6f, sum:%.6f\n",taskId, repeat, globalMax[0], globalSum[0]);
  

}


int main(void)
{
  int num = 1024 * 1024 * 1024;
  //int num = 11;
  cnrtQueue_t queue;
  CNRT_CHECK(cnrtSetDevice(0));
  CNRT_CHECK(cnrtQueueCreate(&queue));

  cnrtDim3_t dim = {4, 1, 1};
  int taskNum = dim.x * dim.y * dim.z;
  cnrtFunctionType_t ktype = CNRT_FUNC_TYPE_UNION1;

  cnrtNotifier_t start, end;
  CNRT_CHECK(cnrtNotifierCreate(&start));
  CNRT_CHECK(cnrtNotifierCreate(&end));

  float* host_dst = (float*)malloc(num * sizeof(float));
  float* host_src1 = (float*)malloc(num * sizeof(float));
  

  for (int i = 0; i < num; i++) {
    host_src1[i] = i%4;
  }

  float* mlu_dst;
  float* mlu_src1;
  float* globalMax;
  float* globalSum;
  CNRT_CHECK(cnrtMalloc((void**)&mlu_dst, num * sizeof(float)));
  CNRT_CHECK(cnrtMalloc((void**)&mlu_src1, num * sizeof(float)));
  CNRT_CHECK(cnrtMalloc((void**)&globalMax, sizeof(float)));
  CNRT_CHECK(cnrtMalloc((void**)&globalSum, sizeof(float)));

  CNRT_CHECK(cnrtMemcpy(mlu_src1, host_src1, num * sizeof(float), cnrtMemcpyHostToDev));
  
  //----------------------------
  CNRT_CHECK(cnrtPlaceNotifier(start, queue));
  softmaxKernel<4><<<dim, ktype, queue>>>(mlu_dst, mlu_src1, globalMax, globalSum, num);
  CNRT_CHECK(cnrtPlaceNotifier(end, queue));
  cnrtQueueSync(queue);

  //---------------------------
  CNRT_CHECK(cnrtMemcpy(host_dst, mlu_dst, num * sizeof(float), cnrtMemcpyDevToHost));
  for(int i = 0; i < 10; i++){
    printf("softmax[%d]:%.6e,origin:%.6f\n", i, host_dst[i], host_src1[i]);
  }
  float timeTotal;
  CNRT_CHECK(cnrtNotifierDuration(start, end, &timeTotal));
  printf("Total Time: %.3f ms\n", timeTotal / 1000.0);

  CNRT_CHECK(cnrtQueueDestroy(queue));

  cnrtFree(mlu_dst);
  cnrtFree(mlu_src1);
  cnrtFree(globalMax);
  cnrtFree(globalSum);
  
  free(host_dst);
  free(host_src1);
  

  return 0;
}

在这里插入图片描述

高维向量的softmax

高维向量的softmax实现更加复杂,回忆之前在英伟达平台上实现高维向量的softmax函数,比如说我们以形状为[1,2,3,4,5,6]的6维向量举例,变换维度假设axis=2,之前英伟达平台的实现,我们计算出变换维度的长度dimsize=3,其他维度的乘积othersize=1×2×4×5×6 = 240,步长stride= 1×6×5×4 = 120,使用othersize=240个线程块,其中每个线程块处理对应一份数据,计算出int tid =blockIdx.x % stride + (blockIdx.x - blockIdx.x % stride) × dimsize;全局索引为tid + threadIdx.x × stride,类似地,我们也按照这个思路来实现寒武纪显卡上的高维向量softmax:
我们利用taskId来处理othersize,但是考虑到taskDim往往是2或者4的倍数,而othersize不一定满足这个条件,因此我们使用for循环来解决,参考for(int otherIdx = taskId; otherIdx < othersize; otherIdx += taskDim)
进入上述for循环以后,我们尝试来处理dimsize,由于寒武纪的函数基本上支持向量操作,无法针对具体某个元素来处理,为此我们仍然把dimsize这份数据按照maxNum长度分成多个小单元,如果不能整除后面特殊处理,特殊处理的方式和上面一维向量一模一样。在代码24行——25行,这里使用两层for循环来加载数据,高维数组导致每次处理的数据不连续,间隔stride,为此必须要不断遍历数组把结果集中到src1数组上处理,后续的处理类似,这里不做赘述。

#include <bang.h>
#include <bang_device_functions.h>
#define EPS 1e-7
const int NRAM_MAX_SIZE = 1024 * 4;//后续树状求和必须保证NRAM_MAX_SIZE为2的幂次
const int maxNum = NRAM_MAX_SIZE/sizeof(float); //NRAM上最多存储maxNum个float元素
const int warpSize = 32;//__bang_reduce_sum每次从src取128字节数据相加,对应32个float元素,并且0-31的结果保存在索引0,32-63的结果保存在索引1

__nram__ float src1[maxNum];//每次搬运maxNum数据到NRAM
__nram__ float destSum[maxNum];//后面数值求和
__nram__ float destSumFinal[warpSize];//将destSum规约到destFinal[0]
__nram__ float srcMax[2];

__mlu_entry__ void softmaxKernel(float* dst, float* source1, int othersize, int dimsize, int stride) {
  int remain = dimsize%maxNum;
  int repeat = (dimsize - remain)/maxNum;
  __nram__ float destOldMax;
  __nram__ float destNewMax;
  //下面利用taskId来处理其他维度
  for(int otherIdx = taskId; otherIdx < othersize; otherIdx += taskDim){
    destOldMax = -INFINITY;
    destNewMax = -INFINITY;
    __bang_write_zero(destSum, maxNum);
    int tid = otherIdx % stride + (otherIdx - otherIdx % stride) * dimsize;
    for(int i = 0; i < repeat; i++){
      for(int j = 0; j < maxNum; j++){//从source1间隔stride读取数据
        __memcpy(src1 + j, source1 + tid + (i * maxNum + j) * stride, sizeof(float), GDRAM2NRAM);
      }
      __bang_argmax(srcMax, src1, maxNum);
      if(destNewMax < srcMax[0]){
        destNewMax = srcMax[0];//更新最大值
      }
      __bang_sub_scalar(src1, src1, destNewMax, maxNum);//src1 = src1 - 最大值
      __bang_active_exp_less_0(src1, src1, maxNum);//src1 = exp(src1 - 最大值)
      if(i > 0){
        __bang_mul_scalar(destSum, destSum, exp(destOldMax - destNewMax), maxNum);
      }
      __bang_add(destSum, destSum, src1, maxNum);//destSum = destSum + exp(src1 - destNewMax)
      destOldMax = destNewMax;
    }
    //-------------------------------------
    if(remain){
      __bang_write_value(src1, maxNum, -INFINITY);//多余部分必须设置负无穷
      for(int j = 0; j < remain; j++){
        __memcpy(src1 + j, source1 + tid + (repeat * maxNum + j) * stride, sizeof(float), GDRAM2NRAM);
      }
      __bang_argmax(srcMax, src1, maxNum);
      if(destNewMax < srcMax[0]){
        destNewMax = srcMax[0];
      }
      __bang_write_value(src1, maxNum, destNewMax);//必须重新初始化为destNewMax
      for(int j = 0; j < remain; j++){
        __memcpy(src1 + j, source1 + tid + (repeat * maxNum + j) * stride, sizeof(float), GDRAM2NRAM);
      }
      __bang_sub_scalar(src1, src1, destNewMax, maxNum);//后面maxNum-remain部分为0
      __bang_active_exp_less_0(src1, src1, maxNum);//相当于多加了maxNum-remain
      if(repeat > 0){
        __bang_mul_scalar(destSum, destSum, exp(destOldMax - destNewMax), maxNum);
      }
      __bang_add(destSum, destSum, src1, maxNum);
      destOldMax = destNewMax;
    }
    
    //--------------------------------
    __bang_write_zero(destSumFinal, warpSize);
    int segNum = maxNum / warpSize;
    for(int strip = segNum/2; strip > 0; strip = strip / 2){
      for(int i = 0; i < strip ; i++){
        __bang_add(destSum + i * warpSize, destSum + i * warpSize, destSum + (i + strip) * warpSize, warpSize);
      } 
    }
    __bang_reduce_sum(destSumFinal, destSum, warpSize);
    
    destSumFinal[0] = destSumFinal[0] - (maxNum - remain);
    //__bang_printf("--max:%.3e,sum:%.6e,:%d\n",destNewMax,destSumFinal[0], maxNum - remain);
    //------------------------------------至此全局最大值为destNewMax,全局数值和为destSumFinal[0]
    float globalSumInv = 1.0/destSumFinal[0];
    for(int i = 0; i < repeat; i++){
      for(int j = 0; j < maxNum; j++){
        __memcpy(src1 + j, source1 + tid + (i * maxNum + j) * stride, sizeof(float), GDRAM2NRAM);
      }
      __bang_sub_scalar(src1, src1, destNewMax, maxNum); 
      __bang_active_exp_less_0(src1, src1, maxNum);
      __bang_mul_scalar(src1, src1, globalSumInv, maxNum);//倒数和另一个向量逐元素相乘得到除法结果
      for(int j = 0; j < maxNum; j++){
        __memcpy(dst + tid + (i * maxNum + j) * stride, src1 + j, sizeof(float), NRAM2GDRAM);
      }
    }
    if(remain){
      __bang_write_value(src1, maxNum, destNewMax);
      for(int j = 0; j < remain; j++){
        __memcpy(src1 + j, source1 + tid + (repeat * maxNum + j) * stride, sizeof(float), GDRAM2NRAM);
      }
      __bang_sub_scalar(src1, src1, destNewMax, maxNum);
      __bang_active_exp_less_0(src1, src1, maxNum);
      __bang_mul_scalar(src1, src1, globalSumInv, maxNum);//倒数和另一个向量逐元素相乘得到除法结果
      for(int j = 0; j < remain; j++){
        __memcpy(dst + tid + (repeat * maxNum + j) * stride, src1 + j, sizeof(float), NRAM2GDRAM);
      }
    }
    
  }
  
  
}


int main(void)
{
  int num = 32 * 16 * 64 * 128;//shape = {32, 16, 64, 128},axis = 2
  int stride = 128;
  int dimsize = 64;
  int othersize = 32 * 16 * 128;
  /***
  int num = 24;//shape = {2,3,2,2}, axis = 1
  int stride = 4;
  int dimsize = 3;
  int othersize = 8;
  ***/
  cnrtQueue_t queue;
  CNRT_CHECK(cnrtSetDevice(0));
  CNRT_CHECK(cnrtQueueCreate(&queue));

  cnrtDim3_t dim = {4, 1, 1};
  int taskNum = dim.x * dim.y * dim.z;
  cnrtFunctionType_t ktype = CNRT_FUNC_TYPE_UNION1;

  cnrtNotifier_t start, end;
  CNRT_CHECK(cnrtNotifierCreate(&start));
  CNRT_CHECK(cnrtNotifierCreate(&end));

  float* host_dst = (float*)malloc(num * sizeof(float));
  float* host_src1 = (float*)malloc(num * sizeof(float));
  

  for (int i = 0; i < num; i++) {
    host_src1[i] = i%4;
    //host_src1[i] = i;
  }

  float* mlu_dst;
  float* mlu_src1;
  
  CNRT_CHECK(cnrtMalloc((void**)&mlu_dst, num * sizeof(float)));
  CNRT_CHECK(cnrtMalloc((void**)&mlu_src1, num * sizeof(float)));
  

  CNRT_CHECK(cnrtMemcpy(mlu_src1, host_src1, num * sizeof(float), cnrtMemcpyHostToDev));
  
  //----------------------------
  CNRT_CHECK(cnrtPlaceNotifier(start, queue));
  softmaxKernel<<<dim, ktype, queue>>>(mlu_dst, mlu_src1, othersize, dimsize, stride);
  CNRT_CHECK(cnrtPlaceNotifier(end, queue));
  cnrtQueueSync(queue);

  //---------------------------
  CNRT_CHECK(cnrtMemcpy(host_dst, mlu_dst, num * sizeof(float), cnrtMemcpyDevToHost));
  for(int i = 0; i < 24; i++){
    printf("softmax[%d]:%.6e,origin:%.6f\n", i, host_dst[i], host_src1[i]);
  }
  float timeTotal;
  CNRT_CHECK(cnrtNotifierDuration(start, end, &timeTotal));
  printf("Total Time: %.3f ms\n", timeTotal / 1000.0);

  CNRT_CHECK(cnrtQueueDestroy(queue));

  cnrtFree(mlu_dst);
  cnrtFree(mlu_src1);
  
  
  free(host_dst);
  free(host_src1);
  

  return 0;
}
                           

高维向量的pingpong流水

此时流水无法加速

#include <bang.h>
#include <bang_device_functions.h>
#define EPS 1e-7
const int NRAM_MAX_SIZE = 1024 * 4;//后续树状求和必须保证NRAM_MAX_SIZE为2的幂次
const int maxNum = NRAM_MAX_SIZE/sizeof(float); //NRAM上最多存储maxNum个float元素
const int warpSize = 32;//__bang_reduce_sum每次从src取128字节数据相加,对应32个float元素,并且0-31的结果保存在索引0,32-63的结果保存在索引1

__nram__ float src1[3 * maxNum];//每次搬运maxNum数据到NRAM
__nram__ float destSum[maxNum];//后面数值求和
__nram__ float destSumFinal[warpSize];//将destSum规约到destFinal[0]
__nram__ float srcMax[2];

__mlu_entry__ void softmaxKernel(float* dst, float* source1, int othersize, int dimsize, int stride) {
  int remain = dimsize%maxNum;
  int repeat = (dimsize - remain)/maxNum;
  __nram__ float destOldMax;
  __nram__ float destNewMax;
  //下面利用taskId来处理其他维度
  for(int otherIdx = taskId; otherIdx < othersize; otherIdx += taskDim){
    destOldMax = -INFINITY;
    destNewMax = -INFINITY;
    __bang_write_zero(destSum, maxNum);
    int tid = otherIdx % stride + (otherIdx - otherIdx % stride) * dimsize;
    for(int i = 0; i < repeat + 1; i++){
      if(i < repeat){
        for(int j = 0; j < maxNum; j++){//从source1间隔stride读取数据
          __memcpy_async(src1 + i%2 * maxNum + j, source1 + tid + (i * maxNum + j) * stride, sizeof(float), GDRAM2NRAM);
        }
      }
      if(i > 0){
        __bang_argmax(srcMax, src1 + (i - 1)%2 * maxNum, maxNum);
        if(destNewMax < srcMax[0]){
          destNewMax = srcMax[0];//更新最大值
        }
        __bang_sub_scalar(src1 + (i - 1)%2 * maxNum, src1 + (i - 1)%2 * maxNum, destNewMax, maxNum);//src1 = src1 - 最大值
        __bang_active_exp_less_0(src1 + (i - 1)%2 * maxNum, src1 + (i - 1)%2 * maxNum, maxNum);//src1 = exp(src1 - 最大值)
        if(i > 1){
          __bang_mul_scalar(destSum, destSum, exp(destOldMax - destNewMax), maxNum);
        }
        __bang_add(destSum, destSum, src1 + (i - 1)%2 * maxNum, maxNum);//destSum = destSum + exp(src1 - destNewMax)
        destOldMax = destNewMax;
      }
      //__sync_all_ipu();
    }
    //-------------------------------------
    if(remain){
      __bang_write_value(src1, maxNum, -INFINITY);//多余部分必须设置负无穷
      for(int j = 0; j < remain; j++){
        __memcpy(src1 + j, source1 + tid + (repeat * maxNum + j) * stride, sizeof(float), GDRAM2NRAM);
      }
      __bang_argmax(srcMax, src1, maxNum);
      if(destNewMax < srcMax[0]){
        destNewMax = srcMax[0];
      }
      __bang_write_value(src1, maxNum, destNewMax);//必须重新初始化为destNewMax
      for(int j = 0; j < remain; j++){
        __memcpy(src1 + j, source1 + tid + (repeat * maxNum + j) * stride, sizeof(float), GDRAM2NRAM);
      }
      __bang_sub_scalar(src1, src1, destNewMax, maxNum);//后面maxNum-remain部分为0
      __bang_active_exp_less_0(src1, src1, maxNum);//相当于多加了maxNum-remain
      if(repeat > 0){
        __bang_mul_scalar(destSum, destSum, exp(destOldMax - destNewMax), maxNum);
      }
      __bang_add(destSum, destSum, src1, maxNum);
      destOldMax = destNewMax;
    }
    
    //--------------------------------
    __bang_write_zero(destSumFinal, warpSize);
    int segNum = maxNum / warpSize;
    for(int strip = segNum/2; strip > 0; strip = strip / 2){
      for(int i = 0; i < strip ; i++){
        __bang_add(destSum + i * warpSize, destSum + i * warpSize, destSum + (i + strip) * warpSize, warpSize);
      } 
    }
    __bang_reduce_sum(destSumFinal, destSum, warpSize);
    
    destSumFinal[0] = destSumFinal[0] - (maxNum - remain);
    //__bang_printf("--max:%.3e,sum:%.6e,:%d\n",destNewMax,destSumFinal[0], maxNum - remain);
    //------------------------------------至此全局最大值为destNewMax,全局数值和为destSumFinal[0]
    float globalSumInv = 1.0/destSumFinal[0];
    for(int i = 0; i < repeat + 2; i++){
      if(i < repeat){
        for(int j = 0; j < maxNum; j++){
          __memcpy_async(src1 + i%3 * maxNum + j, source1 + tid + (i * maxNum + j) * stride, sizeof(float), GDRAM2NRAM);
        }
      }
      if(i > 0){
        __bang_sub_scalar(src1 + (i - 1)%3 * maxNum, src1 + (i - 1)%3 * maxNum, destNewMax, maxNum); 
        __bang_active_exp_less_0(src1 + (i - 1)%3 * maxNum, src1 + (i - 1)%3 * maxNum, maxNum);
        __bang_mul_scalar(src1 + (i - 1)%3 * maxNum, src1 + (i - 1)%3 * maxNum, globalSumInv, maxNum);//倒数和另一个向量逐元素相乘得到除法结果
      }
      if(i > 1){
        for(int j = 0; j < maxNum; j++){
            __memcpy_async(dst + tid + ((i - 2) * maxNum + j) * stride, src1 + (i - 2)%3 * maxNum + j, sizeof(float), NRAM2GDRAM);
        }
      }
      //__sync_all_ipu();
    }
    if(remain){
      __bang_write_value(src1, maxNum, destNewMax);
      for(int j = 0; j < remain; j++){
        __memcpy(src1 + j, source1 + tid + (repeat * maxNum + j) * stride, sizeof(float), GDRAM2NRAM);
      }
      __bang_sub_scalar(src1, src1, destNewMax, maxNum);
      __bang_active_exp_less_0(src1, src1, maxNum);
      __bang_mul_scalar(src1, src1, globalSumInv, maxNum);//倒数和另一个向量逐元素相乘得到除法结果
      for(int j = 0; j < remain; j++){
        __memcpy(dst + tid + (repeat * maxNum + j) * stride, src1 + j, sizeof(float), NRAM2GDRAM);
      }
    }
    
  }
  
  
}


int main(void)
{
  int num = 32 * 16 * 64 * 128;//shape = {32, 16, 64, 128},axis = 2
  int stride = 128;
  int dimsize = 64;
  int othersize = 32 * 16 * 128;
  /***
  int num = 24;//shape = {2,3,2,2}, axis = 1
  int stride = 4;
  int dimsize = 3;
  int othersize = 8;
  ***/
  cnrtQueue_t queue;
  CNRT_CHECK(cnrtSetDevice(0));
  CNRT_CHECK(cnrtQueueCreate(&queue));

  cnrtDim3_t dim = {4, 1, 1};
  int taskNum = dim.x * dim.y * dim.z;
  cnrtFunctionType_t ktype = CNRT_FUNC_TYPE_UNION1;

  cnrtNotifier_t start, end;
  CNRT_CHECK(cnrtNotifierCreate(&start));
  CNRT_CHECK(cnrtNotifierCreate(&end));

  float* host_dst = (float*)malloc(num * sizeof(float));
  float* host_src1 = (float*)malloc(num * sizeof(float));
  

  for (int i = 0; i < num; i++) {
    host_src1[i] = i%4;
    //host_src1[i] = i;
  }

  float* mlu_dst;
  float* mlu_src1;
  
  CNRT_CHECK(cnrtMalloc((void**)&mlu_dst, num * sizeof(float)));
  CNRT_CHECK(cnrtMalloc((void**)&mlu_src1, num * sizeof(float)));
  

  CNRT_CHECK(cnrtMemcpy(mlu_src1, host_src1, num * sizeof(float), cnrtMemcpyHostToDev));
  
  //----------------------------
  CNRT_CHECK(cnrtPlaceNotifier(start, queue));
  softmaxKernel<<<dim, ktype, queue>>>(mlu_dst, mlu_src1, othersize, dimsize, stride);
  CNRT_CHECK(cnrtPlaceNotifier(end, queue));
  cnrtQueueSync(queue);

  //---------------------------
  CNRT_CHECK(cnrtMemcpy(host_dst, mlu_dst, num * sizeof(float), cnrtMemcpyDevToHost));
  for(int i = 0; i < 24; i++){
    printf("softmax[%d]:%.6e,origin:%.6f\n", i, host_dst[i], host_src1[i]);
  }
  float timeTotal;
  CNRT_CHECK(cnrtNotifierDuration(start, end, &timeTotal));
  printf("Total Time: %.3f ms\n", timeTotal / 1000.0);

  CNRT_CHECK(cnrtQueueDestroy(queue));

  cnrtFree(mlu_dst);
  cnrtFree(mlu_src1);
  
  
  free(host_dst);
  free(host_src1);
  

  return 0;
}
                           

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处:/a/331364.html

如若内容造成侵权/违法违规/事实不符,请联系我们进行投诉反馈qq邮箱809451989@qq.com,一经查实,立即删除!

相关文章

STM32标准库开发——USART串口外设

USART外设介绍 USART (Universal Synchronous/AsynchronousReceiver/Transmitter&#xff09;通用同步/异步收发器USART是STM32内部集成的硬件外设&#xff0c;可根据数据寄存器的一个字节数据自动生成数据帧时序&#xff0c;从TX引脚发送出去&#xff0c;也可自动接收RX引脚的…

WebGL中开发AR应用

WebGL在本质上是用于在浏览器中进行3D和2D图形渲染的技术&#xff0c;而增强现实&#xff08;AR&#xff09;通常需要与现实世界的环境进行交互。要在WebGL中开发AR应用&#xff0c;您可以采取以下步骤&#xff0c;希望对大家有所帮助。北京木奇移动技术有限公司&#xff0c;专…

Arm Generic Interrupt Controller v3 and v4(GICv3v4)学习(一)

提示 该博客主要为个人学习&#xff0c;通过阅读官网手册整理而来&#xff08;个人觉得阅读官网的英文文档非常有助于理解各个IP特性&#xff09;。若有不对之处请参考参考文档&#xff0c;以官网参考文档为准。 Arm Generic Interrupt Controller v3 and v4学习一共分为三章&…

RHEL8 Samba服务器详细配置用户模式

任务&#xff1a; 配置server01为samba服务器&#xff0c;samba服务器的/companydata/sales为共享目录&#xff0c;共享名为sales&#xff0c;里面创建测试文件test_share.tar&#xff0c;创建用户组sales&#xff0c;创建组内用户sale1&#xff0c;要求配置用户模式访问&#…

Uniapp多选Popup(弹出层)

uniapp中多选组件很少&#xff0c;故个人简单开发了一个&#xff0c;可简单使用&#xff0c;也可根据个人需求稍微改进 支持的功能 单选多选&#xff08;默认&#xff09;限制选择数量默认选中禁用选项 属性说明 属性默认值说明singlefalsetrue为开启单选&#xff0c;否则为…

无需信用卡注册美区Apple ID指南

第一步 准备工作 1、一个没有注册过AppleID的邮箱&#xff0c;建议最好是Gmail邮箱 2、一个苹果手机&#xff0c;当然这个是必须的 3、需要科学上网 第二步 苹果网站注册 为了避免cookie的干扰&#xff0c;最好是在无痕模式下打开以上网页&#xff0c;创建你的AppleID&#…

rabbitmq-java基础详解

一、rabbitmq是什么&#xff1f; 1、MQ定义 MQ&#xff08;Message Queue&#xff09;消息队列 主要解决&#xff1a;异步处理、应用解耦、流量削峰等问题&#xff0c;是分布式系统的重要组件&#xff0c;从而实现高性能&#xff0c;高可用&#xff0c;可伸缩和最终一致性的架…

Spring+SpringMVC+Mybatis进行项目的整合

Spring SpringMVCM Mybatis 整合 一、 通过idea创建maven工程 二、 引入依赖项以及导入mybatis逆向工程的插件 将如下的文件替换所在工程的pom文件 <?xml version"1.0" encoding"UTF-8"?><project xmlns"http://maven.apache.org/POM/4…

HCIA的访问控制列表ACL

ACL -----access control-list 允许/拒绝 ACL作用&#xff1a; 1.实现访问控制 2.定义感兴趣流量 ACL分类&#xff1a; 标准ACL 2000-2999&#xff08;只关注源IP地址&#xff0c;使用时应该尽量靠近目标&#xff09; 扩展ACL 3000-3999&#xff1a;写ACL不能写在源上&…

反射计数 - 华为OD统一考试

OD统一考试 分值&#xff1a; 200分 题解&#xff1a; Java / Python / C 题目描述 给定一个包含 0 和 1 的二维矩阵, 给定一个初始位置和速度。 一个物体从给定的初始位置触发, 在给定的速度下进行移动, 遇到矩阵的边缘则发生镜面反射无论物体经过 0 还是 1&#xff0c;都不…

2024美赛数学建模思路 - 案例:异常检测

文章目录 赛题思路一、简介 -- 关于异常检测异常检测监督学习 二、异常检测算法2. 箱线图分析3. 基于距离/密度4. 基于划分思想 建模资料 赛题思路 &#xff08;赛题出来以后第一时间在CSDN分享&#xff09; https://blog.csdn.net/dc_sinor?typeblog 一、简介 – 关于异常…

宠物空气净化器真的有用吗?五款猫用宠物空气净化器测评!

作为一个养猫四年的铲屎官&#xff0c;我不得不说&#xff0c;宠物空气净化器是21世纪养猫人最伟大的神器之一&#xff01; 当我刚开始养猫的时候&#xff0c;我并没有意识到猫毛会成为一个如此头疼的问题。虽然朋友们告诉我要做好心理准备&#xff0c;但我并没有想到家里的猫毛…

Apache Zeppelin学习记录2

Apache Zeppelin学习记录2 文章目录 Apache Zeppelin学习记录2前言一、基础调用二、带参数调用1.代码块要增加一行z.textbox("folder_path", "input")2.读取result 总结 前言 上一章讲了如何使用zeppelin来接入python&#xff0c;本节我们来看看如何使用R…

ArcGIS初始化软件界面Normal.mxt

ArcGIS有时候永久了&#xff0c;或者呢突然不自觉软件界面乱了&#xff0c;或者一些窗口打开却找不到&#xff01; 这时候可以去删除arcgis的界面配置文件&#xff0c;Normal.mxt 删除后再打开软件&#xff0c;软件界面就会回到初始化设置了&#xff01; 文件所在的路径&…

3d音响按键怎么建立模型---模大狮模型网

要建立3D音响按键的模型&#xff0c;您可以按照以下步骤进行&#xff1a; 选择建模软件&#xff1a;首先&#xff0c;选择一个三维建模软件&#xff0c;如Blender、3ds Max或Maya。这些软件都提供了丰富的建模工具和功能&#xff0c;适合用于创建复杂的三维模型。 参考图像&am…

IPv6自动隧道---ISATAP隧道

ISATAP隧道 ISATAP(Intra-Site Automatic Tunnel Addressing Protocol)是另外一种自动隧道技术。ISATAP隧道同样使用了内嵌IPv4地址的特殊IPv6地址形式,只是和6to4不同的是,6to4是使用IPv4地址做为网络前缀,而ISATAP用IPv4地址做为接口标识。 站点内自动隧道寻址协议(I…

Web Animation API

工作中经常会遇到需要动画的场景&#xff0c;连贯动画都是用CSS实现&#xff0c;&#xff0c;但是如果遇到需要用户互动介入的动画&#xff0c;那纯CSS很比较吃力&#xff0c;也不是不能实现&#xff0c;需要动态修改CSS变量&#xff0c;而且动画容易被JS代码阻塞&#xff0c;导…

VMP比较正确的编译教程

一、编译环境 1.1 编译整体配置 采用VS2022社区版MSVC2017_xpQT5.6.0WDK7.1&#xff08;编译DDK需要&#xff0c;不需要DDK的可以不用下载&#xff09; 1.1 VS2022安装 1.1.1 除常规勾选桌面C以外&#xff0c;需要勾选win xp支持和支持相应的MSVC版本。教程采用msvc2017和x…

​批量文件夹随机小写字母重命名:文件夹重命名简单步骤,高效结果

在日常工作中&#xff0c;经常要对大量的文件夹重命名进行管理和查找文件。手动重命名每个文件夹不仅耗时&#xff0c;而且容易出错。现在一起来看云炫文件管理器如何给文件夹名称批量随机小写字母重命名&#xff0c;简单的步骤&#xff0c;高效率的结果。 文件夹名称随机小写…

多文件转二维码的两种方式,有兴趣的了解一下

多个文件能一键生成二维码吗&#xff1f;二维码是现在很多人用来展示文件内容的一种手段&#xff0c;在制作二维码图片之后&#xff0c;其他人扫码就可以查看文件或者下载文件&#xff0c;有效的提升文件获取的效率。一般情况下&#xff0c;文件二维码分为多个文件生成一个二维…