nv-cuda-prll-prog-merge-1

Nvidia CUDA 并行编程笔记（二）

  1. 共享内存的本质：它是GPU上可由程序员控制的高速、片上内存，以线程块为单位进行分配和管理。

  2. 声明方法：使用 __shared__ 限定符在内核中声明变量或数组。

  3. 核心价值：用于数据重用以降低延迟，以及实现线程块内的线程间协作与通信。

  4. 关键挑战：并发访问共享内存会引入数据竞争，导致非确定性和错误的结果。

  5. 解决方案：必须使用 __syncthreads() 等同步原语来协调线程块内线程的执行顺序，确保在关键操作点（如数据加载后、计算后）所有线程达成一致。

理解并正确使用共享内存与同步，是编写高效、正确CUDA并行程序的关键一步。在下一节课中，我们将更详细地探讨 __syncthreads() 以及其他同步机制。

009：动态共享内存与存储体冲突 🚀

概述

在本节课中，我们将要学习CUDA编程中的线程同步机制、动态共享内存的分配与管理，并初步了解共享内存的存储体冲突问题。上一节我们介绍了共享内存的基本概念和使用方法，本节中我们来看看如何确保线程间的正确协作，以及如何更灵活地使用共享内存。

线程同步：__syncthreads() 🔄

在上一节中，我们看到了由于线程的异步执行，可能导致数据竞争和不确定的输出。为了确保线程块内所有线程在执行到某个点之前都已完成特定操作，我们需要使用同步机制。

CUDA提供了 __syncthreads() 函数，它是一个线程块内的屏障。其工作原理如下：

  • 所有属于同一线程块的线程在执行到 __syncthreads() 时都会暂停。

  • 线程会在此等待，直到同一线程块内的所有其他线程都到达这个同步点。

  • 一旦所有线程都到达，它们才会一起继续执行之后的代码。

公式/代码描述：

// 线程块内所有线程执行至此
__syncthreads(); // 同步点：所有线程在此等待彼此
// 所有线程同步后，从此处继续执行

同步示例分析

回顾上节课末尾可能产生非确定输出的代码。通过添加 __syncthreads()，我们可以强制线程执行顺序，从而得到确定性的结果（输出始终为3）。

关键点：

  • __syncthreads() 只对同一线程块内的线程有效。不同线程块的线程之间不会通过此函数进行同步。

  • 如果将 __syncthreads() 放在条件语句（如 if）内部，而并非块内所有线程都能满足该条件，那么未能进入该条件的线程将永远无法到达同步点，导致死锁。

  • 对于同一线程束（Warp） 内的线程，由于SIMD（单指令多数据）的锁步执行特性，它们天然地按步骤执行，在某些简单情况下可能不需要显式同步。但涉及不同线程束间的协作时，__syncthreads() 是必需的。

配置L1缓存与共享内存 ⚙️

GPU的L1缓存/共享内存总大小通常为64KB。程序员可以根据内核的需求，调整用于L1缓存和共享内存的空间分配比例。

这通过 cudaDeviceSetCacheConfig 函数实现：

代码描述：

cudaDeviceSetCacheConfig(cudaFuncCachePrefer option);

参数 option 说明：

  • cudaFuncCachePreferNone: 默认。编译器自动根据内核代码决定分配策略。

  • cudaFuncCachePreferShared: 偏好共享内存。分配 48KB共享内存 + 16KB L1缓存。

  • cudaFuncCachePreferL1: 偏好L1缓存。分配 16KB共享内存 + 48KB L1缓存。

  • cudaFuncCachePreferEqual: 均等分配。分配 32KB共享内存 + 32KB L1缓存。

使用示例：

cudaDeviceSetCacheConfig(cudaFuncCachePreferL1); // 为该内核设置L1偏好配置
myKernel<<<grid, block>>>(...); // 启动内核

注意： 如果内核声明的共享内存数组大小超过了配置的共享内存容量，编译器会报错。

动态共享内存 🧠

之前我们使用的共享内存，其大小在编译时就是确定的（例如 __shared__ int s_data[256];）。然而，有时我们需要根据运行时的输入数据来决定共享内存的大小，这时就需要使用动态共享内存。

分配动态共享内存

动态共享内存的分配分为两个步骤：

  1. 内核启动时指定大小：在内核调用（<<< >>>）的第三个执行配置参数中，指定需要分配的共享内存字节数。

  2. 内核内声明：在内核中，使用 extern __shared__ 关键字声明一个未指定大小的共享内存数组。

代码描述：

// 主机端代码
int size = 1024; // 运行时决定的大小
myKernel<<<grid, block, size * sizeof(float)>>>(...);

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/3026efa599658b0243d3bfd2c79f4da7_89.png)

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/3026efa599658b0243d3bfd2c79f4da7_91.png)

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/3026efa599658b0243d3bfd2c79f4da7_93.png)

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/3026efa599658b0243d3bfd2c79f4da7_95.png)

// 设备端（内核）代码
__global__ void myKernel(...) {
    extern __shared__ float s_data[]; // 动态共享内存数组
    // 使用 s_data，其大小为启动内核时指定的 size * sizeof(float)
}

使用动态共享内存

动态共享内存通常作为一个“大块”内存进行分配。如果需要将其用作多个不同大小的数组，可以在内核内通过指针计算来手动划分。

示例：将一块动态共享内存分为两部分

__global__ void myKernel(int n) {
    extern __shared__ int s_data[];
    int *s_part1 = s_data;          // 第一部分起始地址
    int *s_part2 = &s_data[n/2];    // 第二部分起始地址（偏移 n/2）
    // 现在可以独立使用 s_part1 和 s_part2
}

共享内存的存储体冲突 ⚡️（初步了解）

共享内存为了实现高带宽并行访问，被组织成多个存储体。当前GPU架构通常有32个存储体（与线程束大小对应）。

  • 理想情况：一个线程束中的32个线程分别访问32个不同存储体中的地址。这些访问可以同时进行，实现高带宽。

  • 存储体冲突：如果同一个线程束内的多个线程访问同一个存储体中的不同地址，这些访问将被迫串行化，从而显著降低访问效率。

类比：

想象一个有32个柜台的银行（32个存储体）。如果32个人（一个线程束）每人去一个不同的柜台办理业务，所有业务可以同时处理。但如果有多个人都要去同一个柜台，他们就必须排队，导致整体速度变慢。

简单示例：

假设共享内存按4字节（例如int）进行存储体映射。

  • 无冲突访问：线程0访问地址0，线程1访问地址128，线程2访问地址256...（地址间隔较大，可能落入不同存储体）。

  • 有冲突访问：线程0访问地址0，线程1访问地址4，线程2访问地址8...（这些地址可能位于同一存储体，导致冲突）。

存储体冲突是优化共享内存访问性能时需要重点考虑的问题。我们将在后续课程中详细讨论其原理和避免方法。

总结 🎯

本节课中我们一起学习了：

  1. 线程同步：使用 __syncthreads() 函数来同步线程块内的线程，确保数据操作的顺序性和正确性。

  2. 内存配置：如何使用 cudaDeviceSetCacheConfig 来配置L1缓存和共享内存的分配比例，以适应不同内核的需求。

  3. 动态共享内存：如何在内核启动时动态指定共享内存的大小，并在内核中使用 extern __shared__ 进行声明和划分。

  4. 存储体冲突：初步了解了共享内存的存储体组织结构，以及多线程访问同一存储体时可能引发的性能下降问题。

掌握这些知识，你将能够编写出协作更安全、内存使用更灵活、并为后续性能优化打下基础的CUDA内核。下一节我们将更深入地探讨存储体冲突及矩阵计算等高级优化技术。

010：存储体冲突、纹理内存、常量内存及计算能力 🚀

概述

在本节课中，我们将继续深入学习GPU内存体系结构。我们将详细探讨共享内存中的存储体冲突问题，并通过示例理解其成因和影响。接着，我们将介绍两种特殊的只读内存：纹理内存和常量内存，了解它们的特性、声明方式和使用场景。最后，我们将简要介绍CUDA计算能力的概念及其与GPU硬件家族的关系。

存储体冲突详解

上一节我们介绍了共享内存的基本概念和__syncthreads()同步原语。本节中，我们来看看共享内存内部的组织结构——存储体，以及由此引发的“存储体冲突”问题。

共享内存被组织成32个存储体。对共享内存的访问由线程束（32个线程）执行。理解以下关键点至关重要：

  1. 对同一存储体的访问是顺序执行的。

  2. 对同一存储体中不同字的访问会导致顺序化，从而引发存储体冲突。

  3. 对同一存储体中同一字的访问不会导致顺序化，因此不会引发存储体冲突。

  4. 连续的字存储在相邻的存储体中。

以下是理解存储体访问模式的几个核心示例：

无存储体冲突的访问模式

当线程束中的每个线程访问连续且位于不同存储体中的字时，所有访问可以并行完成，没有冲突。

__global__ void noBankConflict(float* output) {
    __shared__ float s_data[1024];
    int tid = threadIdx.x;
    // 每个线程访问不同存储体中的元素
    s_data[tid] = tid * 1.0f; // 乘以1仅为示意，无实际影响
    // ... 其他操作
}

在这个例子中，thread0访问s_data[0]（存储体0），thread1访问s_data[1]（存储体1），以此类推。所有访问并行发生。

导致存储体冲突的访问模式

当线程束中的多个线程访问同一存储体中的不同字时，这些访问会被顺序化，导致性能下降。

__global__ void bankConflict(float* output) {
    __shared__ float s_data[1024];
    int tid = threadIdx.x;
    // 导致冲突的访问模式：多个线程访问同一存储体
    s_data[tid * 32] = tid * 1.0f;
    // ... 其他操作
}

在这个例子中，thread0访问s_data[0]（存储体0），thread1访问s_data[32]（存储体0），thread2访问s_data[64]（存储体0）。由于它们访问的是存储体0中的不同字，这些请求会被顺序处理，形成存储体冲突。这被称为“多路存储体冲突”（例如，本例是32路冲突）。

存储体冲突可视化

为了更直观地理解，我们可以通过以下图示来分析几种访问模式：

  • 场景A：每个线程访问不同存储体中的唯一字。这是无冲突的理想情况。

  • 场景B：两个线程（如thread0和thread16）访问同一存储体中的不同字。这会导致2路存储体冲突。

  • 场景C：访问模式看似杂乱，但每个线程访问的存储体仍然是唯一的。这同样是无冲突的。

  • 场景D & E：多个线程访问同一存储体中的同一个字。根据规则3，这属于特殊情况，访问可以广播给所有请求线程，因此不会引发存储体冲突。

优化优先级：如果一个程序同时存在未合并的全局内存访问和存储体冲突，应优先优化全局内存访问，使其合并。因为全局内存访问延迟（400-800周期）远高于共享内存。在全局内存访问优化好后，再着手解决存储体冲突问题。

纹理内存

纹理内存是一种全局的、只读的内存空间，针对具有空间局部性的访问模式（如图像处理中的2D邻域访问）进行了优化。

声明与绑定

使用纹理内存需要以下步骤：

  1. 声明纹理引用：使用texture类型声明一个变量。

    texture<float, 2, cudaReadModeElementType> texRef;

    其中，float是数据类型，2表示是2D纹理，cudaReadModeElementType表示以元素形式读取。

  1. 绑定到设备内存：将纹理引用绑定到已分配的CUDA数组。

    cudaArray* cuArray;
// ... 分配并初始化cuArray (例如使用cudaMallocArray, cudaMemcpyToArray)
cudaBindTextureToArray(texRef, cuArray);

在核函数中访问

在核函数内部，使用tex2D()函数来读取纹理内存。

float value = tex2D(texRef, x, y);

应用示例：图像旋转

以下是一个简化的核函数，演示如何使用纹理内存实现图像旋转（或仿射变换）：

__global__ void transformKernel(float* output, int width, int height, float theta) {
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;

    if (x < width && y < height) {
        // 计算旋转后的坐标 (u, v)
        float u = x * cosf(theta) - y * sinf(theta);
        float v = x * sinf(theta) + y * cosf(theta);

        // 从纹理内存中读取旋转后坐标处的像素值
        // 注意：tex2D会自动处理边界和插值（如果设置）
        float pixel = tex2D(texRef, u + 0.5f, v + 0.5f);

        // 写入输出数组（全局内存）
        output[y * width + x] = pixel;
    }
}

在这个例子中，输入图像被绑定到纹理内存texRef。每个线程计算其对应输出像素(x, y)在输入图像中对应的源位置(u, v)，然后通过tex2D高效地读取该位置的值。纹理内存硬件会对(u, v)附近的像素进行缓存，非常适合这种访问模式。

重要提示：纹理内存是只读的。计算结果需要写入到全局内存（如output数组）中。

常量内存

常量内存是另一种全局只读内存，大小固定（通常为64KB）。它用于存储需要被所有线程频繁访问的常量数据。常量内存通过缓存实现高速访问。

声明与初始化

  1. 声明常量变量：使用__constant__限定符在全局作用域声明。

    __constant__ int constData[256];

  1. 从主机初始化：常量内存只能从主机（CPU）代码进行初始化，使用cudaMemcpyToSymbol函数。

    int h_data = 10;
cudaMemcpyToSymbol(constData, &h_data, sizeof(int));

在核函数中访问

在核函数中，可以像读取普通全局变量一样读取常量内存。

__global__ void myKernel(int* result) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    result[tid] = constData[0] * tid; // 所有线程读取同一个常量值
}

简单示例

以下代码展示了常量内存的完整使用流程：

// 1. 声明常量内存变量
__constant__ unsigned int myConst;

__global__ void constantKernel(unsigned int* data) {
    // 3. 在核函数中访问常量内存
    int tid = threadIdx.x;
    data[tid] = myConst; // 所有线程都将自己的输出设置为常量值
}

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/fe6f14585acd6e9a70d02305c3769a23_161.png)

int main() {
    unsigned int h_val = 10;
    unsigned int* d_data;
    cudaMalloc(&d_data, 32 * sizeof(unsigned int));

    // 2. 从主机复制数据到常量内存
    cudaMemcpyToSymbol(myConst, &h_val, sizeof(unsigned int));

    constantKernel<<<1, 32>>>(d_data);
    // ... 同步并打印d_data，所有32个元素值都为10
    cudaFree(d_data);
    return 0;
}

CUDA计算能力

计算能力（Compute Capability）是一个由主版本号和次版本号组成的版本号（例如8.6），它代表了GPU硬件的架构和功能集。

  • 硬件特性标识：不同的计算能力版本支持不同的硬件特性（如双精度原子操作、动态并行、张量核心等）。

  • 编译指定：在使用nvcc编译时，可以通过-arch=sm_xx标志指定目标计算能力（例如-arch=sm_86），以确保编译器使用该版本支持的指令和优化。

  • 与CUDA版本区别：CUDA版本（如CUDA 11.0）是软件工具包的版本。计算能力是GPU硬件的属性。高版本的CUDA工具包支持更多计算能力的设备。

主要计算能力版本与特性

| 计算能力 | 架构家族 | 新增关键特性示例 |

| :--- | :--- | :--- |

| 1.x | Tesla | 基础CUDA支持 |

| 2.x | Fermi | 线程束内投票、共享内存原子操作 |

| 3.x | Kepler | 统一内存、动态并行性（从3.5开始） |

| 5.x, 6.x | Maxwell, Pascal | 改进的功耗效率、统一内存增强 |

| 7.x | Volta, Turing | 张量核心（AI加速）、独立的线程调度 |

| 8.x | Ampere | 硬件支持的异步拷贝、第三代张量核心 |

| 9.x | Hopper (最新) | 第四代张量核心、Transformer引擎 |

注意：计算能力4.x被跳过。计算能力10.x（Lovelace）和11.x（Blackwell）对应更新的架构，部分仍在开发或发布初期。

选择编译目标时，需要根据程序使用的特性和目标部署平台的GPU来确定合适的计算能力版本。

总结

本节课中我们一起学习了CUDA内存体系的几个高级主题。

我们深入分析了共享内存的存储体冲突，明白了其产生条件是线程束内多个线程访问同一存储体中的不同字，并通过代码示例学会了如何识别和避免这类冲突。

接着，我们探讨了两种特殊的只读内存：纹理内存和常量内存。纹理内存针对空间局部性访问优化，非常适合图像处理等算法；常量内存则通过缓存为所有线程提供对常量数据的高速访问。我们掌握了它们的声明、初始化和在核函数中访问的方法。

最后，我们介绍了CUDA计算能力的概念，理解了它是标识GPU硬件功能和架构的版本号，并知道了如何在编译时指定目标计算能力。

至此，我们已经完成了对CUDA内存子系统主要组成部分的学习。从全局内存、共享内存、寄存器到纹理和常量内存，这些知识将帮助我们更好地优化CUDA程序，充分发挥GPU的并行计算潜力。

011：数据竞争与同步 🚦

概述

在本节课中，我们将要学习并行编程中的核心概念——数据竞争与同步。我们将从数据竞争的定义和必要条件入手，通过多个代码示例深入理解其产生的原因和潜在影响。接着，我们将探讨如何通过不同的同步机制来避免数据竞争，确保并行程序的正确性。课程内容将涵盖互斥、原子操作等基本概念，并通过实际案例（如成绩统计、最短路径计算）来巩固理解。

什么是同步？

在并行编程的视角下，当多个线程协作完成特定任务时，同步指的是这些线程之间的协调。当计算任务要求线程之间必须相互协调，以确保任务能够正确完成且不产生数据竞争时，就需要同步。

上一节我们介绍了GPU内存的各个方面，本节中我们来看看线程间的协调与同步。

何时不需要同步？

并非所有并行操作都需要同步。一个典型的例子是初始化数组。例如，将数组所有元素设置为零。每个线程独立地将其负责的数组元素赋值为零，线程之间无需通信或协调即可完成任务。这是一个“令人尴尬的并行”操作。

然而，在今天的课程中，我们将看到多个需要线程间协调的场景。

同步的成本

在深入同步机制之前，需要理解一个关键点：同步操作的成本非常高。其开销远高于内存访问和计算操作。

以下是各类操作的大致开销排序（从低到高）：

  1. 计算：最快，通常在一个GPU周期内完成。

  2. 内存访问：较慢，访问寄存器、共享内存、全局内存的延迟依次增加（从几个周期到数百个周期）。

  3. 同步：最慢，开销最大。

因此，在设计CUDA程序时，优化顺序应为：

  1. 首先，尝试设计无需同步的算法。

  2. 其次，优化内存访问模式（如合并访问、避免存储体冲突）。

  3. 最后，再考虑优化计算本身（如避免线程发散）。

课程路线图

在接下来的几讲中，我们将系统学习同步相关的所有内容。以下是我们的学习路线：

以下是未来五讲的核心主题：

  • 数据竞争：理解其定义和产生条件。

  • 互斥：确保临界区一次只被一个线程执行。

  • 实现同步的三种主要方式：

    • 原子操作

    • 锁

    • 屏障（例如，我们已在内存章节见过的 __syncthreads()，它用于同步线程块内的线程）

  • 两种重要的并行计算模式：

    • 归约

    • 前缀和

本节课，我们将重点学习数据竞争，并初步了解两个线程间的互斥实现。

课堂练习：识别数据竞争 (1)

让我们通过一个具体的编程问题来理解数据竞争。

任务描述：

我们有一个结构体 Point，包含属性 x 和 y。声明一个包含 n 个 Point 的数组。每个线程处理四个连续的元素（例如，线程0处理元素0-3，线程1处理元素4-7，依此类推）。每个线程需要：

  1. 计算其负责的四个 x 值的平均值。

  2. 检查这四个点的 y 值：

    • 如果任何一个 y 值大于步骤1计算出的 x 平均值，则将该四个点的所有 y 值替换为该平均值。

    • 否则，将这四个 y 值累加到一个全局变量 global_sum 中。

  3. 程序最后需要打印出被设置为平均值的元素数量。

问题：请尝试构思代码实现，并指出代码中可能存在的潜在问题（特别是数据竞争）。

分析与解答：

  • 独立计算部分：每个线程计算自己四个元素的 x 平均值、检查 y 值、以及可能的替换操作，这些都是独立的，不会引发数据竞争。

  • 数据竞争风险点1 - 全局累加：当多个线程同时执行 global_sum = global_sum + y_value 时，会发生数据竞争。这条语句在底层会被分解为“读取 global_sum -> 计算新值 -> 写回 global_sum”三个步骤。如果两个线程几乎同时读取了旧的 global_sum，分别加上自己的 y_value 后写回，那么后写入的结果会覆盖前一个，导致部分加法丢失。

    • 示例：设 global_sum 初值为4。线程A想加2，线程B想加3。正确结果应为9。但如果它们同时读取到4，线程A计算4+2=6并写回，线程B计算4+3=7并写回（覆盖了6），最终结果可能是6或7，而不是9。

  • 数据竞争风险点2 - 全局计数：同样，当多个线程需要递增一个全局计数器 count（用于统计有多少个元素被设置为平均值）时，count++ 操作也会引发类似的数据竞争。

这个练习清晰地展示了当多个线程并发读写共享变量时，数据竞争如何导致程序结果错误。

数据竞争的必要条件

从上面的例子，我们可以总结出数据竞争发生的四个必要条件，必须全部满足才会发生：

  1. 存在多个线程。

  2. 这些线程访问（读或写）同一个共享内存位置。

  3. 至少有一个线程执行的是写操作。

  4. 这些访问是并发进行的（没有强制的时间顺序）。

如果缺少其中任何一个条件，数据竞争就不会发生。

避免数据竞争的方法

既然数据竞争需要四个条件同时成立，那么避免数据竞争的方法就是破坏其中至少一个条件。以下是四种对应策略：

以下是避免数据竞争的四种主要思路：

  • 顺序执行：只使用一个线程，破坏“多个线程”的条件。这牺牲了并行性。

  • 数据私有化：为每个线程创建共享变量的私有副本，让线程操作自己的数据，破坏“共享内存位置”的条件。最后可能需要合并这些私有结果。

  • 使用屏障分离读写：通过屏障同步点，确保所有读操作完成后，再进行写操作（或反之），破坏“并发访问”的条件。例如，先让所有线程并行读取数据，然后同步(__syncthreads)，再让一个线程执行写操作。

  • 实现互斥：确保对于临界区（访问共享资源的代码段），同一时刻只有一个线程能够进入，破坏“并发访问”的条件。这通常通过锁或原子操作实现。

课堂练习：分析互斥代码 (2)

考虑以下两个线程（T1和T2）的代码片段，其中 flag 是一个初始值未知的共享变量，S1 和 S2 代表两个临界区（例如，更新共享变量的操作）。

// 线程 T1
while(flag); // 循环等待，直到 flag 为 0
S1;          // 临界区
flag = 1;

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/f0bef69f03e0d8ea63e07ce3df5f3d4a_124.png)

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/f0bef69f03e0d8ea63e07ce3df5f3d4a_126.png)

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/f0bef69f03e0d8ea63e07ce3df5f3d4a_128.png)

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/f0bef69f03e0d8ea63e07ce3df5f3d4a_130.png)

// 线程 T2
while(!flag); // 循环等待，直到 flag 为 1
S2;           // 临界区
flag = 0;

问题：这段代码是否存在数据竞争？它保证了什么属性？

分析与解答：

  • 是否存在数据竞争？ 从高层逻辑分析，这段代码的设计目的是实现互斥。通过分析 flag 初始值为0或1的各种执行交错情况，可以发现 S1 和 S2 永远不会同时执行。虽然对 flag 的读写是并发的，但简单的内存赋值（如 flag=1）在硬件层面通常是原子的单条指令，其并发执行不会导致数据竞争引发的不确定性错误。因此，这段代码本身没有会导致程序错误的数据竞争。

  • 它保证了什么？

    1. 互斥：它确实保证了 S1 和 S2 这两个临界区不会同时被执行。

    2. 执行顺序：进一步分析可以发现，该代码还隐含地保证了 S2 总是先于 S1 执行。

  • 潜在问题：这段代码存在一个严重缺陷——可能发生死锁或线程饥饿。例如，若 flag 初始值为1，T2会直接执行 S2 并将 flag 设为0。如果T2结束后T1才开始运行，它会将 flag 设为1，然后陷入 while(flag); 的无限循环，因为再也没有线程会将 flag 改回0。

这个练习说明，即使避免了数据竞争，不正确的同步设计也可能导致死锁等其他并发问题。

课堂练习：成绩统计与数据竞争 (3)

任务描述：有80个学生，学号即数组索引。给定一个存储学生分数的数组 marks[]。评分标准为：S(>=90), A(80-89), B(70-79), C(60-69), D(50-59), E(40-49), U(<40)。需要统计每个等级的学生人数。

问题：如何用CUDA内核实现？会存在数据竞争吗？如何避免？

分析与解答：

  • 实现思路1（按学生并行）：启动80个线程（或更多），每个线程处理一个学生的分数。线程根据分数判断等级，然后递增对应等级的全局计数器（例如，count_S++, count_A++）。

    • 数据竞争：存在。如果多个学生属于同一等级，多个线程就会同时尝试递增同一个计数器，导致数据竞争。

    • 避免方法：需要使用原子操作（如 atomicAdd）来保护对每个等级计数器的递增操作。

  • 实现思路2（按等级并行）：启动7个线程（对应7个等级）。每个线程遍历整个80人的分数数组，统计属于自己等级的人数。

    • 数据竞争：不存在。每个线程独立操作自己的私有计数器，最后写入结果。线程间没有共享的写入位置。

    • 缺点：每个线程都要遍历整个数组，计算量是思路1的7倍，当学生数量极大时效率低下。

  • 综合策略：可以结合两种思路。例如，使用多个线程块，每个块内先按思路2进行局部归约（避免块内数据竞争），然后再使用原子操作或更高级的归约方法在块间进行全局累加。

案例分析：并行最短路径算法

任务描述：使用GPU计算从源城市（如图中的A）到所有其他城市的最短路径。假设使用类似Bellman-Ford的算法，在每轮迭代中，每个顶点尝试通过其邻居来松弛到其他顶点的距离。

核心代码逻辑：

__global__ void shortestPathKernel(Graph graph, int* distance) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x; // 线程ID对应顶点ID
    if (tid >= graph.numVertices) return;

    int my_dist = distance[tid];
    for (each neighbor ‘n’ of vertex ‘tid’) {
        int weight = getEdgeWeight(tid, n); // 获取边权
        int alt_dist = my_dist + weight;
        if (alt_dist < distance[n]) { // 如果找到更短路径
            distance[n] = alt_dist;    // 更新邻居的距离
        }
    }
}

问题：找出上述内核代码中至少两处可能发生数据竞争的地方，并分析其影响。

分析与解答：

  1. 竞争点1（影响正确性）：if (alt_dist < distance[n]) 中的读取 distance[n] 与 distance[n] = alt_dist; 中的写入 distance[n] 可能并发发生。

    • 场景：顶点A和顶点D同时试图更新顶点C的距离。C的初始距离为无穷大。A计算的 alt_dist 为3，D计算的为7。两者都读取到 distance[C] = 无穷大，并都判断 alt_dist < 无穷大 成立。如果它们随后并发地执行写操作，最终 distance[C] 的值可能是3或7，取决于谁最后写入。而正确的最短路径是3。这导致了错误的结果。

  2. 竞争点2（不影响最终正确性，但影响效率）：int my_dist = distance[tid]; 中的读取与另一个线程正在执行的 distance[tid] = ... 写入可能并发发生。

    • 场景：顶点A正在更新 distance[B] 为7。同时，顶点B正在读取自己的 distance[B] 以计算其邻居的距离。B可能读到旧的无穷大值，从而基于错误信息进行松弛。然而，在Bellman-Ford的多轮迭代中，下一轮迭代B会读到更新后的正确值7。因此，这不会影响算法的最终正确性，但可能会增加达到收敛所需的迭代轮数，影响性能。

这个案例说明，数据竞争不仅可能导致结果错误，也可能导致性能下降。在并行图算法中，同步至关重要。

总结

本节课中我们一起学习了并行编程中数据竞争与同步的基础知识。

  • 数据竞争 发生在多个线程并发读写共享内存且至少有一个写操作时，它会导致程序行为不确定和结果错误。

  • 我们通过多个实例（累加、计数、成绩统计、最短路径）深入分析了数据竞争的产生场景和后果。

  • 数据竞争有四个必要条件，破坏任一条件即可避免竞争。

  • 避免数据竞争的主要方法包括顺序化、数据私有化、使用屏障和实现互斥。

  • 我们初步探讨了通过标志变量实现简单互斥的方法，并指出了其可能存在的死锁缺陷。

  • 最后，我们以并行最短路径算法为例，分析了实际应用中复杂的数据竞争问题及其对正确性和性能的影响。

理解数据竞争是编写正确、高效CUDA程序的基础。在接下来的课程中，我们将学习更健壮、更高效的同步机制，如原子操作和锁。

012：互斥锁 🔒

概述

在本节课中，我们将学习如何使用控制流和数据流（即不使用原子操作或硬件屏障等特殊指令）来实现两个线程间的互斥锁。我们将分析几种不同的实现方法，探讨它们各自在互斥性、进展性和死锁方面的优缺点，并最终介绍经典的Peterson算法。

上节回顾

在上一节中，我们介绍了同步、互斥锁和临界区的概念。我们通过数据库的例子了解了数据竞争，分析了导致数据竞争的四个必要条件，并探讨了如何通过消除这些条件来避免数据竞争。课程最后，我们分析了一个用于计算城市间最短距离的GPU内核代码，并指出了其中因多个线程同时读写共享数组而可能引发的数据竞争问题。

同步的三种方式

实现同步通常有三种主要方式：

  1. 控制流 + 数据流：仅使用程序逻辑（如循环、条件判断）和普通变量进行同步。

  2. 原子操作：使用硬件支持的原子指令。

  3. 屏障：使用硬件支持的同步屏障。

其中，原子操作和屏障需要硬件支持，而控制流+数据流的方式则完全在软件层面实现。

任务：实现两线程互斥锁

我们的目标是：仅使用控制流和数据流（不使用原子操作、锁或屏障），为两个线程实现互斥锁机制。

以下是几种尝试方案及其分析。

方案一：简单标志法

我们首先尝试一个简单的方案，使用一个共享的布尔标志 flag。

// 线程 T1 的代码
while(!flag); // 等待
S1; // 临界区
flag = false;

// 线程 T2 的代码
while(flag); // 等待
S2; // 临界区
flag = true;

假设：flag 初始值为 false。

分析：

  • 互斥性：❌ 不保证。如果 flag 初始值为 true，两个线程的 while 循环条件会同时为假，导致它们同时进入临界区。

  • 进展性：❌ 不保证。执行顺序被固定：必须先执行 S2，然后才能执行 S1。如果 T2 不执行，T1 将永远等待。

  • 死锁：✅ 不会发生。因为 flag 的值非真即假，不会导致两个线程都卡在 while 循环中。

结论：此方案存在严重缺陷，无法保证互斥性。

方案二：对称等待法

我们尝试修改方案一，让两个线程对称地等待对方。

// 线程 T1 的代码
while(!flag);
S1;
flag = false;

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/849f9634ffa9fef9cf255d58f370c34c_85.png)

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/849f9634ffa9fef9cf255d58f370c34c_87.png)

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/849f9634ffa9fef9cf255d58f370c34c_89.png)

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/849f9634ffa9fef9cf255d58f370c34c_91.png)

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/849f9634ffa9fef9cf255d58f370c34c_93.png)

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/849f9634ffa9fef9cf255d58f370c34c_95.png)

// 线程 T2 的代码
while(flag);
S2;
flag = true;

假设：flag 初始值为 false。flag 被声明为 volatile，确保读写直接作用于主存，避免缓存一致性问题。

分析：

  • 互斥性：✅ 保证。flag 在同一时刻只能有一个值（true 或 false），因此只有一个线程能通过 while 循环。

  • 进展性：❌ 不保证。如果 flag 初始为 true 且 T1 永远不想进入临界区，那么想进入临界区的 T2 将永远在 while(flag); 处等待。

  • 死锁：❌ 可能发生。考虑以下执行顺序：

    1. T1 执行 flag = false;（假设之前已执行完 S1）。

    2. T2 执行 flag = true;。

    3. T1 和 T2 几乎同时到达各自的 while 循环。

    此时，flag 的最终值为 true（后写入者生效）。T1 看到 while(!true) 为假，进入 S1。T2 看到 while(true) 为真，陷入无限等待。这虽然不是一个典型的“双方互相等待”的死锁，但导致了 T2 饥饿，从效果上看类似于死锁。更精确地说，在并发执行时，可能导致一个线程永远等待。

结论：此方案保证了互斥性，但无法保证进展性，且在某些执行顺序下会导致线程饥饿/死锁。

方案三：数组声明法 (Lock Version 1)

我们尝试实现通用的 lock() 和 unlock() 函数，使用一个标志数组 flag[2]。

volatile bool flag[2] = {false, false};

void lock(int tid) {
    int me = tid; // 0 或 1
    int other = 1 - me;
    flag[me] = true; // 声明自己希望进入临界区
    while(flag[other] == true); // 如果对方也希望进入，则等待
}

void unlock(int tid) {
    flag[tid] = false; // 声明自己离开临界区
}

分析：

  • 互斥性：✅ 保证。只有当对方线程的标志为 false 时，当前线程才能进入临界区。

  • 进展性：✅ 保证。即使一个线程不参与（其 flag 为 false），另一个线程也能顺利进入临界区。

  • 死锁：❌ 可能发生。考虑以下执行顺序：

    1. T0 执行 flag[0] = true;

    2. T1 执行 flag[1] = true;

    3. 现在 flag[0] 和 flag[1] 都为 true。

    4. T0 执行 while(flag[1] == true); 陷入等待。

    5. T1 执行 while(flag[0] == true); 陷入等待。

    双方都在等待对方将标志置为 false，形成死锁。

结论：此方案保证了互斥性和进展性，但存在死锁风险。

方案四：牺牲变量法 (Lock Version 2)

我们引入一个共享的 victim 变量来打破对称性，避免死锁。

volatile int victim;

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/849f9634ffa9fef9cf255d58f370c34c_140.png)

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/849f9634ffa9fef9cf255d58f370c34c_142.png)

void lock(int tid) {
    int me = tid;
    victim = me; // 主动“牺牲”自己，让对方优先
    while(victim == me); // 如果自己仍然是“牺牲者”，则等待
}

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/849f9634ffa9fef9cf255d58f370c34c_144.png)

void unlock(int tid) {
    // 解锁操作简单，无需修改 victim
    // victim 的值会在下一次 lock 竞争时被覆盖
}

分析：

  • 互斥性：✅ 保证。victim 只能为 0 或 1。最后执行 victim = me; 的线程会成为真正的“牺牲者”，在其 while 循环中等待。

  • 进展性：❌ 不保证。如果只有一个线程（如 T0）调用 lock()，它会执行 victim = 0;，然后发现 victim == 0 为真，从而在 while 循环中无限等待，即使没有竞争者。

  • 死锁：✅ 不会发生。victim 的最终值会确定一个唯一的等待线程，另一个线程总能进入临界区。

结论：此方案保证了互斥性且无死锁，但无法保证进展性（单个线程无法独立工作）。

最终方案：Peterson算法 (Lock Version 3)

Peterson算法巧妙地结合了声明意向 (flag[]) 和礼貌谦让 (victim) 两种机制。

volatile bool flag[2] = {false, false};
volatile int victim;

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/849f9634ffa9fef9cf255d58f370c34c_158.png)

void lock(int tid) {
    int me = tid;
    int other = 1 - me;
    flag[me] = true; // 步骤1：声明自己希望进入
    victim = me;      // 步骤2：主动让对方优先
    // 步骤3：等待条件：当对方想进入，并且自己是被礼让者时，继续等待
    while(flag[other] == true && victim == me);
}

void unlock(int tid) {
    flag[tid] = false; // 离开临界区，取消声明
}

分析：

  • 互斥性：✅ 保证。进入临界区的条件是：flag[other] == false（对方不想进）或 victim == other（对方是牺牲者）。两个线程不可能同时使对方的条件为假。

  • 进展性：✅ 保证。如果只有一个线程想进入（flag[other] == false），它将立即通过 while 循环。即使两个线程都想进入，victim 变量也能确保其中之一（最后设置 victim 的线程）会等待，另一个则进入。

  • 死锁：✅ 不会发生。victim 变量确保了至少有一个线程能通过 while 循环。不会出现双方无限等待的情况。

核心思想：

  1. flag[me] = true：大声宣布“我想进去”。

  2. victim = me：礼貌地说“您先请”。

  3. while(...)：等待的条件是“您想进，并且您还让我先请”，那我就继续等。否则（您不想进，或者该我进了），我就进去。

总结

本节课我们一起学习了如何仅用控制流和数据流实现两线程互斥锁。我们经历了多个方案的迭代：

  1. 简单标志法：互斥性和进展性均无法保证。

  2. 对称等待法：保证互斥性，但无进展性，可能饥饿。

  3. 数组声明法 (V1)：保证互斥性和进展性，但存在死锁风险。

  4. 牺牲变量法 (V2)：保证互斥性且无死锁，但无进展性。

  5. Peterson算法 (V3)：完美地同时保证了互斥性、进展性和无死锁。

Peterson算法是一个经典的软件互斥解决方案，它优雅地展示了如何通过共享变量和程序逻辑来解决复杂的同步问题。需要注意的是，它只适用于两个线程的场景。在下一节课中，我们将探讨如何将互斥锁的概念扩展到多个线程。

013：互斥锁（续）及CUDA中的原子操作 🔒⚛️

概述

在本节课中，我们将要学习如何为多个线程（n个线程）实现互斥锁，并深入了解CUDA中的原子操作。我们将从适用于多线程的“面包店算法”开始，然后探讨原子操作的概念、其与互斥锁的区别，以及几种常见的原子指令。

面包店算法 🍞

上一节我们介绍了如何为两个线程实现互斥锁。本节中我们来看看如何为n个线程实现互斥锁，这需要使用“面包店算法”。

面包店算法由Leslie Lamport提出，它模拟了面包店或服务台取号排队的机制。每个线程在进入临界区前会获取一个“号码”（标签），然后根据号码大小顺序进入临界区，从而保证互斥和先进先出（FIFO）的顺序。

以下是该算法的核心伪代码：

bool flag[n]; // 初始化为 false
int label[n]; // 初始化为 0

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/40d1258203816d1c71a3a1177d535055_39.png)

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/40d1258203816d1c71a3a1177d535055_41.png)

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/40d1258203816d1c71a3a1177d535055_43.png)

void lock(int me) {
    flag[me] = true; // 表示线程 me 希望进入临界区
    label[me] = 1 + max(label[0], ..., label[n-1]); // 获取当前最大标签值加1
    while ( (存在 k != me 且 flag[k] == true) &&
            ( (label[k], k) < (label[me], me) ) ) {
        // 等待：如果存在其他希望进入且“号码”更小的线程，则循环等待
    }
}

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/40d1258203816d1c71a3a1177d535055_45.png)

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/40d1258203816d1c71a3a1177d535055_46.png)

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/40d1258203816d1c71a3a1177d535055_48.png)

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/40d1258203816d1c71a3a1177d535055_50.png)

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/40d1258203816d1c71a3a1177d535055_52.png)

void unlock(int me) {
    flag[me] = false; // 表示线程 me 已离开临界区
}

核心概念解释：

  • flag[me]: 表示线程me是否希望进入临界区。

  • label[me]: 线程me的“排队号码”。

  • 字典序比较 (label[k], k) < (label[me], me): 先比较标签label，如果标签相同，则比较线程ID k 和 me。这确保了当多个线程获得相同标签时，线程ID较小的线程优先。

算法要点与讨论

  1. 标签生成非原子性：label[me] = 1 + max(...) 这一步不是原子操作。如果多个线程同时计算最大值并获取标签，它们可能得到相同的标签值。

  2. 字典序比较的必要性：如果只比较标签 (label[k] < label[me])，当多个线程获得相同标签时，它们可能同时进入临界区，破坏互斥性。因此需要引入线程ID作为决胜局。

  3. 内存可见性：flag和label数组应声明为volatile，或在支持缓存一致性的系统中确保写入能及时对其他线程可见，否则可能因缓存导致线程读取到旧值。

  4. 在GPU上的局限性：面包店算法在GPU的同一个Warp内执行会导致死锁。因为Warp内的32个线程是锁步（SIMD）执行的。如果所有线程都执行lock函数，它们会同时设置flag、计算相同的label，然后大部分线程会在while循环中等待。而唯一满足条件（如线程ID最小）的线程也因为Warp锁步执行而无法前进，从而导致整个Warp死锁。

从互斥锁到原子操作 🔄

上一节我们探讨了基于软件的互斥锁实现。本节中我们来看看GPU上更常用、更底层的同步原语——原子操作。

为什么需要原子操作？

考虑一个简单的例子：两个线程（来自不同的线程块）同时对一个全局内存变量x进行自增操作 x++。

// 假设 x 指向全局内存，初始值为0
// 内核启动配置: <<<2, 1>>> 表示2个线程块，每个块1个线程
__global__ void kernel(int *x) {
    x[0]++; // 这不是原子操作！
}

x++ 在底层可能被编译成三条机器指令：

  1. LOAD: 将x[0]的值从内存加载到寄存器R。

  2. ADD: 将寄存器R的值加1。

  3. STORE: 将寄存器R的值存回x[0]所在的内存。

由于两个线程的执行可能交错进行，最终x的值可能是1或2，而不是预期的2。这就是数据竞争。

原子操作可以解决这个问题。它将“读取-修改-写入”这一系列操作打包成一个不可分割的、连续的整体操作。在执行原子操作期间，其他线程无法访问目标内存位置，从而保证了结果的正确性。

原子操作 vs. 互斥锁

  • 互斥锁：是一种更高级别的软件抽象，用于保护一段代码（临界区）。线程在进入前获取锁，退出时释放锁。在GPU上，由于线程数量极多，锁的开销很大，容易导致性能下降和死锁。

  • 原子操作：是硬件支持的底层指令，直接作用于单个内存位置（如加、减、比较交换等）。它更轻量，适用于简单的数据更新场景。在GPU编程中，原子操作比互斥锁更常用。

原子操作的特性与应用场景

  • 硬件支持：原子操作需要硬件支持，CUDA GPU提供了丰富的原子指令。

  • 内存范围：原子操作可用于全局内存和共享内存。

  • 常见指令：

    • atomicAdd(address, val): 原子加法。

    • atomicCAS(address, compare, val): 原子比较并交换（可实现锁）。

    • atomicMin/atomicMax(address, val): 原子求最小/最大值。

  • 应用场景：统计直方图、求和、求极值、实现无锁数据结构等。

现实世界类比：银行ATM取款 💳

假设一个联合账户有1000元存款，双胞胎兄弟同时从两台ATM机各取1000元。ATM机的逻辑是：

  1. 检查余额 >= 取款金额。

  2. 如果满足，则扣款并吐出现金。

如果步骤1和2不是原子操作，可能发生以下交错执行：

  • 两台ATM都检查余额（1000 >= 1000，通过）。

  • 两台ATM都扣款（1000 - 1000 = 0）。

  • 两台ATM都吐出现金1000元。

结果：账户余额为0，但用户总共取走了2000元，银行损失1000元。

原子操作可以确保“检查余额”和“扣款”作为一个不可分割的单元执行，从而避免此问题。

CUDA原子加法示例 ➕

为了解决之前x++的数据竞争问题，我们可以使用CUDA的atomicAdd函数。

__global__ void safe_kernel(int *x) {
    atomicAdd(&x[0], 1); // 原子地将 x[0] 的值增加1
}

函数原型：int atomicAdd(int* address, int val);

  • address: 要进行原子操作的目标内存地址。

  • val: 要加上的值。

  • 返回值：返回该地址在操作前的旧值。

使用atomicAdd后，无论启动多少个线程块和线程，也无论执行顺序如何，最终x[0]的值都将是所有线程执行次数的总和（例如，启动配置为<<<K1, K2>>>，则最终值为 K1 * K2）。

性能注意：原子操作会导致对同一内存地址的访问串行化，即多个线程必须排队依次执行该操作。这会成为性能瓶颈，应尽量避免在热点循环中频繁使用原子操作。

总结

本节课中我们一起学习了：

  1. 面包店算法：一种为n个线程实现互斥锁的算法，通过取号和字典序比较来保证互斥与顺序。但它在GPU的Warp内执行会导致死锁。

  2. 原子操作：硬件支持的、不可分割的单一内存操作。它解决了多线程下的数据竞争问题，是GPU上重要的同步机制。

  3. 原子操作与互斥锁的区别：原子操作更底层、更轻量，适用于简单的数据更新；互斥锁更高级，用于保护代码段，但在GPU上开销大。

  4. atomicAdd指令：CUDA中用于原子加法的函数，它能确保计数的正确性，但会引入串行化开销。

下一节课，我们将深入探讨更强大的atomicCAS（比较并交换）指令，并学习如何使用它来实现自旋锁。

014：CUDA中的多种原子函数与屏障 🚀

在本节课中，我们将学习如何使用CUDA中的原子函数来解决数据竞争问题，并深入探讨屏障（Barrier）这一重要的同步机制。我们将从回顾单源最短路径算法中的数据竞争问题开始，然后学习如何使用原子操作（如 atomicMin 和 atomicCAS）来确保代码的正确性。接着，我们将了解如何利用 atomicCAS 实现锁（Lock）和单线程执行区（Single Section）。最后，我们将正式定义屏障，并介绍CUDA中不同级别的屏障同步。

回顾：单源最短路径算法中的数据竞争 🔄

在之前的课程中，我们学习了单源最短路径算法。我们注意到，在并行执行时，代码中存在数据竞争问题，这可能导致计算出错误的距离值。

问题出现在以下代码段中：

if (alt_dist < dist[neighbor]) {
    dist[neighbor] = alt_dist;
}

当多个线程同时读取和更新同一个邻居节点（例如节点C）的距离时，可能会发生数据竞争。例如，线程A和线程B可能都读取到节点C的初始距离为无穷大，然后分别尝试将其更新为3和6。最后写入的值（可能是6）将覆盖掉更短的正确距离（3），从而导致错误。

上一节我们介绍了数据竞争的问题，本节中我们来看看如何使用原子操作来解决它。

使用原子操作解决数据竞争 ⚛️

为了消除数据竞争，我们需要确保“读取-比较-写入”这一系列操作是原子性的，即不可分割的。CUDA提供了 atomicMin 函数，它正是用于此目的。

我们可以将上述易出错的代码替换为：

atomicMin(&dist[neighbor], alt_dist);

atomicMin 函数会原子性地比较 dist[neighbor] 的当前值与 alt_dist 的值，并将 dist[neighbor] 更新为两者中的较小值。这样，无论线程执行顺序如何，每个节点的距离都会被正确地更新为所有可能路径中的最小值。

深入原子比较与交换（atomicCAS） 🔧

atomicCAS（Compare And Swap）是一个功能强大的原子操作，其语法如下：

int atomicCAS(int* address, int compare, int val);

它的工作流程是：

  1. 原子性地读取 address 指针处的值。

  2. 将该值与 compare 参数进行比较。

  3. 如果相等，则将 address 处的值设置为 val。

  4. 无论是否交换，函数都返回 address 处原来的旧值。

atomicCAS 有两个典型应用：

  1. 实现锁（Lock），用于保护临界区。

  2. 实现单线程执行区（Single Section），确保只有任意一个线程执行特定代码块。

使用 atomicCAS 实现锁 🔒

简单的 atomicCAS 调用本身并不能实现锁。我们需要将其与循环结合，构建一个自旋锁。

以下是使用 atomicCAS 在GPU上实现一个能正确处理线程束（Warp）内同步的锁的示例：

// 假设 lock_var 初始化为 0
__shared__ int lock_var;

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/b570460d95f5d51d21e4c35ea3a489a9_104.png)

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/b570460d95f5d51d21e4c35ea3a489a9_106.png)

// 尝试获取锁
int old;
do {
    old = atomicCAS(&lock_var, 0, 1); // 尝试将锁从0设为1
} while (old != 0); // 如果old不是0，说明锁已被占用，继续循环

// 临界区代码
// ...

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/b570460d95f5d51d21e4c35ea3a489a9_108.png)

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/b570460d95f5d51d21e4c35ea3a489a9_110.png)

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/b570460d95f5d51d21e4c35ea3a489a9_112.png)

// 释放锁
lock_var = 0;

重要说明：这个基础实现在CPU或不同Warp的线程间可以工作，但在同一个Warp的线程间会导致死锁。因为Warp内的线程是同步执行的，一个线程获得锁后，其他线程会在循环中等待，而获得锁的线程也在等待循环中的其他线程，形成死锁。

为了解决Warp内死锁问题，需要结合条件判断和线程掩码，确保同一时刻只有一个Warp线程进入临界区，而其他线程被正确屏蔽。这通常涉及更复杂的逻辑。

使用 atomicCAS 实现单线程执行区 🎯

单线程执行区要求只有任意一个线程执行某段代码，其他线程无需等待，直接跳过。

使用 atomicCAS 可以轻松实现：

// 假设 single_flag 初始化为 0
__shared__ int single_flag;

![](https://github.com/OpenDocCN/dsai-notes-pt1-zh/raw/master/docs/nv-cuda-prll-prog/img/b570460d95f5d51d21e4c35ea3a489a9_130.png)

int old = atomicCAS(&single_flag, 0, 1); // 尝试将标志从0设为1
if (old == 0) {
    // 只有一个线程会成功进入此区域
    // 单线程执行区代码
    // ...
}
// 其他线程直接跳过，继续执行后续代码

在这个实现中，single_flag 在首次被成功后就不再是0，因此后续线程的 atomicCAS 操作都会失败，从而不会进入if块。线程执行完单线程区后，也无需重置 single_flag。

理解屏障（Barrier） 🚧

屏障是一种同步点，要求所有参与线程都必须到达该点后，才能有任何线程继续执行后续代码。

类比：想象一群朋友约定去踢足球。大家从各自家中（异步）出发，到一个集合点汇合（屏障）。只有所有人都到达集合点后，大家才一起前往足球场。

在CUDA中，存在不同级别的屏障：

  1. 隐式全局屏障：内核（Kernel）的结束点本身就是一个隐式屏障，主机代码会等待所有GPU线程完成。

  2. 线程块内屏障：使用 __syncthreads() 函数。它同步同一个线程块内的所有线程。

  3. 全局屏障：从CUDA 9开始，支持 __syncgrid()（或类似机制，具体API需查阅文档），可以同步网格（Grid）中所有线程块内的线程。

  4. 线程束内屏障：对于Warp内的线程，由于它们以锁步方式执行指令，因此本身就存在隐式的执行屏障，通常不需要显式同步。

总结 📚

本节课我们一起学习了以下核心内容：

  • 原子函数：我们回顾了使用 atomicMin 解决单源最短路径算法中的数据竞争问题。

  • atomicCAS：深入学习了其原理，并探讨了如何用它来实现锁（解决临界区互斥访问）和单线程执行区。特别要注意在GPU上实现锁时需考虑Warp内线程的死锁问题。

  • 屏障：我们正式定义了屏障作为同步点的概念，并介绍了CUDA中不同层次的同步机制，包括隐式内核屏障、__syncthreads() 线程块内屏障以及更新的全局屏障支持。

通过掌握原子操作和屏障，你可以编写出正确、高效且同步良好的CUDA并行程序，避免数据竞争和协调线程间的执行顺序。

015：屏障、归约与CUDA中的前缀和 🚀

概述

在本节课中，我们将要学习CUDA编程中的三个核心概念：屏障（Barrier）、归约（Reduction）和前缀和（Prefix Sum）。我们将了解屏障如何实现线程间的同步，如何利用归约高效地计算数组的总和，以及如何实现和应用前缀和算法。

屏障（Barrier）回顾与深入

上一节我们介绍了原子操作，本节我们继续探讨屏障这一同步机制。

屏障是一个程序点，要求所有线程都必须到达该点后，任何线程才能继续执行后续指令。在CUDA中，内核（Kernel）的结束处对所有GPU线程来说是一个隐式的全局屏障。从CUDA 9开始，grid.sync()函数可以作为跨GPU所有线程的全局屏障。而__syncthreads()函数则充当线程块（Thread Block）内所有线程的屏障。

屏障示例分析

以下是使用__syncthreads()的一个示例代码：

__global__ void kernel(unsigned int* vector, unsigned int vector_size) {
    unsigned int id = threadIdx.x;
    vector[id] = id; // 每个线程根据其ID初始化数组元素
    __syncthreads(); // 线程块内屏障
    if (id < vector_size - 1 && vector[id + 1] != id + 1) {
        printf("__syncthreads doesn't work\n");
    }
}

以下是代码逻辑的逐步分析：

  1. 每个线程根据其threadIdx.x（即ID）为数组vector的对应位置赋值。

  2. 调用__syncthreads()，确保线程块内所有线程都完成了初始化操作。

  3. 每个线程（除了最后一个）检查其下一个元素的值是否等于id + 1。

潜在问题：如果数组大小vector_size超过了一个线程块的最大线程数（例如1024），则需要启动多个线程块。__syncthreads()仅同步同一个线程块内的线程。因此，可能出现以下情况：

  • 第一个线程块（线程ID 0-1023）的线程完成了初始化，并通过了屏障。

  • 第二个线程块（线程ID 1024-2047）的线程仍在执行初始化。

  • 此时，第一个线程块中ID为1023的线程执行检查：它试图访问vector[1024]（即id+1），但这个位置的值尚未被第二个线程块的线程初始化。这可能导致检查条件成立，从而打印错误信息。

解决方案：如果使用grid.sync()替代__syncthreads()，则可以确保所有GPU线程（跨所有线程块）都完成初始化后才进行检查，从而避免此问题。

屏障的数据同步作用

__syncthreads()不仅提供控制流同步（所有线程到达后才继续），还提供数据同步。它确保在屏障点之前，线程块内所有线程对内存的写入操作，对屏障点之后所有线程的读取操作是可见的。

这是通过内存栅栏（Memory Fence） 操作实现的。内存栅栏强制将线程的写入刷新到主内存，并确保后续读取从主内存获取最新值，从而保证内存操作的顺序性和可见性。

  • __threadfence_block(): 在__syncthreads()内部隐式调用，确保线程块内的数据同步。

  • __threadfence(): 确保GPU内所有线程的数据同步。

归约（Reduction）

上一节我们介绍了屏障，本节中我们来看看归约。归约是一种将一组数据（例如一个数组）聚合为单个值或少量值的操作，例如计算总和、最大值、最小值等。

为什么需要归约？

考虑计算一个包含N个元素的数组的总和。一种简单的方法是使用原子操作：

__global__ void atomic_sum(int* array, int* sum, int n) {
    int id = threadIdx.x;
    if (id < n) {
        atomicAdd(sum, array[id]); // 原子加法
    }
}

这种方法的问题是严重的顺序性：尽管有多个线程，但atomicAdd操作是串行的，同一时刻只有一个线程能成功更新sum。这无法充分利用GPU的并行能力，性能甚至可能差于CPU串行代码。

归约的基本思想

归约通过树形结构并行地合并数据。其核心思想是：在每一步，将数据两两分组进行合并，使数据量减半，经过 log₂N 步后得到最终结果。

操作前提：归约操作（如加法、求最大值）必须是可结合（Associative） 的。即 (a ⊕ b) ⊕ c = a ⊕ (b ⊕ c)。加法、乘法、求最大值、最小值、按位与/或/异或等操作满足结合律。而减法、除法则不满足。

归约算法实现

以下是一种常见的归约实现方式（假设数组大小为2的幂，且只使用一个线程块）：

__global__ void reduction_sum(int* input, int* output, int n) {
    // 假设input数组已在GPU上，且n为2的幂
    // 使用共享内存或全局内存进行原地操作
    for (int offset = n / 2; offset > 0; offset >>= 1) {
        int id = threadIdx.x;
        if (id < offset) {
            input[id] += input[id + offset]; // 将相隔offset的元素相加
        }
        __syncthreads(); // 每一步都需要同步
    }
    if (threadIdx.x == 0) {
        *output = input[0]; // 最终结果在input[0]
    }
}

算法步骤：

  1. 初始时，有N个元素，使用 N/2 个线程。

  2. 每个线程将位置id的元素与位置id + offset的元素相加，结果存回id位置。offset初始为N/2。

  3. 调用__syncthreads()确保本步所有加法完成。

  4. offset减半（offset >>= 1），重复步骤2-3，直到offset为0。

  5. 此时，input[0]中存储的就是所有元素的总和。

图解：对于数组 [2, 3, 5, 6, 1, 9, 4, 7]

  • 第1步 (offset=4): [2+1, 3+9, 5+4, 6+7, 1, 9, 4, 7] -> [3, 12, 9, 13, ...]

  • 第2步 (offset=2): [3+9, 12+13, 9, 13, ...] -> [12, 25, ...]

  • 第3步 (offset=1): [12+25, ...] -> [37, ...]

结果37即为总和。

优势：

  • 并行度高：每一步都使用大量线程。

  • 步数少：仅需 log₂N 步，而非串行的 N 步。

  • 无需原子操作：通过巧妙的索引计算和同步避免数据竞争。

posted @ 2026-10-17 13:17  绝不原创的飞龙  阅读(6)  评论(0)    收藏  举报