CUDA从入门到放弃（五）：CUDA线程模型 CUDA Thread Hierarchy

1 简介

CUDA线程模型精心构建了三个层次：线程（Thread）、线程块（Block）和线程网格（Grid），它们协同工作，共同实现高效的并行计算。

2 线程（Thread）

线程（Thread）：作为计算的基本单元，线程承载着实际的计算任务。每个线程都拥有独一无二的ID，这一ID使得线程能够精确地确定自己在所属线程块中的位置。CUDA线程模型巧妙地结合了二维线程块和三维线程网格，通过这种层次化的组织结构，使得大规模并行计算成为可能，从而极大地提升了计算效率。

3 线程块（Block）

线程块（Block）：线程块是由多个线程组成的并行执行单元，这些线程共享一些资源，如共享内存，这种共享机制大大简化了线程间的通信与协作。线程块的大小既可以在编译时预先确定，也可以根据实际情况在运行时进行动态调整。每个线程块都具有唯一的块ID，这使得线程能够准确识别自己在块中的位置。

4 线程网格（Grid）

网格（Grid）：网格是线程块的集合体，它负责管理和组织线程块。在同一个线程网格中，线程块之间共享全局资源，如全局内存，这使得网格内的线程块能够高效地协同工作，共同处理大规模的数据集。网格的大小同样具备灵活性，既可以在编译时设定，也可以在运行时根据需求进行动态调整。每个线程网格都拥有独特的网格ID，通过这个ID，线程能够清晰地定位自己在整个网格中的位置。

5 线程块集群（Thread Block Clusters）

线程块集群（Thread Block Clusters）由线程块组成。类似于线程块中的线程保证在流多处理器上协同调度，集群中的线程块也保证在GPU中的GPU处理集群（GPC）上协同调度。

与线程块类似，集群也组织成一维、二维或三维结构。集群中的线程块数量可由用户定义，在CUDA中，集群的最大线程块数支持为8，作为一个可移植的集群大小。

6 统计直方图示例

#include <stdio.h>  
#include <cuda_runtime.h>  #define N 1000000 // 数组大小  
#define BIN_COUNT 256 // 直方图条目数量  // CUDA内核函数，用于计算直方图  
__global__ void histogramKernel(int *data, int *histogram, int binCount) {  int index = threadIdx.x + blockIdx.x * blockDim.x;  int binIndex = data[index] / (N / binCount); // 计算bin索引  if (index < N) {  atomicAdd(&histogram[binIndex], 1); // 原子操作增加对应bin的计数  }  
}  int main() {  int *h_data, *d_data, *h_histogram, *d_histogram;  h_data = (int*)malloc(N * sizeof(int));  h_histogram = (int*)calloc(BIN_COUNT, sizeof(int));  // 初始化数据（这里仅为示例，实际应用中应有实际数据）  for (int i = 0; i < N; i++) {  h_data[i] = rand() % (BIN_COUNT * (N / BIN_COUNT));  }  // 分配GPU内存  cudaMalloc((void**)&d_data, N * sizeof(int));  cudaMalloc((void**)&d_histogram, BIN_COUNT * sizeof(int));  // 将数据从主机传输到GPU  cudaMemcpy(d_data, h_data, N * sizeof(int), cudaMemcpyHostToDevice);  // 设置CUDA内核的线程块和网格大小  int blockSize = 256;  int gridSize = (N + blockSize - 1) / blockSize;  // 执行CUDA内核函数  histogramKernel<<<gridSize, blockSize>>>(d_data, d_histogram, BIN_COUNT);  // 等待CUDA操作完成  cudaDeviceSynchronize();  // 将直方图结果从GPU传回主机  cudaMemcpy(h_histogram, d_histogram, BIN_COUNT * sizeof(int), cudaMemcpyDeviceToHost);  // 打印直方图结果（仅为示例，实际应用中可能有其他处理方式）  for (int i = 0; i < BIN_COUNT; i++) {  printf("Bin %d: %d\n", i, h_histogram[i]);  }  // 释放内存  free(h_data);  free(h_histogram);  cudaFree(d_data);  cudaFree(d_histogram);  return 0;  
}

参考资料
1 CUDA编程入门
2 CUDA编程入门极简教程
3 CUDA C++ Programming Guide
4 CUDA C++ Best Practices Guide
5 NVIDIA CUDA初级教程视频
6 CUDA专家手册 [GPU编程权威指南]
7 CUDA并行程序设计：GPU编程指南
8 CUDA C编程权威指南