cuda实现矩阵相加

Posted 2021-11-26 wangha

cuda实现向量相加

博客最后附上整体代码

如果有说的不对的地方还请前辈指出, 因为cuda真的接触没几天

一些总结(建议看)

1. cuda 并不纯GPU在运行程序, 而是 cpu 与 gpu 一起在运行程序, cpu负责调度, gpu 负责运算, cpu称为HOST , gpu 称为 DEVICE
2. 记住三个东西 grid block thread ,关系分别是 grid 包含多个 block , block 包含多个 thread
3. 一个block中thread个数选取一般为32的整数倍, 原因和warp有关, 有兴趣自行查阅
4. 一个grid中block的个数选取和你的kernel函数以及thread数量有关, 举个例子, int a[1000] 加上 int b[1000] , 你的thread为64, 那么, block = 1000/64 = 16个合适
5. __global__函数一般表示一个内核函数，是一组由GPU执行的并行计算任务,由cpu调用
6. __host__一般是由CPU调用，由CPU执行的函数,
7. __device1__一般表示由GPU中一个线程调用的函数

代码实现

引入

#include <stdio.h>
#include <cuda_runtime.h>

kernel函数

__global__ void
vectorAdd(float *a, float *b, float *c, int num){
        int i = blockDim.x * blockIdx.x + threadIdx.x; //vector is 1-dim, blockDim means the number of thread in a block
        if(i < num){
                c[i] = a[i] + b[i];
        }
}

int i = blockDim.x * blockIdx.x + threadIdx.x;

这句代码解释一下:

blockDim.x 表示block的size行数(如果是一维的block的话,即一行有多少个thread)

blockIdx.x 表示当前运行到的第几个block(一维grid的话,即该grid中第几个block)

threadIdx.x 表示当前运行到的第几个thread (一维的block的话.即该block中第几个thread)

画个图解释一下

比如上面这个图的话, ABCDE各代表一个block, 总的为一个Grid, 每个block中有四个thread, 图中我花了箭头的也就是代表着第1个block中的第0个thread.

那么 i = blockDim.x * blockIdx.x + threadIdx.x 就是指 i = 4 * 1 + 0

申请内存空间与释放

host中申请内存

float *a = (float *)malloc(size);
float *b = (float *)malloc(size);
float *c = (float *)malloc(size);

free(a);
free(b);
free(c);

device中申请内存

float *da = NULL;
float *db = NULL;
float *dc = NULL;

cudaMalloc((void **)&da, size);
cudaMalloc((void **)&db, size);
cudaMalloc((void **)&dc, size);

cudaFree(da);
cudaFree(db);
cudaFree(dc);

host中内存copy到device

cudaMemcpy(da,a,size,cudaMemcpyHostToDevice);
cudaMemcpy(db,b,size,cudaMemcpyHostToDevice);
cudaMemcpy(dc,c,size,cudaMemcpyHostToDevice);

上面的cudaMemcpyHostToDevice用于指定方向有四种关键词

cudaMemcpyHostToDevice | cudaMemcpyHostToHost | cudaMemcpyDeviceToDevice | cudaMemcpyDeviceToHost

启动 kernel函数

int threadPerBlock = 256;                        
int blockPerGrid = (num + threadPerBlock - 1)/threadPerBlock;
vectorAdd <<< blockPerGrid, threadPerBlock >>> (da,db,dc,num)

此处确定了block中的thread数量以及一个grid中block数量

利用kernel function <<< blockPerGrid, threadPerBlock>>> (paras,...) 来实现在cuda中运算

参考

https://zhuanlan.zhihu.com/p/345877391

https://docs.nvidia.com/cuda/cuda-c-programming-guide/

源码展示

#include <stdio.h>

#include <cuda_runtime.h>

// vectorAdd run in device
__global__ void 
vectorAdd(float *a, float *b, float *c, int num){
    int i = blockDim.x * blockIdx.x + threadIdx.x; //vector is 1-dim, blockDim means the number of thread in a block
    if(i < num){
        c[i] = a[i] + b[i];
    }
}

// main run in host
int
main(void){
    int num = 10000; // size of vector
    size_t size = num * sizeof(float);

    // host memery
    float *a = (float *)malloc(size);
    float *b = (float *)malloc(size);
    float *c = (float *)malloc(size);

    // init the vector
    for(int i=1;i<num;++i){
        a[i] = rand()/(float)RAND_MAX;
        b[i] = rand()/(float)RAND_MAX;
    }

    // copy the host memery to device memery
    float *da = NULL;
    float *db = NULL;
    float *dc = NULL;

    cudaMalloc((void **)&da, size);
    cudaMalloc((void **)&db, size);
    cudaMalloc((void **)&dc, size);

    cudaMemcpy(da,a,size,cudaMemcpyHostToDevice);
    cudaMemcpy(db,b,size,cudaMemcpyHostToDevice);
    cudaMemcpy(dc,c,size,cudaMemcpyHostToDevice);

    // launch function add kernel
    int threadPerBlock = 256;
    int blockPerGrid = (num + threadPerBlock - 1)/threadPerBlock;
    printf("threadPerBlock: %d \\nblockPerGrid: %d \\n",threadPerBlock,blockPerGrid);

    vectorAdd <<< blockPerGrid, threadPerBlock >>> (da,db,dc,num);

    //copy the device result to host
    cudaMemcpy(c,dc,size,cudaMemcpyDeviceToHost);

    // Verify that the result vector is correct
    for (int i = 0; i < num; ++i){
        if (fabs(a[i] + b[i] - c[i]) > 1e-5){
            fprintf(stderr, "Result verification failed at element %d!\\n", i);
            return 0;
        }
    }

    printf("Test PASSED\\n");

    // Free device global memory
    cudaFree(da);
    cudaFree(db);
    cudaFree(dc);
    // Free host memory
    free(a);
    free(b);
    free(c);

    printf("free is ok\\n");
    return 0;
}