×
分享到微信

打开微信,使用扫一扫进入页面后,点击右上角菜单,

点击“发送给朋友”或“分享到朋友圈”完成分享

CUDA转BANG C示例指南系列连载之一 BANG C编程基础 chap72021-01-13 20:36:36 回复 2 查看 BANG语言
CUDA转BANG C示例指南系列连载之一 BANG C编程基础
分享到:

写在前面

可能很多人已经对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


版权所有 © 2021 寒武纪 Cambricon.com 备案/许可证号:京ICP备17003415号-1
关闭