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核上限时会达到瓶颈。