CUDA-3

CUDA-3

上一次已经见到了gridDim这个的意思是每个grid中我们想要用多少个blocks来运行我们的并行程序,而blockDim指的是每个block中我们想要开多少个线程来运行我们的程序。 这一节就准备好好理解一下这些名词。

blockDim

这个是built-in variable, 对于所有的blocks来说都是一个constant, 代表着每个block中的线程数量。 其实blockDim有三个维度, 分别是 blockDim.x,blockDim.y, blockDim.z, 而gridDim有两个维度gridDim.x, gridDim.y,总共加起来有5个维度,已经能够处理大多数的问题了,而且许多问题用不了那么多的线程,只是用了这种方法之后能够提高运行的效率,节约时间。

int tid = threadIdx.x + blockIdx.x*blockDim.x

这就像是把一个维的矩阵拉成一维的向量的时候的操作,其中blockDim.x是个常值,比如核函数中我们给定的如果是<<<N, 10>>> 的话,那么blockDim.x=10, threadIdx.xblockIdx.x的取值都为 0,1,...,9.

就相当于把一个 N×10的矩阵换成一个长为N×10的一维的向量一样。容易看出,按上面的操作最多能开 N×10个线程来算,但是如果向量太大的时候呢,或者向量的长度不能够整除一些数的时候。

不够整除的时候很简单 N-batch/batch +1就可以了,再加上 while(tid<N) 就比较好了。

如果向量实在太大的话,还可以用另外一个办法,我把它称为tile操作。

int tid = threadIdx.x +blockIdx.x* blockDim.x;
while(tid<N){
    c[tid] = a[tid]+b[tid];
    tid += blockDim.x * gridDim.x;
}

就可以了。下面对此加以说明,假设现在N很大,但是由于硬件限制,核函数只能开成 <<<128, 1024>>>即每个grid可以有128个block,每个block的维度可以有1024个线程。

那么 tid在没有经过while之前最大应该是 1023+127*1024=128*1024-1即第一块砖的最后个id,即第一次tid可以取0,1,...,128*1024-1, 那么下一次就可以从

128*1024, 128*1024+1,..., 2*128*1024-1, M次加操作之后最后一个就变成了 (M+1)*128*1024-1了,即是一个shape为 M+1, 128, 1024的三维矩阵了。如果再不够,可以5个维度全部都用上。

下面是一个例子。

#include <stdio.h>

#define N 123456789

//our kernel function

__global__ void add(int *a, int *b, int *c){
    
    int tid = threadIdx.x+blockIdx.x*blockDim.x;  // handle the data at this index
    while(tid<N){
        c[tid] = a[tid] + b[tid];
        tid += blockDim.x*gridDim.x;
    }   

}

int main(void){
    int a[N], b[N], c[N];
    int *dev_a, *dev_b, *dev_c;
    
    //allocate the memory on the GPU
    cudaMalloc((void**)&dev_a, N*sizeof(int));
    cudaMalloc((void**)&dev_b, N*sizeof(int));
    cudaMalloc((void**)&dev_c, N*sizeof(int));
    
    
    // fill the arrays on the cpu

    for(int i=0;i<N;i++){
        a[i] = -i;  
        b[i] = i*i;
    }   

    // copy the arrays to the gpu
    cudaMemcpy(dev_a, a, N*sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, N*sizeof(int), cudaMemcpyHostToDevice);

    // run the kernle
    add<<<1024, 1024>>>(dev_a, dev_b, dev_c);
    

    // copy the result c to cpu
    cudaMemcpy(c, dev_c, N*sizeof(int), cudaMemcpyDeviceToHost);
    
    //free the memory on the device
    cudaFree(dev_a);
    cudaFree(dev_b);
    cudaFree(dev_c);

    return 0;

}

打赏,谢谢~~

取消

感谢您的支持,我会继续努力的!

扫码支持
扫码打赏,多谢支持~

打开微信扫一扫,即可进行扫码打赏哦