写在前面

可能很多人已经对CUDA编程比较熟悉，而对于BANG C编程还有些陌生。在学习BANG C的过程中我不由得思考：如果我们已经拥有CUDA版本的算子，要如何以BANG C实现相同功能的算子呢？由此我计划通过一系列连载的形式来讨论这个问题，希望可以给大家一些帮助。

BANG C编程基础

我会对同一个算子以不同的实现方法来介绍 CUDA 算子转 BANG C 算子开发过程的大致流程。 就以最简单的两个张量相加的算子为例，说明BANG C开发的基本流程。

CUDA编程

首先在 CUDA 编程中，我们将需要处理的数据拷贝到GPU内存上，即可直接调用kernel函数。Kernel函数执行完成后， 将结果拷出到CPU内存上即可。

• kernel code：

  __global__ void VecAdd(float* A, float* B, float* C, int N){
      int i = blockDim.x * blockIdx.x + threSadIdx.x;
      if (i < size) C[i] = A[i] + B[i];}

• Host code：

  int main(){
      ...
      float* h_A = (float*)malloc(size);
      float* h_B = (float*)malloc(size);
  
      // Allocate vectors in device memory
      float* d_A;
      cudaMalloc(&d_A, size);
      float* d_B;
      cudaMalloc(&d_B, size);
      float* d_C;
      cudaMalloc(&d_C, size);
  
      // Copy vectors from host memory to device memory
      cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
      cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
  
      VecAdd<<>>(d_A, d_B, d_C, size);
  
      // Copy result from device memory to host memory
      // h_C contains the result in host memory
      cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
      ...}

BANG C kernel编程

• 基础编程

  注：由于部分BANG C的内置算子不支持 float32 类型的输入，以及出于对获得最佳性能的需要， 个人一般推荐使用 half 或 int8 型数据作为BANG C kernel的输入，如果需要更高精度的结果且使用的BANG C算子支持 float32类型，可以尝试修改kernel的数据类型。

  对于BANG C来说，在编写kernel函数时，因为MLU的硬件结构与GPU的不同，BANG C并不是通过将计算分成多个thread执行来获得性能提升的。 因此一般情况下，BANG C kernel的写法与C++类似，如下所示：

  __mlu_entry__ void VecAdd(half* A, half* B, half* C, int N) {
     for (int i = 0; i < N; i++) {
         C[i] = A[i] + B[i];
     }}

• 使用BANGC算子

  为了充分利用MLU的特性来获得最大的性能提升，个人强烈建议使用BANG C提供的内置算子来实现函数。比如上面的向量相加， 可以利用BANG C算子改写如下。需要注意的是，在使用BANG C内置算子时，传入的指针参数所指向的内存所在位置均有限制， 因此需要按照限制将数据先转移到相应的位置再使用内置算子。同时，内置算子对每次调用时计算的数据量有对齐的要求 （即为某个数值的整数倍），所以在调用时还需要考虑对齐。

  // 设定最大nram size
  #define NRAM_SIZE 1024 * 472
  #define NRAM_HALF_SIZE NRAM_SIZE / sizeof(half)
  // 设定对齐函数
  #define ALIGIN_UP(x, num) (((x - 1) / num + 1) * num)
  #define ALIGN_DOWN_TO(x, num) (((x - 1) / num) * num)
  __mlu_entry__ void VecAdd(half* A, half* B, half* C, int N) {
      // nram 存储空间不支持类似malloc的操作，只能通过定义数组的形式来获取
      __nram__ half nram_buffer[NRAM_HALF_SIZE];
  
      // __bang_add 算子要求计算的数据总量 size 需要满足 (size * sizeof(type)) % 128 = 0, 同时考虑到nram大小的局限，
      // 需要对输入的数据进行拆分后再进行计算
  
      // 获取A, B, C 三部分每次进行计算的最大数据量
      int data_size = ALIGN_DOWN_TO(NRAM_HALF_SIZE / 3, 64);
      half* N_A = &nram_buffer[0];
      half* N_B = N_A + data_size;
      half* N_C = N_B + data_size;
      int loop = N / data_size;
      int remain = N % data_size;
      for (int i = 0; i < loop; ++i){
          // 将数据由gdram拷贝至nram
          __memcpy(N_A, A + i * data_size, data_size * sizeof(half), GDRAM2NRAM);
          __memcpy(N_B, B + i * data_size, data_size * sizeof(half), GDRAM2NRAM);
          __bang_add(N_C, N_A, N_B, data_size);
          // 将结果拷出至gdram
          __memcpy(N_C, C + i * data_size, data_size * sizeof(half), GDRAM2NRAM);
      }
      // 处理余数部分
      if(remain != 0){
          __memcpy(N_A, A + loop * data_size, remain * sizeof(half), GDRAM2NRAM);
          __memcpy(N_B, B + loop * data_size, remain * sizeof(half), GDRAM2NRAM);
          // 即使需要计算的数据量是remain，仍需要计算 data_size 个数据以进行对齐
          __bang_add(N_C, N_A, N_B, data_size);
          // 只拷出剩余所需数据，注意此处若拷出data_size个数据可能会导致出错
          __memcpy(N_C, C + loop * data_size, remain * sizeof(half), GDRAM2NRAM);
      }}// 考虑到__bang_add 支持原位运算，上面的代码可以优化为...// 省去了存储计算结果的空间，增大了可以用来计算的空间// int data_size = ALIGN_DOWN_TO(NRAM_HALF_SIZE / 3, 64);int data_size = ALIGN_DOWN_TO(NRAM_HALF_SIZE / 2, 64);...// 计算结果可以存储在其中一个输入的位置上// __bang_add(N_C, N_A, N_B, data_size);// __memcpy(N_C, C + i * data_size, data_size * sizeof(half), GDRAM2NRAM);__bang_add(N_A, N_A, N_B, data_size);__memcpy(N_A, C + i * data_size, data_size * sizeof(half), NRAM2GDRAM);

• 并行计算

  在MLU设备上，一般都含有多个计算单元，通常情况下可以通过调用更多的计算单元来提高计算效率。 在此我先展示如何在kernel函数内进行并行计算编程，后面在host端代码中会介绍如何使用多个计算单元。

  __mlu_entry__ void VecAdd(half* A, half* B, half* C, int N) {
      __nram__ half nram_buffer[NRAM_HALF_SIZE];
      int data_size = ALIGN_DOWN_TO(NRAM_HALF_SIZE / 2, 64);
      half* N_A = &nram_buffer[0];
      half* N_B = N_A + data_size;
      // taskDim为mlu的内置变量，定义了kernel调用的总任务数，在进行数据分割时需要考虑按照任务总数进行分割
      int loop = N / (data_size * taskDim);
      int remain = N % (data_size * taskDim);
      for (int i = 0; i < loop; ++i){
          // taskId为mlu的内置变量，定义了当前的任务编号，根据任务编号来定位需要拷贝的数据位置
          __memcpy(N_A, A + (i * taskDim + taskId) * data_size, data_size * sizeof(half), GDRAM2NRAM);
          __memcpy(N_B, B + (i * taskDim + taskId) * data_size, data_size * sizeof(half), GDRAM2NRAM);
          __bang_add(N_A, N_A, N_B, data_size);
          __memcpy(C + (i * taskDim + taskId) * data_size, N_A, data_size * sizeof(half), NRAM2GDRAM);
      }
      if(remain != 0){
          int loop2 = remain / data_size;
          if(taskId < loop2){
              __memcpy(N_A, A + (loop * taskDim + taskId) * data_size, data_size * sizeof(half), GDRAM2NRAM);
              __memcpy(N_B, B + (loop * taskDim + taskId) * data_size, data_size * sizeof(half), GDRAM2NRAM);
              __bang_add(N_A, N_A, N_B, data_size);
              __memcpy(C + (loop * taskDim + taskId) * data_size, N_A, data_size * sizeof(half), NRAM2GDRAM);
          }
          int remain2 = remain % data_size;
          if(remain2 != 0 && taskId ==loop2){
              int remain2_up = ALIGIN_UP(remain2, 64);
              __memcpy(N_A, A + (loop * taskDim + taskId) * data_size, remain2 * sizeof(half), GDRAM2NRAM);
              __memcpy(N_B, B + (loop * taskDim + taskId) * data_size, remain2 * sizeof(half), GDRAM2NRAM);
              __bang_add(N_A, N_A, N_B, remain2_up);
              __memcpy(C + (loop * taskDim + taskId) * data_size, N_A, remain2 * sizeof(half), NRAM2GDRAM);
          }
      }}

BANG C host编程

首先需要将数据转为 half 型，再拷贝到MLU上进行计算。

#include typedef uint16_t half;
int main(){

    float* h_A = (float*)malloc(size);
    float* h_B = (float*)malloc(size);
    float* h_C = (float*)malloc(size);

    int size_h = size * sizeof(half) / sizeof(float);
    int N = size_h / sizeof(half);

    // Allocate vectors in half in host memory
    half* h_A_h = (half*)(malloc(size_h));
    half* h_B_h = (half*)(malloc(size_h));

    // Convert input to half
    for (int i = 0; i < N; i++) {
        cnrtConvertFloatToHalf(h_A_h + i, h_A[i]);
        cnrtConvertFloatToHalf(h_B_h + i, h_B[i]);
    }

    // Allocate vectors in device memory
    half* d_A;
    half* d_B;
    half* d_C;
    cnrtMalloc((void**)(&d_A), size_h);
    cnrtMalloc((void**)(&d_B), size_h);
    cnrtMalloc((void**)(&d_C), size_h);

    // Copy vectors from host memory to device memory
    cnrtMemcpy(d_A, h_A_h, size_h, CNRT_MEM_TRANS_DIR_HOST2DEV);
    cnrtMemcpy(d_B, h_B_h, size_h, CNRT_MEM_TRANS_DIR_HOST2DEV);

    cnrtQueue_t queue;
    cnrtCreateQueue(&queue);
    // 设定kernel所用到的任务线程数，dim.x * dim.y * dim.z 即为总任务线程数taskDim
    cnrtDim3_t dim;
    dim.x = 1;
    dim.y = 1;
    dim.z = 1;
    // 设定调用的计算单元数，注意总任务线程数必须为计算单元数的整数倍，具体变量列表可参考CNRT开发手册
    cnrtFunctionType_t ktype = CNRT_FUNC_TYPE_BLOCK;


    //调用BANG C kernel
    VecAdd<<>>(d_A, d_B, d_C, N);
    cnrtSyncQueue(queue);

    // 拷出结果至h_C
    half* h_C_h = (half*)(malloc(size_h));
    cnrtMemcpy(h_C_h, d_C, size_h, CNRT_MEM_TRANS_DIR_DEV2HOST);

    for (int i = 0; i < N; i++) {
        cnrtConvertHalfToFloat(h_C_h + i, h_C[i]);
    }
    // 释放mlu资源
    cnrtFree(d_A);
    cnrtFree(d_B);
    cnrtFree(d_C);
    cnrtDestroyQueue(queue);
    cnrtDestroyKernelParamsBuffer(params);
    //释放cpu资源
    ...}

BANG C 程序编译

BANG C程序需要使用neuware提供的cncc编译器进行编译，并在参数中指定编译的平台，如下为一个简单的示例。

export NEUWARE_HOME=/usr/local/neuware
cncc -std=c++11 --bang-mlu-arch=MLU270 --bang-device-only -c  add_kernel.mlu -o add_kernel.o
g++ -std=c++11 -c add.cc -I ${NEUWARE_HOME}/include -L ${NEUWARE_HOME}/lib64 -lcnrt -lopenblas
g++ -std=c++11 add_kernel.o add.o -o add_test -L ${NEUWARE_HOME}/lib64 -lcnrt -lopenblas