并行编程实践——CUDA编程的多线程编程

原创已于 2025-09-15 13:40:16 修改 · 1.1k 阅读

23 ·

本内容遵循CC 4.0 BY-SA版权协议

GEO检测

收录于

C++

并行编程

于 2025-09-14 09:00:16 首次发布

一、线程相关的术语

其实在前面已经对CUDA中线程的初步一些知识进行了说明，本文将开始正式分析CUDA中的多线程编程，首先需要对其中几个重要的相关术语进行说明：
Thread：这个很好理解，和普通的计算机中线程类似，它是最小的运算单元，负责Kernel函数中定义的运算
Block：由多个Thread组成的运算单元，每个块只能由GPU中的单个SM执行。同一块内的线程可共享内存并交换数据。单个块最多支持1024个线程
Clusters：即Thread Block Clusters，由线程块组成的线程簇。可以理解成类似于Thread与Block之间抽象关系，它是一个可选层次结构。保证集群中的线程块在GPU中的GPU处理集群上共同调度。CUDA支持集群上最多8个线程块作为可移植集群大小。
Grid：由多个块组成，即可以理解为一个GPU抽象的对所有块的管理层
通过上面的术语描述其实就可以看出来其实在CUDA中对GPU的线程管理是分层进行处理的。除了簇这个概念层是在较高版本（9.0）提出的外，其实主流的就是三层即线程、块到网格。这也符合GPU本身的硬件特点。
然后再说一下Kernel：Kernel是在GPU上运行的函数即在GPU上并行执行的任务。当程序从Host调用Kernel时，它被加载到GPU上启动，在多个线程上并行运行。一个Kernel对应一个Grid（GPU）。一个Grid对应多个Block（或Clusters）即提到的1-3维，一个Block有多个Thread，即1~3维。

二、CUDA的多线程

在CUDA的多线程编程中，有两个重要的参数需要明白，这个在前面的代码中也分析过：threadIdx.x和blockIdx.x。看一个简单的代码：

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

__global__ void gpu_println() {
    printf("gpu thread [%d,%d] \n", threadIdx.x,blockIdx.x);
}

int main(void) {
    printf("host start......!\n");
    gpu_println<<<1,1>>>();
    //gpu_println<<<1,2>>>();
    //gpu_println<<<2,1>>>();
    cudaDeviceSynchronize();
    return 0;
}

在CUDA编程中，可以有多种方式创建多个线程，一种是创建多个块，一种是创建多个线程，还有就是可以混合创建多线程。也就是通过“<<<1,2>>>”来指定，第一个数字表示块的数量，第二个表示线程的数量。那么索引一个线程的ID怎么计算呢：

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

看上面的代码，其实就是块的维度（1，2，3）方向上的索引乘以块的索引，然后再加上线程中的索引，最终形成一个线程的全局索引。这个很好理解，可以理解成在空间定位一个点的度假村。在类似CPU一维线程中，就是一条线，直接加就可以了。二维平面中就需要定位一下x,y两个方向，那么三维呢，这个都是入门的几何知识。
再举一个复杂的例子：

dim3 block(3, 4, 5);
dim3 grid(2, 2, 2);

上述的代码定义一个345=60线程的块，又定义了一个含有222=8个块的网格。可以将它们理解成一个立方体。块定义中，在x方向有3个线程（元素），y方向有4个，z方向有5个。也可以理解成一个三维数组，对应着x,y,z三个方向。Grid的定义也是可以如此理解。在CUDA中的访问定义方式如下：

threadIdx.x, threadIdx.y, threadIdx.z：线程的块内索引
blockIdx.x, blockIdx.y, blockIdx.z：块的网格索引
blockDim.x, blockDim.y, blockDim.z：块的维度
gridDim.x, gridDim.y, gridDim.z：网格的维度

举一个形象的例子，大家可以想象成多个魔方组合成一个立方体，查找其中一个魔方中的一个小单元格的过程，就容易理解了。

三、CUDA多线程示例

看下面的例程


#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <stdlib.h>

// CUDA核函数：向量加法
__global__ void vecAdd(const float *a, const float *b, float *c, int num) {
	int i = blockDim.x * blockIdx.x + threadIdx.x;
	if (i < num) {
		c[i] = a[i] + b[i];
	}
}

int main(void) {
	// 设置向量大小
	int num = 100000;
	size_t size = num * sizeof(float);
	printf("[Vector addition of %d elements]\n", num);

	// 分配主机内存
	float *h1 = (float *)malloc(size);
	float *h2 = (float *)malloc(size);
	float *h3 = (float *)malloc(size);

	// 初始化主机数组
	for (int i = 0; i < num; ++i) {
		h1[i] = rand() / (float)RAND_MAX;
		h2[i] = rand() / (float)RAND_MAX;
	}

	// 分配设备内存
	float *d1 = NULL;
	float *d2 = NULL;
	float *d3 = NULL;
	cudaMalloc((void **)&d1, size);
	cudaMalloc((void **)&d2, size);
	cudaMalloc((void **)&d3, size);

	// 拷贝数据到设备
	cudaMemcpy(d1, h1, size, cudaMemcpyHostToDevice);
	cudaMemcpy(d2, h2, size, cudaMemcpyHostToDevice);

	// 启动核函数
	int tdPerBlock = 256;
	int bPerGrid = (num + tdPerBlock - 1) / tdPerBlock;
	printf("CUDA set: %d blocks of %d threads\n", bPerGrid, tdPerBlock);
	vecAdd << <bPerGrid, tdPerBlock >> >(d1, d2, d3, num);

	// 拷贝结果回主机
	cudaMemcpy(h3, d3, size, cudaMemcpyDeviceToHost);

	// 验证结果
	for (int i = 0; i < num; ++i) {
		if (fabs(h1[i] + h2[i] - h3[i]) > 1e-5) {
			fprintf(stderr, "verification err id: %d!\n", i);
			exit(EXIT_FAILURE);
		}
	}
	printf("Test OK\n");

	// 释放内存
	free(h1);
	free(h2);
	free(h3);
	cudaFree(d1);
	cudaFree(d2);
	cudaFree(d3);
	cudaDeviceSynchronize();
	system("pause");
	return 0;
}

这段代码利用CUDA的并行框架实现了向量的加法。体现了CUDA多线程编程的用法和流程即从主机内存到设备内存的分配，然后到数据拷贝、Kernel调用，最后将相关数据拷贝回主机程序。

四、CUDA多线程分析

CUDA开发为什么需要这种线程的层次结构，毕竟它增加了相关的编程的复杂度。原因就在于，线程之间往往需要进行数据的交换以及同步。正常情况下，只有在一个块内的线程进行通信，而不在同一收款人的线程在Kernel执行时无法进行通信。这样做的好处就是允许在同一SM中允许程序独立的进行块的调度。而如果日后发布了有更多SM的新的硬件（GPU）而且代码的并行性设计的很好的话，那么代码就可以具有非常良好的可扩展性。而这恰恰是开发者最需要的一种结果。