第 3 章：GPU 设计中的并行编程模型

            网格
       ________________
      /   /   /   /   /
     /   /   /   /   / 
    /   /   /   /   /
   /   /   /   /   /
  /__/__/__/__/__/
    |   |   |   |
    |   |  块  |
    |   |   |   |
  线程 线程 ...

         Warp
   ________________
  /   /   /   /   /
 /   /   /   /   / 
/   /   /   /   /
   |   |   |
   | 分支 |
   |   |   |
  / \ / \ / \
 /   X   \   \
/   / \   \   \
   /   \   \
  /     \   \
 /       \   \
/         \   \
           \
```重新集中

|---|---|
| Shared Memory | 共享内存 |
| Local Memory | 局部内存 |
| Figure 3.3: GPU memory hierarchy. | 图 3.3: GPU 内存层次结构。 |
| CUDA Programming Model and APIs | CUDA 编程模型和 API |
| CUDA (Compute Unified Device Architecture) is a parallel computing platform and programming model developed by NVIDIA for general-purpose computing on GPUs. CUDA provides a set of extensions to standard programming languages, such as C, C++, and Fortran, that enable programmers to express parallelism and leverage the computational power of NVIDIA GPUs. | CUDA (Compute Unified Device Architecture) 是 NVIDIA 为通用 GPU 计算开发的并行计算平台和编程模型。CUDA 为标准编程语言(如 C、C++ 和 Fortran)提供了一系列扩展,使程序员能够表达并行性并利用 NVIDIA GPU 的计算能力。 |
| CUDA Programming Model | CUDA 编程模型 |
| The CUDA programming model is based on the concept of kernels, which are functions executed in parallel by a large number of threads on the GPU. The programmer specifies the number of threads to be launched and their organization into a grid of thread blocks. | CUDA 编程模型基于内核的概念,内核是在 GPU 上并行执行的大量线程函数。程序员指定要启动的线程数量及其在线程块网格中的组织方式。 |
| Thread: The basic unit of execution in CUDA. Each thread has its own program counter, registers, and local memory. | 线程:CUDA 中执行的基本单元。每个线程都有自己的程序计数器、寄存器和局部内存。 |
| Block: A group of threads that can cooperate and synchronize with each other. Threads within a block are executed on the same streaming multiprocessor and can communicate via shared memory. | 块:一组可以相互协作和同步的线程。块内的线程在同一个流多处理器上执行,可以通过共享内存进行通信。 |
| Grid: A collection of thread blocks that execute the same kernel. The grid represents the entire problem space and can be one-, two-, or three-dimensional. | 网格:执行同一个内核的一组线程块。网格代表整个问题空间,可以是一维、二维或三维的。 |
| CUDA also provides built-in variables (e.g., threadIdx, blockIdx, blockDim, gridDim) that allow threads to identify themselves and compute memory addresses based on their position in the thread hierarchy. | CUDA 还提供了内置变量(如 threadIdx、blockIdx、blockDim 和 gridDim),允许线程根据在线程层次结构中的位置来识别自己并计算内存地址。 |
| Figure 3.4: CUDA programming model. | 图 3.4: CUDA 编程模型。 |下面是这个 Markdown 文件的中文翻译版本,其中代码部分的注释进行了翻译,代码本身没有翻译。

CUDA 将 GPU 内存层次结构暴露给程序员,允许对数据放置和移动进行显式控制。CUDA 中的主要内存空间有:

- 全局内存：可被内核中的所有线程访问,并在内核启动之间持续存在。全局内存延迟较高,通常用于大型数据结构。
- 共享内存：一种快速的片上内存,由同一个块中的所有线程共享。共享内存用于线程间通信和数据共享。
- 常量内存：一种只读内存空间,用于将只读数据广播到所有线程。常量内存已缓存,提供低延迟访问。
- 纹理内存：一种针对空间局部性进行优化的只读内存空间,通过纹理缓存访问。纹理内存更常用于图形工作负载。
- 局部内存：每个线程的私有内存空间,用于寄存器溢出和大型数据结构。局部内存通常映射到全局内存。

程序员可以使用 CUDA 运行时 API（如 cudaMalloc、cudaMemcpy 和 cudaFree）在主机（CPU）和设备（GPU）内存之间分配和传输数据。

图 3.5 说明了 CUDA 的内存层次结构。

图 3.5：CUDA 内存层次结构。

### CUDA 同步和协调

CUDA 提供了同步和协调原语,以使线程之间能够进行合作和通信:

- 屏障同步：__syncthreads()以下是文件的中文翻译:

s()函数充当一个屏障,确保块中的所有线程在继续之前都到达了同一个点。

- 原子操作：CUDA支持原子操作(例如,atomicAdd,atomicExch)，允许线程在没有其他线程干扰的情况下对共享或全局内存执行读-改-写操作。
- 线程束级原语：CUDA提供了线程束级内置函数(例如，__shfl,__ballot)，可以在线程束内进行高效的通信和同步。

正确使用同步和协调原语对于在CUDA中编写正确和高效的并行程序至关重要。

示例3.1显示了一个简单的CUDA内核,它执行向量加法。

```c
// 执行向量加法的CUDA内核
__global__ void vectorAdd(int *a, int *b, int *c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

int main() {
    int *a, *b, *c;
    int n = 1024;
    
    // 在主机上分配内存
    a = (int*)malloc(n * sizeof(int));
    b = (int*)malloc(n * sizeof(int));
    c = (int*)malloc(n * sizeof(int));
    
    // 初始化输入向量
    for (int i = 0; i < n; i++) {
        a[i] = i;
        b[i] = i * 2;
    }
    
    // 在设备上分配内存
    int *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, n * sizeof(int));
    cudaMalloc(&d_b, n * sizeof(int));
    cudaMalloc(&d_c, n * sizeof(int));
    
    // 将输入向量从主机复制到设备
    cudaMemcpy(d_a, a, n * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, n * sizeof(int), cudaMemcpyHostToDevice);
    
    // 启动内核
    int blockSize = 256;
    int numBlocks = (n + blockSize - 1) / blockSize;
    vectorAdd<<<numBlocks,blockSize>>>(d_a, d_b, d_c, n);
    
    // 将结果向量从设备复制到主机
    cudaMemcpy(c, d_c, n * sizeof(int), cudaMemcpyDeviceToHost);
    
    // 释放设备内存
    cudaFree(d_a);
    cudaFree(d_b); 
    cudaFree(d_c);
    
    // 释放主机内存
    free(a); 
    free(b);
    free(c);
    
    return 0;
}
```以下是提供的 Markdown 文件的中文翻译。代码部分中,注释被翻译,但代码本身没有被翻译。

```c
int n = 0;
}

// 创建两个流
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
 
// 异步将输入数据拷贝到设备
cudaMemcpyAsync(d_a, a, size, cudaMemcpyHostToDevice, stream1);
cudaMemcpyAsync(d_b, b, size, cudaMemcpyHostToDevice, stream2);
 
// 在不同的流中启动内核
kernelA<<<blocks, threads, 0, stream1>>>(d_a);
kernelB<<<blocks, threads, 0, stream2>>>(d_b);
 
// 异步将结果拷贝回主机
cudaMemcpyAsync(a, d_a, size, cudaMemcpyDeviceToHost, stream1);
cudaMemcpyAsync(b, d_b, size, cudaMemcpyDeviceToHost, stream2);
 
// 同步流
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);

__kernel void vector_add(const __global int *a, const __global int *b, __global int *c, int n) {
    int i = get_global_id(0);
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}