Cuda编程 之 块(block)计算

#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>
#include <time.h>

#define N 512

__global__ void vec_block_add(int *a, int *b, int *c)
{ 
    c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x];
}

void rand_ints(int *arr, int count)
{
    srand(time(NULL));
    for(int i=0; i<count; i++)
    {
        arr[i] = rand()%100;
    }
}

int main()
{
    int *a, *b, *c;
    int *d_a, *d_b, *d_c;
    int size = N*sizeof(int);
    clock_t start, finish;

    cudaMalloc((void**)&d_a, size);
    cudaMalloc((void**)&d_b, size);
    cudaMalloc((void**)&d_c, size);

    a=(int*)malloc(size);
    rand_ints(a, N);
    b=(int*)malloc(size);
    rand_ints(b, N);
    c=(int*)malloc(size);
    rand_ints(c, N);

    cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice);

    start = clock();
    vec_block_add<<<100, 1>>>(d_a, d_b, d_c);
    finish = clock();
    double time = (finish - start)/(double)CLOCKS_PER_SEC;

    cudaMemcpy(c, d_c, size, cudaMemcpyDeviceToHost);

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

    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);

    free(a);
    free(b);
    free(c);
    return 0;
}

程序说明

程序有两个关键部分：

c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x];

vec_block_add<<<100, 1>>>(d_a, d_b, d_c);

add<<<N,1>>>N表示同时调用N次add函数，这是实现并行的向量相加的语句。每个被并行调用的add函数称之为一个块(block)。

块的集合称之为网格(grid).
每个块可以使用索引值blockIdx.x

函数并行执行的过程，相当于由GPU这个加速器使用硬件的方式进行了循环展开，而blockIdx.x就是用来定位当前执行的是循环的哪个部分。通过使用blockIdx.x作为索引，每个块可以处理数组元素中的一部分。从硬件的角度看，相当于同时有多个块在并行执行：

块0: c[0]=a[0]+b[0]
块1: c[1]=a[1]+b[1]
块2: c[2]=a[2]+b[2]
块3: c[3]=a[3]+b[3]
….

程序中通过时间函数比较了并行计算的块数目不同时的计算时间。

---

本文参考自hipercome博文，向作者致谢！