CUDA(六):多种方法实现向量加

---

博主CUDA学习系列汇总传送门（持续更新）：编程语言|CUDA入门

---

---

向量加法是高性能运算中最简单的运算，
本章节将通过CUDA中的几种加法例子来理解CUDA中thread/block的概念

理论转自https://blog.csdn.net/canhui_wang/article/details/51730264

• Grid，Block和Thread三者的关系
  
  其中，一个grid包含多个blocks，这些blocks的组织方式可以是一维，二维或者三维。任何一个block包含有多个Threads，这些Threads的组织方式也可以是一维，二维或者三维。举例来讲：比如上图中，任何一个block中有10个Thread，那么，Block(0,0)的第一个Thread的ThreadIdx是0，Block(1,0)的第一个Thread的ThreadIdx是10；Block(2,0)的第一个Thread的ThreadIdx是20，…，依此类推，不难整理出其中的映射公式（表达式已在代码中给出）。

原文链接：https://blog.csdn.net/canhui_wang/article/details/51730264

一、单block 单thread相加

GPU端的向量加法与CPU端类似，只是核函数申明需要用__global__限定符标识；
核函数调用时需要<<< X, X>>>配置grid和block维度。

/**====================================================
1、单block单thread向量加
=====================================================**/
__global__ void vector_add_gpu_1(int *a, int *b, int *c, int n)
{
    for(int i=0; i<n; ++i)
    {
        c[i] = a[i] + b[i];
    }
}

二、单block多thread相加

kernel 函数启动时，1个grid拥有多个block， 而一个block又拥有多个thread
针对向量加法，kernel函数中 利用threadIdx.x获取线程索引号，通过<<<1, threadnum>>来指定一个block内threadnum个thread同时运算， 每个线程完成1次向量加法索引号根据线程总数（blockDim.x）进行跳步。

/**====================================================
2、单block多thread向量加
=====================================================**/
__global__ void vector_add_gpu_2(int *a, int *b, int *c, int n)
{
    int tid = threadIdx.x; //线程索引号
    const int t_n = blockDim.x; // 一个block内的线程总数
    while (tid < n)
    {
        c[tid] = a[tid] + b[tid];
        tid += t_n;
    }
}

三、多Block多thread相加

全局线程索引： blockIdx.x * blockDim.x + threadIdx.x;
线程数量（跳步步长）： gridDim.x * blockDim.x

/**====================================================
3、多block多thread向量加
=====================================================**/
__global__ void vector_add_gpu_3(int *a, int *b, int *c, int n)
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x; // 获取线程索引
    const int t_n = gridDim.x * blockDim.x; // 跳步的步长，所有线程的数量
    printf("gridDim.x = %d\n", gridDim.x);
    printf("blockDim.x = %d\n", blockDim.x);
    printf("t_n = %d\n", t_n);
    while (tid < n)
    {
        c[tid] = a[tid] + b[tid];
        tid += t_n;
    }
}

四、主函数

#include <iostream>
#define  N 10
int main() {
    int a[N], b[N], c[N];
    int *dev_a, *dev_b, *dev_c;
    for(int i=0; i<N; ++i) // 为数组a、b赋值
    {
        a[i] = i;
        b[i] = i * i;
    }
    cudaMalloc(&dev_a, sizeof(int) * N);
    cudaMemcpy(dev_a, a, sizeof(int) * N, cudaMemcpyHostToDevice);

    cudaMalloc(&dev_b, sizeof(int) * N);
    cudaMemcpy(dev_b, b, sizeof(int) * N, cudaMemcpyHostToDevice);

    cudaMalloc(&dev_c, sizeof(int) * N);
    cudaMemcpy(dev_c, c, sizeof(int) * N, cudaMemcpyHostToDevice);

//    vector_add_gpu_1<<<1, 1>>>(dev_a, dev_b, dev_c, N);
//    vector_add_gpu_2<<<1, 4>>>(dev_a, dev_b, dev_c, N);
    vector_add_gpu_3<<<2, 4>>>(dev_a, dev_b, dev_c, N);
    cudaMemcpy(c, dev_c, sizeof(int) * N, cudaMemcpyDeviceToHost);

    for(int i=0; i<N; ++i)
    {
        printf("%d + %d = %d \n", a[i], b[i], c[i]);
    }

    cudaFree(dev_a);
    cudaFree(dev_b);
    cudaFree(dev_c);

    std::cout << "Hello, World!" << std::endl;
    return 0;
}