CUDA编程：对线程模型的理解

一.核函数的性质

1. 核函数在GPU上并行执行
2. 注意：（1）核函数前要有限定词__global__修饰；（2）核函数的返回值必须是void
3. 核函数只能访问GPU内存
4. 核函数不能使用变长参数
5. 核函数不能使用静态变量
6. 核函数不能使用函数指针
7. 核函数具有异步性

二.线程模型

线程模型的重要概念

• 网格 grid
• 线程块 block

线程分块是逻辑上的划分，物理上线程不分块

配置线程

<<<grid_size,block_size>>>
//<<<线程块个数，线程块中线程个数>>>

最大允许线程块大小：1024
最大允许网格大小：2^31-1（针对一维网格）

三.一维线程模型

1. 每个线程在核函数中都有一个唯一的身份标识
2. 每个线程的唯一标识由这两个<<<grid_size,block_size>>>确定；grid_szie，block_size保存在内建变量（bulid-in variable)，目前考虑的是一维的情况
   gridDim.x：该变量的数值等于执行配置中变量grid_size的值
   blockDim.x：该变量的数值等于执行配置中变量block_size的值
3. 线程索引保存成内建变量（build-in variable）：
   （1）blockldx.x：该变量指定一个线程在一个网格中的线程块的索引值，范围为0~gridDim.x-1；
   （2）threadldx.x：该变量指定一个线程在一个线程块中的线程索引值，范围为0~blockDim.x-1;

线程唯一标识公式：
Idx = threadIdx.x + blockIdx.x * blockDim.x
通俗的解释一下：

• ldx是这个线程的坐标
• threadIdx.x是某一个线程块中线程的序号，在图中的范围是0~3
• blockIdx.x是一个网格中线程块的序号，在图中的范围是0~1
• blockDim.x是一个线程块的大小，在图中为的值4

举个例子：
红色块所在位置，该线程的唯一标识为：ldx = 1+1*4 = 5

//test.cu
#include <stdio.h>

__global__ void hello_from_gpu()
{
	printf("hello world from the gpu\n");
}

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

test.cu编译后运行结果：打印八行：hello world from the gpu

//test01.cu
#include <stdio.h>

__global__ void hello_from_gpu()
{	
	const int bid = blockIdx.x;
	const int tid = threadIdx.x;
	
	const int id =  threadldx.x + blockldx.x * blockDim.x
	printf("hello world from block %d and thread %d,global id %d\n",bid,tid,id);
}

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

test01.cu编译后运行结果：会打印出八行，每一行都明确了该线程所处的线程块，在线程块中的位置（第几个线程），以及该线程的唯一标识

四.推广到多维线程

1. cuda可以组织三维的网格和线程块
2. blockIdx和threadIdx是类型为uint3的变量，该类型是一个结构体，具有x，y，z三个成员（三个成员都由无符号类型的成员构成）
3. blockIdx.x的范围————[0,gridDim.x-1]，threadIdx.x的范围————[0,threadIdx.x-1]，y和z同理
   PS：内建变量只在核函数内有效，且无需定义！！！
4. gridDim和blockDim没有指定的维度默认为1
   gridDim.x = grid_size blockDim.x = block_size
   gridDim.y = 1 blockDim.y = 1
   gridDim.z = 1 blockDim.z = 1
5. 定义多维网格和线程块（C++构造函数语法）：
   dim3 grid_size(Gx,Gy,Gz);
   dim3 block_size(Bx,By,Bz);
   举个例子定义一个221的网格，531的线程块，代码中定义如下：
   dim3 grid_size(2,2)————等价于dim3 grid_size(2,2,1)
   dim3 block_size(5,3)————等价于dim3 grid_size(5,3,1)
6. 多维网格和多维线程块本质是一维的，GPU物理上不分块
7. 线程块大小的限制，限制x×y×z最大值为1024，也就是说一个线程块中最多有1024个线程