参加CUDA线上训练营·（三）cuda的线程层次

CUDA中的线程组织为三个层次Grid、Block、Thread。假设需要对两个256*1的向量进行加法操作，各部分如下设置：

Thread

thread对应的是需要执行的线程数，对应的是GPU中的cuda核。此处需要对256个数进行计算，因此threadIdx = 256。

Block

block一开始比较难以理解，其计算公式如下：

block_size = threads_all[总线程数] / thread_num[每个block含有多少个thread)] （此处向上取整）

因此，假设thread_num设置为128，则block_size = 256 / 128 = 2.

在下图中可以看到，若干个cuda核加上共享内存和其余计算单元，共同组成了一个SM。因此，block其实就是运行在SM上，也与相应的硬件特性相符：

• SM中有共享内存，内部的每个cuda核都可以访问，这就对应了编程中每个block内的thread都可以使用shared memory。但因为这是在SM内的存储空间，所以其空间是十分有限的，且与GPU的架构相关（一般在几十KB左右）。
• SM中的cuda core数量是有限的，目前最多到1024个，且与硬件架构有关。因此设置合适的thread_num数就显得尤为重要。过小会产生大量的block，导致核心没有充分利用，过大则会导致cuda重复计算，并行效果不好。
• 每个thread产生的一般类型的数据（如定长数组等），一般会被放置在寄存器中，以达到最大的访问速度。但如果数据量超过存储限制，则会被放置在local memory甚至global memory中，这会导致计算速度大幅下降。因此，如果thread需要存储大量数据，建议削减thread_num以平衡存储消耗。
  

测试

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <math.h>


void __global__ add(const double *x, const double *y, double *z, int count)
{
    const int n = blockDim.x * blockIdx.x + threadIdx.x;
	if( n < count)
	{
	    z[n] = x[n] + y[n];
	}

}

void check(const double *z, const int N)
{
    bool error = false;
    for (int n = 0; n < N; ++n)
    {
        if (fabs(z[n] - 3) > (1.0e-10))
        {
            error = true;
        }
    }
    printf("%s\n", error ? "Errors" : "Pass");
}


int main(void)
{
    const int N = 1e7;
    const int M = sizeof(double) * N;
    double *h_x = (double*) malloc(M);
    double *h_y = (double*) malloc(M);
    double *h_z = (double*) malloc(M);

    for (int n = 0; n < N; ++n)
    {
        h_x[n] = 1;
        h_y[n] = 2;
    }

    double *d_x, *d_y, *d_z;
    cudaMalloc((void **)&d_x, M);
    cudaMalloc((void **)&d_y, M);
    cudaMalloc((void **)&d_z, M);
    cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);
    cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice);

    const int thread_size = 8;
    const int block_size = (N + thread_size - 1) / thread_size;
    add<<<block_size, thread_size>>>(d_x, d_y, d_z, N);

    //计算时间
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    cudaEventRecord(start, 0);
    float time;



    add<<<block_size, thread_size>>>(d_x, d_y, d_z, N);


    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&time, start, stop);
    printf("gpu time:%.4f\n", time);

    cudaEventDestroy(start);
    cudaEventDestroy(stop);
    //计算时间



    cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost);
    check(h_z, N);

    free(h_x);
    free(h_y);
    free(h_z);
    cudaFree(d_x);
    cudaFree(d_y);
    cudaFree(d_z);
    return 0;
}

在此处分别设置thread_size为8，16，32和64，其运行时间分别为2.46，1.25，0.77，0.78ms ，硬件平台是RTX3060。由此可见，thread_size并不是越大越好，当达到cuda核上限时会达到瓶颈。