cuda程序的后缀：.cu
编译：nvcc hello_world.cu
执行：./hello_world.cu
使用语言还是C++。
1. 核函数
__global__ void add(int *a, int *b, int *c) {
*c = *a + *b;
}
核函数只能访问GPU的内存。也就是显存。CPU的存储它是碰不到的。
并且核函数不能使用变长参数、静态变量、函数指针。
核函数具有异步性。GPU无法控制CPU，CPU也不会去等GPU，所以需要同步，也就是显式调用同步函数。有些线程也是需要同步的。
编写CUDA程序：
int main(void){
主机代码
核函数调用
主机代码
return 0;
}
核函数不支持C++的iostream。
#include<stdio.h>
__global__ void hello_from_gpu(){
printf("Hello from GPU!\n");
__syncthreads();// 显式同步
}
int main(){
hello_from_gpu<<<1,1>>>();// 显式调用核函数
cudaDeviceSynchronize();// 显式同步
return 0;
}
2. 线程块
int main() {
int a = 1;
int b = 2;
int c;
add<<<1, 1>>>(&a, &b, &c);
return 0;
}
线程模型重要概念：
- grid网格
- block线程块
线程分块是逻辑上的划分，物理上线程不分块。
配置线程：<<<grid_num, block_num>>>
第一个参数代表着我们有M个线程块，第二个参数代表着我们的每个线程块中有N个线程。他们都是一维的。这昂个参数保存在内建变量(build-in variable)中。
gridDim.x: 该变量的数值等于执行配置中变量grid_num的值。
blockDim.x: 该变量的数值等于执行配置中变量block_num的值。
最大允许线程块的大小为1024。最大允许的网格大小是 2 3 1 − 1 2^31-1 231−1（针对一维网格）。
实际使用中，总线程数大于实际使用的线程数能更好地利用计算资源，因为这样可以使得GPU在计算的时候内存访问同时进行，节省计算机计算的时间。使得核心一直处于计算中。
启动核函数后，CPU并不会等待核函数执行完毕，立马去执行主机中其他程序。所以我们要做的就是使得这两部分时间重叠。
3. 线程块的索引
int main() {
int a = 1;
int b = 2;
int c;
add<<<1, 1>>>(&a, &b, &c);
return 0;
}
线程索引保存成内s建变量(build-in variable)：
- blockIdx.x: 该变量指定一个线程在一个网格中的线程块索引值，范围0-girdDim.x-1。
- threadIdx.x: 该变量指定一个线程在线程块中的索引值，范围0-blockDim.x-1。
线程具有唯一标识：
I d x = t h r e a d I d x . x + b l o c k D i m . x ∗ b l o c k I d x . x ; Idx = threadIdx.x + blockDim.x * blockIdx.x; Idx=threadIdx.x+blockDim.x∗blockIdx.x;
4. 推广到多维线程
-
CUDA可以组织三维的网格和线程块；
-
blockIdx和threadIdx是类型为uint3的变量，该类型是一个结构体，具有x,y,z三个成员（3个成员都为无符号类型的成员构成）:
- 定义多维网格和线程块（C++构造函数语法）：
dim3 grid_num(Gx,Gy,Gz);
dim3 block_num(Bx,By,Bz);
dim3 grid_num(2,2); // 等价于dim3 grid_num(2,2,1);
dim3 block_num(5,3); // 等价于dim3 block_num(5,3,1);
5. 一维网格 一维线程块
定义grid和block尺寸：
dim3 grid_num(4);
dim3 block_num(8);
调用核函数：
kernel_fun<<<grid_num, block_num>>>(…);
具体的线程索引方式如图所示。
blockIdx.x从0到3，threadIdx.x从0到7。
计算方式：
I d x = t h r e a d I d x . x + b l o c k D i m . x ∗ b l o c k I d x . x ; Idx = threadIdx.x + blockDim.x * blockIdx.x; Idx=threadIdx.x+blockDim.x∗blockIdx.x;
6. 二维网格 二维线程块
定义grid和block尺寸：
dim3 grid_num(2,2);
dim3 block_num(5,3);
调用核函数：
kernel_fun<<<grid_num, block_num>>>(…);
具体的线程索引方式如图所示。
blockIdx.x从0到1，threadIdx.y从0到1。
blockIdx.x从0到1，threadIdx.y从0到3。
计算方式：
i n t b l o c k I d = b l o c k I d x . x + g r i d D i m . x ∗ b l o c k I d x . y ; i n t t h r e a d I d = t h r e a d I d x . x + b l o c k D i m . x ∗ t h r e a d I d x . y ; i n t i d = b l o c k I d ∗ ( b l o c k D i m . x ∗ b l o c k D i m . y ) + t h r e a d I d ; int blockId = blockIdx.x + gridDim.x * blockIdx.y; int threadId = threadIdx.x + blockDim.x * threadIdx.y; int id = blockId * (blockDim.x * blockDim.y) + threadId; intblockId=blockIdx.x+gridDim.x∗blockIdx.y;intthreadId=threadIdx.x+blockDim.x∗threadIdx.y;intid=blockId∗(blockDim.x∗blockDim.y)+threadId;
7. 三维网格 三维线程块
定义grid和block尺寸：
dim3 grid_num(2,2,2);
dim3 block_num(5,3,1);
调用核函数：
kernel_fun<<<grid_num, block_num>>>(…);
具体的线程索引方式如图所示。
blockIdx.x、blockIdx.y和blcokIdx.z从0到1，
threadIdx.x、threadIdx.y从0到3，threadIdx.z从0到1。
计算方式：
i n t b l o c k I d = b l o c k I d x . x + g r i d D i m . x ∗ b l o c k I d x . y + g r i d D i m . x ∗ g r i d D i m . y ∗ b l o c k I d x . z ; i n t t h r e a d I d = ( t h r e a d I d x . z ∗ ( b l o c k D i m . x ∗ b l o c k D i m . y ) ) + ( t h r e a d I d x . y ∗ b l o c k D i m . x ) + t h r e a d I d x . x ; i n t i d = b l o c k I d ∗ ( b l o c k D i m . x ∗ b l o c k D i m . y ∗ b l o c k D i m . z ) + t h r e a d I d ; int blockId = blockIdx.x + gridDim.x * blockIdx.y + gridDim.x * gridDim.y * blockIdx.z; int threadId= (threadIdx.z * (blockDim.x * blockDim.y) ) + (threadIdx.y * blockDim.x) + threadIdx.x; int id = blockId * (blockDim.x * blockDim.y * blockDim.z) + threadId; intblockId=blockIdx.x+gridDim.x∗blockIdx.y+gridDim.x∗gridDim.y∗blockIdx.z;intthreadId=(threadIdx.z∗(blockDim.x∗blockDim.y))+(threadIdx.y∗blockDim.x)+threadIdx.x;intid=blockId∗(blockDim.x∗blockDim.y∗blockDim.z)+threadId;
三维网格、三维线程块如图所示：
https://github.com/user-attachments/assets/c57924c1-2157-4c73-87ea-36f6842e9eff
Reference
[1]. 权双.CUDA编程基础入门系列（持续更新）[M/OL](2023-07-14)[2024-08-21].https://www.bilibili.com/video/BV1sM4y1x7of/?p=7&share_source=copy_web&vd_source=8b2bc57e71349607b55c9fde6b078ebd
