CUDA编程（二）CUDA中的线程组织

grace 1314

于 2023-10-22 21:08:23 发布

阅读量165

点赞数

文章标签： c++ windows visual studio

本文链接：https://blog.csdn.net/qq_43616651/article/details/133974278

版权

1.写一个Hello World from GPU

本文使用win11+CUDA11.6

安装好cuda后，打开vs2022，创建项目时有项目模板，打开

进行程序编写。

C++中的Hello World程序

流程，在IDE中写源代码，用g++编译得到.out可执行文件，执行，得到显示。

而在 CUDA中的HelloWorld程序：.cu文件

#include <stdio.h>

int main(void) {
	printf("Hello World!\n");
	return 0;
}

只有主机函数的CUDA程序，

CUDA程序的编译器驱动（compilerdriver）nvcc支持编译纯粹的C++代码。

CUDA程序的编译器驱动nvcc在编译一个CUDA程序时，会将纯粹的C++代码交给C++的编译器（如前面提到的g++或cl）去处理，它自己则负责编译剩下的部分。CUDA程序源文件的后缀名默认是.cu，所以我们可以将上面写好的源文件更名为hello1.cu，然后用nvcc编译。

上述程序只有主机Host代码。

一个真正利用了GPU的CUDA程序既有主机代码,也有设备代码（需要设备执行的代码）

主机对设备的调用时通过和函数（kernel function）来实现的。

所以，一个典型的、简单的CUDA程序的结构具有下面的形式：

CUDA的核函数，两个重要的点：

1.必须备限定词__global__修饰。其中global前后是双下划线。

2.核函数的返回类型必须是空类型，即void。

写一个打印字符串的核函数。代码win11，vs2022，cuda11.6

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

//__global__ 关键字需要包含头文件cuda_runtime.h
 void __global__  hello_from_gpu() {

	 //printf函数需要包含头文件<stdio.h>
	 printf("Hello World from the GPU!\n");

}

int main(void) {
	//核函数调用，有一对3括号，还有逗号隔开两个数字。
	//三括号中的数用来指明核函数中的线程数目和排列情况。
	//核函数中的线程thread常组织为若干线程块thread block；
	//三括号中的第一个数字。可以看作线程块的个数，第二个数字可以看作每个线程块中的线程数
	//一个核函数的全部线程块组成一个网格grid，而线程块的个数即为网格大小（grid size）。
	//线程块中含有同样数目的线程，称为线程块大小（block size）
	//所以核函数中的总线程数就等于网格大小乘以线程块大小。
	//而三括号中的两个数字分别就是网格和线程大小。即<<<网格大小，线程块大小>>>。
	//这里主机只指派了一个线程，网格大小和线程块大小都是1，即1X1=1；
	hello_from_gpu << <1, 1 >> > ();
	//cudaDeviceSynchronize，是CUDA runtime 的API函数。
	/*
	这是因为调用输出函数时，输出流是先存放在缓冲区的，
	而缓冲区不会自动刷新。只有程序遇到某种同步操作时缓冲区才会刷新。
	函数cudaDeviceSynchronize的作用是同步主机与设备，所以能够促使缓冲区刷新。Synchronize 同步
	*/
	cudaDeviceSynchronize();
	return 0;
}

在函数名hello_from_gpu 和括号 () 之间有一对三括号<<<1,1> >>，里面还有用逗号隔开的两个数字。调用核函数时为什么需要这对三括号里面的信息呢？这是因为，一块GPU中有很多（例如，TeslaV100中有5120个）计算核心，从而可以支持很多线程（thread）。主机在调用一个核函数时，必须指明需要在设备中指派多少个线程，不然设备不知道如何工作。三括号中的数就是用来指明核函数中的线程数目以及排列情况的。核函数中的线程常组织为若干线程块（threadblock）：三括号中的第一个数字可以看作线程块的个数，第二个数字可以看作每个线程块中的线程数。一个核函数的全部线程块构成一个网格（grid），而线程块的个数就记为网格大小（gridsize）。每个线程块中含有同样数目的线程，该线程的数目称为线程块大小（blocksize）。所以，核函数中总的线程数就等于网格大小乘以线程块大小，而三括号中的两个数字分别就是网格大小和线程块大小，即<<<网格大小,线程块大小>>>。所以，在上述程序中，主机只指派了设备的一个线程，网格大小和线程块大小都是1，即1X1=1 。

上述程序注意事项，

在vs中，用__global__关键字需要cuda_runtime头文件。

在vs中<<< >>>会显示错误，但不用担心，可以运行，所以一般都是在linux中运行。

运行结果

以上都是在.cu文件中书写。

2.CUDA中的线程组织

核函数中允许指派很多线程，GPU往往有几千个计算核心，而总的线程数必须至少等于计算核心书时，才能充分利用GPU中的全部计算资源。

实际上总的线程数大于计算和行书才能更充分地利用GPU中的计算资源，因为这会让计算和内存访问之间及不同的计算之间合理地重叠，从而减小计算核心空闲的时间。

根据需要，在调用核函数时可以指定使用多个线程。

对于以下示例

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

 void __global__  hello_from_gpu() {

	 //printf函数需要包含头文件<stdio.h>
	 printf("Hello World from the GPU!\n");

}

int main(void) {

	hello_from_gpu << <2, 4 >> > ();
	cudaDeviceSynchronize();
	return 0;
}

网格大小2，线程块大小4，总的线程数时2X4=8。即该程序中的核函数调用将指派8个线程。核函数中代码的执行方式时“单指令，多线程”，即每个线程都回执行同一串指令。既然核函数的指令是打印一个字符串，那么编译运行就得到下面结果：

8行同样的文字。

其中，每一行对应一个指派的线程。每一行分别是哪一个线程输出的呢？

使用线程索引：

一个个函数可以指引多个线程，而这些线程的组织结构是由执行配置（execution cinfiguration）

<<<grid_size,block_size>>>决定。这里的grid_size网格大小和block_size（线程块大小）一般来说是一个结构体的变量，但也可以是一个普通的整型变量。

首先考虑整型变量，这两个整型变量的乘积就是被调用核函数中的总的线程数。

开普勒架构的最大允许线程块是1024，最大允许的网格大小是2^31-1。所以上述的执行配置做多可以指派大约两亿个线程。这通常是远大于一般编程问题中常用的线程数目的。

一般来说只要线程数比GPU中的计算核心数（几百至几千个）多几倍时，就有可能充分地利用GPU中的全部计算资源。

总之，一个核函数允许指派的线程数目是巨大的，能够满足几乎所有应用程序的要求。需要指出的是，一个核函数中虽然可以指派如此巨大数目的线程数，但在执行时能够同时活跃（不活跃的线程处于等待状态）的线程数是由硬件（主要是CUDA核心数）和软件（即核函数中的代码）决定的。

每个线程在核函数中都有一个唯一的身份标识，由于我们用两个参数指定了线程数目，那么自然地，每个线程的身份可由两个参数确定。在核函数内部，程序是知道执行配置参数grid_size和block_size的值的。这两个值分别保存于如下内建变量（built-in variable）：

gridDim.x:该变量的数值等于执行配置中变量grid_size的数值。

blockDim.x:该变量的数值等于执行配值中变量block_size的数值。

类似地，在核函数中预定义了如下标识线程地内建变量：

blockIdx.x:该变量指定一个线程在一个网格中线程块的指标，其取值范围是从0~gridDim-1。

threadIdx.x:该变量指定一个线程在一个线程块中的线程指标，其取值范围是从0~blockDim.x-1。

上面.x指只有一个维度。

举一个具体的例子。假如某个核函数的执行配置是<<<10000,256>>>，那么网格大小gridDim.x的值为10000，线程块大小blockDim.x的值为256。线程块指标 blockIdx.x 可以取0到9999之间的值，而每一个线程块中的线程指标threadIdx.x可以取0到255之间的值。

当blockIdx.x 等于 0 时，所有 256 个 threadIdx.x 的值对应第 0 个线程块；当blockIdx.x 等于 1 时，所有 256 个 threadIdx.x 的值对应于第 1 个线程块；依此类推。

改写程序，利用该身份标识判断哪一行是由哪个线程输出的。

#include <stdio.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
/*
这里调用 blockIdx.x，threadIdx.x，在win11，vs2022中需要加 device_launch_parameters.h 头文件。
*/

 void __global__  hello_from_gpu() {

	 const int bid = blockIdx.x;  //bid,block id
	 const int tid = threadIdx.x; //tid,thread id
	 printf("Hello World from the %d block and the %d thread the GPU!\n",bid,tid);
}

int main(void) {

	hello_from_gpu << <2, 4 >> > ();

	cudaDeviceSynchronize();
	return 0;
}

输出结果为：

现在将一维索引推广至多维：

gridDim,blockDim,blockIdx,threadIdx,本身是结构体或者类。

gridDim和blockDim是Dim3类型，是一个结构体，具有x，y，z，这3个成员。

gridDim.x gridDim,y gridDim,z

blockDim.x blockDim.y blockDim.z

结构体dim3也在头文件vector_types.h定义，除了和结构体uint3有同样的3个成员之外，还在使用C++程序的情况下定义了一些成员函数，如下面使用的构造函数。

blockIdx和threadIdx是uint3类型，是一个结构体，具有x，y，z，这3个成员。

即，blockIdx.x blockIdx.y blockIdx.z

threadIdx.x threadIdx.y threadIdx.z

也就是说，该结构体由3个无符号整数类型的成员构成。

以上内建变量都只在核函数中有效，并有一下关系。

blockIdx.x的取值范围是从0到gridDim.x-1。

blockIdx.y的取值范围是从0到gridDim.y-1。

blockIdx.z的取值范围是从0到gridDim.z-1。

threadIdx.x的取值范围是从0到blockDim.x-1。

threadIdx.y的取值范围是从0到blockDim.y-1。

threadIdx.z的取值范围是从0到blockDim.z-1。

网格大小和线程块大小是在调用核函数时通过执行配置指定的。在之前的例子中，我们用的执行配置仅仅用了两个整数。

我们知道，这两个整数的值将分别赋给内建变量gridDim.x和blockDim.x。此时，gridDim 和blockDim中没有被指定的成员取默认值1。在这种情况下，网格和线程块实际上都是“一维”的。

也可以用结构体Dim3定义多维的网格和线程块。

即：

dim3 grid_size(Gx,Gy,Gz);
dim3 block_size(Bx,By,Bz);

当第3个维度是1时，可以写作：

dim3 grid_size(Gx,Gy);
dim3 block_size(Bx,By);

多维的网格和线程块本质上还是一维的，就像多维数组本质上也是数组一样，与一个多维线程指标threadIdx.x、threadIdx.y,threadIdx.z对应的一维指标或者索引为；

int tid = threadIdx.z*blockDim.x*blockDim.y+threadIdx.y*blockDim.x+threadIdx.x;

例如下图，

cuda中threadIdx、blockIdx、blockDim和gridDim的使用 - 知乎图来自

gridDim.x=3    gridDim.y=3   gridDim.z=3

blockDim.x=4   blockDim.y=4  blockDim.z=4

blockIdx.x、blockIdx.y、blockIdx.z分别表示当前线程块所处的线程格的坐标位置

threadIdx.x、threadIdx.y、threadIdx.z分别表示当前线程所处的线程块的坐标位置

grid中总的线程数为：

N = gridDim.x * gridDim.y * gridDim.z * blockDim.x * blockDim.y * blockDim.z

线程块索引：

blockId = blockIdx.x + blockIdx.y*gridDim.x + blockIdx.z*gridDim.x*gridDim.y;

当前线程位于线程块中的哪一个线程threadId

threadId = threadIdx.x + threadIdx.y*blockDim.x + threadIdx.z*blockDim.x*blockDim.y;

标题 cuda中threadIdx、blockIdx、blockDim和gridDim的使用 - 知乎

grid是一个三维结构

对于红色block块，三维结构最里面最下面的块为起始，正方向为向上z，向外x，向右y，那么

blockIdx=（1，1，2）；

blockIdx.z*gridDim.x*gridDim.y

指从下往上计算z位置下面层的block块数，从0开始，共2*3*3=18个。

blockIdx.y*gridDim.x

指从左往右计算y位置的块数，为1*3

blockIdx.x

最后加上x的位置。

得到block索引22，block索引一共是0到26。

threadId同理。

计算一个线程块中一共有多少个线程M

M = blockDim.x*blockDim.y*blockDim.z

那么通过blockid和threadid即可得到当前线程的序列号

idx = threadId + M*blockId;

本质就是一个3维堆叠。第多少个线程按照三维展开成一维序列，得到当前线程序列号。

总结，求一个线程在一个Grid中的序列号。

需要知道gridDim，blockDim。

需要知道blockIdx，threadIdx。

均是3维的。

1.求该线程所在block在grid中的序列号blockid，通过gridDim与blockIdx求得。

blockId = blockIdx.x + blockIdx.y*gridDim.x + blockIdx.z*gridDim.x*gridDim.y;

2.求该线程在block中的序列号threadid，通过blockDim与threadx求得。

threadId = threadIdx.x + threadIdx.y*blockDim.x + threadIdx.z*blockDim.x*blockDim.y;

3.计算每个block中thread的个数

M = blockDim.x*blockDim.y*blockDim.z

4.求当前线程的序列号

idx = threadId + M*blockId;

线程束：线程->线程束->线程块->线程网格

一个线程块中的线程还可以细分为不同的线程束（threadwarp）。

一个线程束（即一束线程）是同一个线程块中相邻的warpSize个线程。warpSize也是一个内建变量，表示线程束大小，其值对于目前所有的GPU架构都是32。

所以，一个线程束就是连续的32个线程。具体地说，一个线程块中第0到第31个线程属于第0个线程束，第32到第63个线程属于第1个线程束，依此类推。

代码实践：

#include <stdio.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
/*
这里调用 blockIdx.x，threadIdx.x，在win11，vs2022中需要加 device_launch_parameters.h 头文件。
*/

 void __global__  hello_from_gpu() {

	 const int bid = blockIdx.x;  //bid,block id
	 const int tidx = threadIdx.x; //tid,thread id
	 const int tidy = threadIdx.y; //tid,thread id
	 printf("Hello World from the %d block and the (%d,%d) thread the GPU!\n",bid,tidx,tidy);
}

int main(void) {
	const dim3 block_size(2, 4);

	hello_from_gpu << <1, block_size >> > ();

	cudaDeviceSynchronize();
	return 0;
}

输出结果：

因为grid_size=1，所以block为0；

block_size = 2x4,threadIdx.x从0到1，threadIdx.y从0到4；threadIdx.x位于最内层，变化慢。

网格与线程块大小的限制：

CUDA中对能够定义的网格大小和线程块大小做了限制。对任何从开普勒到图灵架构的GPU来说，网格大小在x、y和z这3个方向的最大允许值分别为31、和；线程块大小在x、y和z这3个方向的最大允许值分别为2^31-1、65535和65535。

另外还要求线程块总的大小，即blockDim.x、blockDim.y和blockDim.z的乘积不能大于1024。也就是说，不管如何定义，一个线程块最多只能1024有个线程。这些限制是必须牢记的。

grid_Size（3个方向最大允许值分别为2^31-1、65535和65535）

block_Size（3个方向最大值31，31，31）

每个线程块最多只能有1024个线程。

一个线程块最多只能1024有个线程

CUDA中的头文件

使用nvcc编译器驱动编译.cu文件时，将自动包含必要的CUDA头文件，如和。因为包含了，故用nvcc编译CUDA程序时甚至不需要在.cu文件中包含。

windows需要自己加头文件。一般是以下两个

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

nvcc编译cuda程序

CUDA的编译器驱动（compilerdriver）nvcc先将全部源代码分离为主机代码和设备代码。

主机代码完整地支持C++语法，但设备代码只部分地支持C++。

nvcc先将设备代码编译为PTX（ParallelThreadeXecution，并行线程执行）伪汇编代码，再将PTX代码编译为二进制的cubin目标代码。在将源代码编译为PTX代码时，需要用选项-arch=compute_XY指定一个虚拟架构的计算能力，用以确定代码中能够使用的CUDA功能。在将PTX代码编译为cubin代码时，需要用选项-code=sm_ZW指定一个真实架构的计算能力，用以确定可执行文件能够使用的GPU。真实架构的计算能力必须等于或者大于虚拟架构的计算能力。

nvcc 有一种称为即时编译（just-in-time compilation）的机制，可以在运行可执行文件时从其中保留的PTX代码临时编译出一个cubin目标代码。

grace 1314

关注

0
点赞
踩
3

收藏

觉得还不错? 一键收藏
1
评论
CUDA编程（二）CUDA中的线程组织

假如某个核函数的执行配置是<<<10000,256>>>，那么网格大小gridDim.x的值为10000，线程块大小blockDim.x的值为256。线程块指标 blockIdx.x 可以取0到9999之间的值，而每一个线程块中的线程指标threadIdx.x可以取0到255之间的值。所以，核函数中总的线程数就等于网格大小乘以线程块大小，而三括号中的两个数字分别就是网格大小和线程块大小，即<<<网格大小,线程块大小>>>。网格大小2，线程块大小4，总的线程数时2X4=8。这些限制是必须牢记的。
复制链接

扫一扫