nv-cuda-prll-prog-merge-2
Nvidia CUDA 并行编程笔记（三）
注意事项：
-
需要处理数组大小非2的幂的情况（例如，在数组末尾填充0）。
-
当使用多个线程块时，需要先在各块内进行归约，然后再对块的结果进行归约，并可能使用
grid.sync()或启动多个内核。
前缀和（Prefix Sum）
在学习了归约之后，我们来看一个与之相关但功能不同的算法——前缀和，也称为扫描（Scan）。
什么是前缀和？
对于一个输入数组 [x₀, x₁, x₂, ..., xₙ₋₁]，其前缀和输出数组 [y₀, y₁, y₂, ..., yₙ₋₁] 定义为：
-
y₀ = x₀ -
y₁ = x₀ + x₁ -
y₂ = x₀ + x₁ + x₂ -
...
-
yₙ₋₁ = x₀ + x₁ + ... + xₙ₋₁
即，输出数组的每个元素是输入数组从起始到当前位置的累积和。
前缀和的应用
一个典型应用是工作项分配。假设有多个线程，每个线程需要向一个全局工作列表（数组）中推送不同数量的数据项。我们想知道每个线程推送的数据在全局数组中的起始和结束索引。
例如：
-
线程0推送4个项。
-
线程1推送3个项。
-
线程2推送9个项。
-
线程3推送5个项。
我们可以构建一个计数数组 counts = [4, 3, 9, 5]。计算其前缀和（从0开始）得到 prefix = [0, 4, 7, 16]。
-
线程0的项位于索引
prefix[0]到prefix[1]-1，即[0, 3]。 -
线程1的项位于索引
prefix[1]到prefix[2]-1，即[4, 6]。 -
以此类推。前缀和提供了每个线程数据在全局数组中的“偏移量”。
前缀和算法实现（Blelloch Scan）
以下是一种高效的前缀和并行算法实现（上行扫描）：
__global__ void prefix_sum(int* data, int n) {
// 假设 data 是全局数组，n 为2的幂，且只使用一个线程块
// 上行（Reduce）阶段：构建求和树
for (int offset = 1; offset < n; offset <<= 1) {
int id = threadIdx.x;
int temp = 0;
if (id >= offset) {
temp = data[id - offset]; // 读取前一个偏移量的值
}
__syncthreads(); // 分离读和写，避免数据竞争
if (id >= offset) {
data[id] += temp; // 累加到当前位置
}
__syncthreads(); // 同步本次迭代
}
// 此时 data 数组的最后一个元素 data[n-1] 是总和
// 下行（Downsweep）阶段（此处省略）：将部分和分散，得到最终前缀和
// ... (需要额外的步骤将 exclusive scan 转换为 inclusive scan，或反之)
}
算法逻辑（以上行阶段为例）：
-
offset从1开始，每次迭代翻倍（1, 2, 4, ...）。 -
对于每个
offset，每个线程（除了前offset个）读取当前位置向前offset步的值。 -
使用
__syncthreads()确保所有读取操作在写入开始前完成，这是避免数据竞争的关键。 -
将读取的值累加到当前元素。
-
再次同步，确保本步所有计算完成，再进行下一步。
图解（简化上行阶段）：
初始: [1, 2, 3, 4, 5, 6, 7, 8]
offset=1: [1, 1+2, 2+3, 3+4, 4+5, 5+6, 6+7, 7+8] -> [1, 3, 5, 7, 9, 11, 13, 15]
offset=2: [1, 3, 1+5, 3+7, 5+9, 7+11, 9+13, 11+15] -> [1, 3, 6, 10, 14, 18, 22, 26]
offset=4: [1, 3, 6, 10, 1+14, 3+18, 6+22, 10+26] -> [1, 3, 6, 10, 15, 21, 28, 36]
最终data[7]=36是总和。要得到完整的前缀和数组[1,3,6,10,15,21,28,36]，还需要一个下行阶段。
关键点：
-
数据竞争：在读取
data[id - offset]和写入data[id]时，必须用屏障分离，防止其他线程同时修改正在读取的数据。 -
屏障位置：代码中使用了两个
__syncthreads()。第一个确保“读”在“写”之前完成；第二个确保本步所有计算完成，属于迭代间的控制同步。 -
复杂度：同样为 O(log N) 步。
其他应用
前缀和不仅用于上述偏移计算，还广泛应用于：
-
流压缩（Stream Compaction）：过滤数组中的有效元素。
-
基数排序（Radix Sort）的关键步骤。
-
构建数据结构，如芬威克树（Fenwick Tree）。
-
计算直方图（Histogram）的累积分布。
总结
本节课中我们一起学习了CUDA编程中三个重要的高级主题：
-
屏障：深入理解了
__syncthreads()和grid.sync()的同步范围，并通过例子认识到屏障选择不当可能导致逻辑错误。同时，明白了屏障兼具控制同步和数据同步（通过内存栅栏）的功能。 -
归约：掌握了将大规模数据并行聚合为单一值的高效方法。理解了归约的树形结构思想，学习了其并行实现代码，并认识到其相对于原子操作的巨大性能优势。
-
前缀和：学习了前缀和的概念及其在并行计算中的关键应用（如工作分配）。分析了一种并行前缀和算法的实现，特别注意了其中使用屏障来避免数据竞争的技巧。
这些算法是许多并行计算应用的基础构件，理解它们对于编写高效、正确的CUDA程序至关重要。
016：CUDA中的函数 🚀
概述
在本节课中，我们将要学习CUDA中不同类型的函数。到目前为止，我们只见过一种使用 __global__ 关键字定义的函数，我们称之为内核。但CUDA还提供了其他几种定义函数的方式。我们将详细探讨这些不同类型的函数，并通过多个示例来理解它们的设计理念和具体用法。如果时间允许，我们还将简要介绍CUDA的并行算法库——Thrust。
函数类型介绍
上一节我们介绍了CUDA编程的基本概念，本节中我们来看看CUDA中定义函数的几种不同方式。
在CUDA中，主要有三种方式声明函数，每种方式决定了函数在哪里被调用以及在哪里执行。
-
__global__函数（内核）-
使用
__global__关键字声明。 -
从主机（CPU）调用。
-
在设备（GPU）上执行。
-
必须返回
void类型。
-
-
__device__函数（设备函数）-
使用
__device__关键字声明。 -
从设备（GPU）调用。
-
在设备（GPU）上执行。
-
通常由内核函数调用。
-
-
__host__函数（主机函数）-
使用
__host__关键字声明。 -
从主机（CPU）调用。
-
在主机（CPU）上执行。
-
这与普通的C++函数行为一致。
__host__关键字的主要用途是与__device__结合，定义可在主机和设备上执行的函数。
-
一个程序可以包含多种类型的函数，并且同一种函数可以被多次调用。
函数调用规则与示例
理解了基本类型后，我们通过一个具体的代码示例来深入理解这些函数之间的调用关系。
以下是示例代码的第一部分，展示了四种不同方式声明的函数：
// 一个既可在主机运行也可在设备运行的函数
__host__ __device__ void dh_func() {
printf("I can run on both CPU and GPU\n");
}
// 一个设备函数
__device__ void d_func() {
dh_func(); // 设备函数可以调用主机-设备函数
}
// 一个内核函数
__global__ void d_kernel() {
dh_func(); // 内核可以调用主机-设备函数
d_func(); // 内核可以调用设备函数
}
// 一个主机函数
__host__ void host_func() {
dh_func(); // 主机函数可以调用主机-设备函数
}
以下是代码的第二部分，即主函数：
int main() {
// 从主机（CPU）调用内核
d_kernel<<<1, 1>>>();
cudaDeviceSynchronize();
// 从主机调用主机函数
host_func();
// 从主机调用主机-设备函数
dh_func();
// 以下调用会导致编译错误，因为设备函数只能从设备调用
// d_func();
return 0;
}
基于以上代码，我们可以总结出函数调用的核心规则：
-
__device__函数：只能被__global__内核或其他__device__函数调用。 -
__global__内核：只能从主机（CPU）调用。 -
__host__ __device__函数：可以被主机或设备代码调用，但其函数体内的代码必须是主机和设备都能执行的“交集”。
为了更清晰地展示，以下是可能的函数调用关系图：
-
main->d_kernel(有效) -
main->host_func(有效) -
main->dh_func(有效) -
host_func->dh_func(有效) -
d_kernel->dh_func(有效) -
d_kernel->d_func(有效) -
d_func->dh_func(有效)
重要限制：__host__ __device__ 函数内部不能调用纯 __device__ 函数。因为 __host__ __device__ 函数可能从CPU被调用，而 __device__ 函数只能在GPU上执行，这会导致运行时错误。同理，其内部也不能使用 threadIdx.x 或 __syncthreads() 等GPU特有的变量或函数。
统一内存与函数
上一节我们探讨了函数调用的基本规则，本节中我们来看看函数如何与一种特殊的内存——统一内存进行交互。
统一内存（使用 cudaMallocManaged 分配）提供了一种在CPU和GPU之间自动同步的数据视图。这对于 __host__ __device__ 函数特别有用，因为它们可能从两端访问同一数据。
请看以下示例：
__host__ __device__ void func(int *counter) {
(*counter)++;
printf("Counter: %d\n", *counter);
}
__global__ void print_k(int *counter) {
func(counter); // 在GPU上调用func
}
int main() {
int *counter;
// 分配统一内存
cudaMallocManaged(&counter, sizeof(int));
*counter = 0; // 在CPU上初始化
printf("Main: %d\n", *counter); // 输出: 0
print_k<<<1, 1>>>(counter); // 启动内核
cudaDeviceSynchronize(); // 等待GPU完成
func(counter); // 在CPU上再次调用func
printf("Main: %d\n", *counter); // 输出: 2
cudaFree(counter);
return 0;
}
程序输出：
Main: 0
Counter: 1
Counter: 2
Main: 2
解释：
-
计数器在CPU上初始化为0。
-
内核
print_k在GPU上执行func，将计数器增加到1并打印。 -
由于是统一内存，GPU的修改对CPU可见。
-
CPU再次调用
func，此时看到的值是1，将其增加到2并打印。 -
最终在
main中打印的值为2。
这个例子展示了统一内存如何简化CPU和GPU之间的数据共享。如果使用普通的 cudaMalloc（仅在GPU分配）或 malloc（仅在CPU分配），并尝试从另一端访问，则会导致段错误或读取垃圾值。
实践：数组增量示例
理论需要结合实践，本节我们通过一个具体的编程任务来巩固所学知识：编写一个既能被CPU串行执行，也能被GPU并行执行的数组元素自增函数。
目标：实现一个函数，对数组的所有元素执行 arr[i]++ 操作。
初始方案（CPU串行，GPU单线程串行）：
__host__ __device__ void func(int *arr, int n) {
for (int i = 0; i < n; ++i) {
arr[i]++;
}
}
int main() {
int n = 1024;
int *h_arr = new int[n];
int *d_arr;
cudaMalloc(&d_arr, n * sizeof(int));
// 初始化数组...
// CPU端串行执行
func(h_arr, n);
// GPU端串行执行（仅使用1个线程）
func<<<1, 1>>>(d_arr, n);
// ...
}
此方案在GPU上也是串行的，效率低下。
优化方案（GPU并行）：
为了利用GPU的并行能力，我们修改调用方式，让每个GPU线程处理一个数组元素。
__host__ __device__ void func(int *elem) {
(*elem)++; // 只操作单个元素
}
__global__ void d_parallel(int *arr, int n) {
int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx < n) {
func(&arr[idx]); // 每个线程调用func处理一个元素
}
}
int main() {
// ... 分配和初始化内存
// CPU端仍串行
for (int i = 0; i < n; ++i) func(&h_arr[i]);
// GPU端并行
int threadsPerBlock = 256;
int blocks = (n + threadsPerBlock - 1) / threadsPerBlock;
d_parallel<<<blocks, threadsPerBlock>>>(d_arr, n);
// ...
}
现在GPU执行是并行的，但CPU端仍需显式循环。
最终方案（统一的函数接口）：
我们可以让 func 函数自身根据调用者决定行为，提供更统一的接口。
__host__ __device__ void func(int *arr, int n) {
// 当从GPU并行调用时，n=1，每个线程增加一个元素
// 当从CPU串行调用时，n=数组长度，循环增加所有元素
for (int i = 0; i < n; ++i) {
arr[i]++;
}
}
__global__ void d_parallel_unified(int *arr, int n) {
int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx < n) {
// 每个线程告诉func只增加它负责的那个元素
func(&arr[idx], 1);
}
}
int main() {
// ... 分配和初始化内存
// CPU端调用
func(h_arr, n); // n是数组长度，内部循环
// GPU端调用
d_parallel_unified<<<blocks, threadsPerBlock>>>(d_arr, n);
// ...
}
这个设计让 func 函数更加灵活，能适应不同的执行环境。
扩展知识：Thrust库简介 ⚡
在课程的最后，我们简要介绍一下CUDA的Thrust库。Thrust是一个类似于C++标准模板库的并行算法库，它提供了丰富的数据结构（如向量）和算法（如排序、归约、前缀和）。
Thrust的特点：
-
高层抽象：程序员无需关心内核启动、线程管理等底层细节。
-
可移植性：相同的Thrust代码可以在CPU（使用STL后端）或GPU（使用CUDA后端）上运行，编译器根据设备选择实现。
-
高性能：库内部经过优化，能有效利用GPU的并行计算能力。
简单示例（归约求和）：
使用Thrust进行数组求和比手动编写归约内核简单得多：
#include <thrust/reduce.h>
#include <thrust/device_vector.h>
// ...
thrust::device_vector<int> d_vec(n);
// ... 填充数据
int sum = thrust::reduce(d_vec.begin(), d_vec.end(), 0, thrust::plus<int>());
Thrust库包含大量此类实用算法，能极大提升CUDA程序的开发效率。
总结
本节课中我们一起学习了CUDA中不同类型的函数及其调用规则。我们深入探讨了 __global__、__device__ 和 __host__ 函数的定义、调用限制以及它们与统一内存的协作。通过数组增量等示例，我们实践了如何编写既能用于CPU串行执行也能用于GPU并行执行的灵活函数。最后，我们简要了解了能简化并行编程的Thrust库。掌握这些函数类型是构建高效、复杂CUDA应用程序的基础。
017：CUDA中的支持功能
在本节课中，我们将要学习CUDA编程中的一些支持功能，包括如何使用Makefile自动化编译和运行流程，如何使用CUDA-GDB调试程序，以及如何使用性能分析工具（如nvprof）来评估和优化CUDA程序的性能。
Makefile基础
上一节我们介绍了CUDA中的函数类型。本节中，我们来看看如何利用Makefile来管理CUDA项目的编译和运行。
Makefile是一组包含目标和变量的命令集合，用于自动化编译、链接和清理项目文件。它特别适用于管理大型项目中文件之间的依赖关系。
以下是Makefile的基本结构：
target: prerequisites
command
-
target：可以是生成的文件名，也可以是要执行的操作名称。
-
prerequisites：是执行该目标所依赖的文件或其他目标。
-
command：是生成目标或执行操作所需的一系列shell命令。注意：命令前必须使用Tab字符缩进，而不是空格。
一个简单的Makefile示例
以下是一个打印“Hello”的简单Makefile规则：
hello:
echo Hello
当在终端运行 make hello 命令时，它会执行 echo Hello，输出“Hello”。
用于CUDA程序的Makefile示例
让我们为一个打印“Hello”32次的简单CUDA程序编写Makefile。我们希望它能自动完成编译、运行和清理。
hello:
nvcc hello.cu -o hello
run: hello
./hello
clean: run
rm hello
这个Makefile包含三个规则：
-
hello：编译hello.cu文件，生成可执行文件hello。 -
run：依赖于hello目标。在hello编译完成后，运行./hello。 -
clean：依赖于run目标。在程序运行后，删除生成的可执行文件hello。
执行 make clean 将按顺序触发这三个目标。
在Makefile中使用变量
我们可以使用变量来使Makefile更清晰、更易于维护。
FILE1 = hello
FILE2 = run
FILE3 = clean
$(FILE1):
nvcc hello.cu -o hello
$(FILE2): $(FILE1)
./hello
$(FILE3): $(FILE2)
rm hello
变量通过 = 赋值，并通过 $(变量名) 引用。
使用 all 目标管理多个任务
all 目标常用于指定默认要构建的一系列目标。
all: hello1 hello2 hello3
hello1:
nvcc hello1.cu -o hello1
./hello1
hello2:
nvcc hello2.cu -o hello2
./hello2
hello3:
nvcc hello3.cu -o hello3
./hello3
clean:
rm hello1 hello2 hello3
运行 make all 将按顺序编译并运行 hello1、hello2 和 hello3，最后可以运行 make clean 进行清理。
Makefile中的条件与循环
Makefile支持条件判断和循环，以增加灵活性。
条件判断示例：
TARGET_CPU = x86_64
ifeq ($(TARGET_CPU), x86_64)
ARCH_FLAG = -m64
else
ARCH_FLAG = -m32
endif
循环示例：
run_tests:
for number in 1 2 3 4 5; do \
./my_program input$$number.txt output$$number.txt; \
done
综合示例：自动化测试多个图数据
假设我们有一个CUDA程序 vertex_removal.cu，需要对 test_bench.zip 中的多个图文件（t1.mtx 到 t10.mtx）进行测试。
all: vertex_removal
unzip -o examples/test_cases/test_bench.zip -d .
number=1; \
while [ $$number -le 10 ]; do \
if [ $$number -ne 5 ]; then \
./vertex_removal t$$number.mtx vr_graph_$$number.txt; \
fi; \
number=$$((number + 1)); \
done
vertex_removal:
nvcc vertex_removal.cu -o vertex_removal
clean:
rm -f vertex_removal *.txt
这个Makefile自动化了以下流程：
-
编译
vertex_removal.cu。 -
解压测试数据。
-
使用循环对除
t5.mtx外的所有测试图运行程序，并将结果输出到对应文件。 -
clean目标用于清理生成的可执行文件和输出文件。
使用CUDA-GDB调试
调试并行程序比调试串行程序更具挑战性，因为线程执行顺序的不确定性可能导致不同的结果。CUDA-GDB是GNU调试器(GDB)的扩展，专门用于调试运行在GPU硬件上的CUDA程序。
调试示例：空指针解引用
考虑以下存在错误的CUDA内核：
__global__ void myKernel(int *x) {
*x = 0; // 潜在的空指针解引用
}
int main() {
int *d_x; // 仅在主机端声明指针，未在设备上分配内存
myKernel<<<2, 10>>>(d_x);
cudaDeviceSynchronize();
return 0;
}
此代码中，设备指针 d_x 未被分配设备内存，导致内核中尝试解引用非法地址。
使用 cudaGetLastError 捕获错误
一种调试方法是使用 cudaGetLastError 函数。
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("CUDA error: %s\\n", cudaGetErrorString(err));
}
运行有问题的程序后，此代码会输出类似 CUDA error: an illegal memory access was encountered 的信息，提示存在非法内存访问。
使用CUDA-GDB进行交互式调试
更强大的方法是使用CUDA-GDB。
-
编译调试版本：使用
-G和-g标志编译CUDA代码，生成调试信息。nvcc -G -g my_program.cu -o my_program_debug -
启动调试：
cuda-gdb my_program_debug -
设置断点并运行：在
(cuda-gdb)提示符下：(cuda-gdb) break main (cuda-gdb) run -
分析错误：程序崩溃后，CUDA-GDB会显示错误位置（如
my_program.cu:6）、引发错误的线程信息（网格、块、线程ID）以及指针的值（可能为0x0），从而帮助定位空指针解引用问题。
常用的CUDA-GDB命令
-
info cuda kernels：列出当前所有CUDA内核的信息。 -
info cuda threads：显示当前内核中线程的状态。 -
cuda thread block thread：切换调试焦点到指定的线程块和线程。 -
info cuda lanes：显示特定线程（通道）的详细信息。 -
break filename.cu:23 if threadIdx.x == 1：在特定文件和行号设置条件断点。
性能分析：nvprof
性能分析是优化CUDA程序的关键步骤。nvprof 是NVIDIA提供的命令行性能分析工具，属于非侵入式分析，无需修改源代码。
使用nvprof进行性能分析
假设我们有一个程序调用了三个内核（kernel1、kernel2、kernel3）各100次。
-
直接分析：
nvprof ./my_cuda_programnvprof会输出每个内核的执行时间、调用次数以及占总时间的百分比。==PROF== Profiling result: Time(%) Time Calls Kernel 39.2% 1.23ms 100 kernel2 33.1% 1.04ms 100 kernel3 26.5% 0.83ms 100 kernel1 95.0% -- 300 CUDA API (cudaLaunchKernel)从结果可知，
kernel2耗时最多，是首要的优化目标。同时，CUDA内核启动开销也占据了显著时间。 -
分析特定代码区域：可以使用
nvprof的API在代码中标记分析范围。#include <nvToolsExt.h> nvtxRangePushA("MyCodeSection"); // ... 要分析的代码 ... nvtxRangePop();然后运行：
nvprof --analysis-metrics -o profile.nvvp ./my_cuda_program生成的
profile.nvvp文件可以用更可视化的工具（如NVIDIA Nsight Systems）进行深入分析。
nvprof 可以提供丰富的指标，包括内存吞吐量、缓存命中率、分支分化等，帮助开发者全面了解程序在GPU上的行为。
本节课中我们一起学习了CUDA编程中的三项重要支持功能：使用Makefile自动化项目构建流程，使用CUDA-GDB调试GPU内核代码，以及使用nvprof工具进行程序性能分析。掌握这些工具能显著提升CUDA程序的开发效率和运行性能。
018：动态并行、多GPU与线程束投票 🚀
概述
在本节课中，我们将学习CUDA编程中的三个高级主题：动态并行、多GPU编程以及线程束投票。动态并行允许从GPU内核内部启动新的内核，多GPU编程则涉及如何利用系统中的多个GPU进行协同计算，而线程束投票则提供了一种在同一个线程束内进行高效通信和决策的机制。
动态并行 🔄
上一节我们介绍了CUDA程序的构建、调试和性能分析。本节中，我们来看看动态并行。
动态并行在需要嵌套并行的情况下非常有用。例如，从CPU启动一个包含N个线程的内核，然后该内核的每个线程再启动另一个包含M个线程的内核。这种模式可以高效地模拟嵌套循环等场景，特别是当内层循环的迭代次数在运行时才能确定时。
动态并行的适用场景
以下是动态并行典型的适用场景：
-
分层数据结构：例如树结构，每个节点的子节点数量可能不同。遍历所有节点（外层循环），并对每个节点的可变数量的子节点进行处理（内层循环）。
-
递归算法：递归调用可以建模为递归树，树中每个节点的分支数量（递归调用次数）可能不同。
-
非均匀工作负载：任务可以自然地划分为独立的批次，但每个批次的工作量并不均匀。
注意：并非所有嵌套并行都需要动态并行。如果内外层循环的迭代次数在编译时已知，可以使用静态并行，即直接启动一个包含 N * M 个线程的单一内核，而不是使用动态并行。
动态并行示例
以下是一个简单的动态并行示例代码：
__global__ void child(int father) {
printf("Parent %d, Child %d\n", father, threadIdx.x);
}
__global__ void parent() {
printf("Parent %d\n", threadIdx.x);
child<<<1, 5>>>(threadIdx.x);
}
int main() {
parent<<<1, 3>>>();
cudaDeviceSynchronize();
return 0;
}
编译与运行：编译动态并行代码需要指定计算能力（至少3.5）并启用可重定位设备代码模式。
nvcc -arch=sm_35 -rdc=true dynamic_example.cu -o dynamic_example
程序输出分析：parent 内核由3个线程执行，每个线程打印自己的ID，然后启动一个包含5个线程的 child 内核。child 内核打印其父线程ID和自己的线程ID。由于内核启动是异步的，且多个 parent 线程并发执行，child 内核的输出顺序可能与 parent 线程的启动顺序交错，但属于同一个 parent 线程启动的5个 child 线程的输出顺序是固定的。
动态并行的计算与内存特性
-
计算特性：
-
父内核与子内核关联着各自的网格。
-
一个父内核可以启动多个子内核网格。
-
父内核和子内核可能异步执行。父内核线程在启动子内核后不会等待其完成，而是继续执行后续语句。
-
父网格在其所有子网格完成之前不会被视为完成。
-
-
内存特性：
-
父内核和子内核共享相同的全局内存和常量内存。
-
父内核和子内核拥有不同的本地内存和共享内存。不能将父内核的共享内存指针传递给子内核直接访问。
-
在子内核启动之前，父内核完成的所有全局内存操作对子内核可见。
-
子内核完成的所有全局内存操作，只有在父内核通过
cudaDeviceSynchronize()等方式同步等待子内核完成后，才对父内核可见。
-
多GPU编程 🖥️🖥️
上一节我们介绍了动态并行。本节中，我们来看看如何利用多个GPU进行并行计算。
使用多个GPU的主要优势在于可以获得更多的计算核心和更大的总内存容量，从而能够加速计算或处理更大规模的数据。
多GPU系统架构
典型的系统包含一个CPU和多个GPU。GPU之间可能存在两种连接方式：
-
点对点访问：某些GPU对（例如，通过PCIe桥接器或NVLink直接连接）可以直接访问彼此的内存，无需经过CPU，数据传输延迟较低。
-
通过CPU中转：对于没有直接连接的GPU对，数据交换需要通过CPU内存中转（GPU A -> CPU内存 -> GPU B），这会带来额外的开销。
多GPU编程基础
在CUDA中，系统中的GPU被编号为0到N-1。核心API是 cudaSetDevice(int device_id)，用于设置当前操作的GPU设备。在此调用之后，所有的内核启动、内存分配和数据传输操作都将在该设备上进行。
以下代码展示了基本的双GPU使用模式：
cudaSetDevice(0); // 切换到 GPU 0
kernel1<<<...>>>(...); // 在 GPU 0 上启动 kernel1
cudaMemcpyAsync(..., cudaMemcpyDeviceToHost); // 从 GPU 0 异步拷贝数据
cudaSetDevice(1); // 切换到 GPU 1
kernel2<<<...>>>(...); // 在 GPU 1 上启动 kernel2
cudaMemcpyAsync(..., cudaMemcpyDeviceToHost); // 从 GPU 1 异步拷贝数据
关键点：为了确保 kernel1 和 kernel2 能够真正在GPU 0和GPU 1上并发执行，必须使用异步内存拷贝函数 cudaMemcpyAsync，而不是阻塞式的 cudaMemcpy。否则，CPU会在第一个 cudaMemcpy 处等待，导致内核串行执行。
多GPU相关API
以下是用于多GPU编程的一些重要API：
-
cudaGetDeviceCount(int* count)：获取系统中可用的GPU数量。 -
cudaDeviceCanAccessPeer(int* canAccessPeer, int device, int peerDevice)：检查device能否直接访问peerDevice的内存（点对点访问）。 -
cudaDeviceEnablePeerAccess(int peerDevice, unsigned int flags)：启用从当前设备到peerDevice的点对点访问。
注意：点对点访问在硬件层面是对称的，但软件API需要显式地双向启用。一个设备最多只能与8个其他设备建立点对点连接。
枚举设备属性示例
以下代码展示了如何获取并打印系统中每个GPU设备的属性：
int device_count;
cudaGetDeviceCount(&device_count);
for (int i = 0; i < device_count; ++i) {
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, i);
printf("Device %d: %s\n", i, prop.name);
printf(" Compute Capability: %d.%d\n", prop.major, prop.minor);
// ... 打印其他属性
}
线程束投票 ✋
上一节我们介绍了多GPU编程。本节中，我们来看看线程束投票，这是一种在同一个线程束（Warp，通常32个线程）内进行高效协作的机制。
线程束投票函数对于需要根据某个条件在 warp 内进行集体决策的场景非常有用，例如，当需要根据数据依赖性为整个 warp 分配任务时。
线程束投票函数
CUDA 提供了以下内置的线程束投票函数：
-
__all(int predicate)：如果 warp 内所有活动线程的predicate值非零，则对所有线程返回1（真），否则返回0（假）。 -
__any(int predicate)：如果 warp 内任意活动线程的predicate值非零，则对所有线程返回1，否则返回0。 -
__ballot(int predicate)：返回一个32位的无符号整数掩码（mask），其中第n位对应于 warp 内第n个线程的predicate值（非零则为1，否则为0）。该掩码对所有线程返回相同的值。
__ballot 是 __all 和 __any 的泛化形式。
投票函数示例
考虑以下内核，它使用 __all 检查 warp 内所有线程的 threadIdx.x 是否都小于100：
__global__ void kernel() {
int tid = threadIdx.x;
int pred = (tid < 100);
int result = __all(pred);
if (tid % 32 == 0) { // 每个warp选一个代表线程打印
printf("Warp %d: __all(tid<100) = %d\n", tid/32, result);
}
}
// 启动配置：<<<1, 128>>> (4个warps)
对于前3个warps（线程0-95），所有线程都满足条件，result 为1。对于第4个warp（线程96-127），只有部分线程满足条件，result 为0。输出可能为：1, 1, 1, 0。
如果将 __all 替换为 __any，则第4个warp中也有线程满足条件，因此所有warps的 result 都为1。
__ballot 示例：
unsigned int mask = __ballot(threadIdx.x % 2 == 0); // 检查是否为偶数线程
对于每个warp，mask 是一个32位数，偶数位为1，奇数位为0。其十六进制表示为 0x55555555（二进制0101...）。
使用投票优化原子操作
线程束投票的一个典型应用是优化高争用的原子操作。考虑以下需要条件原子递增的代码：
// 原始代码：高争用
if (condition) {
atomicAdd(&counter, value);
}
如果warp中有多个线程满足 condition，它们会各自调用 atomicAdd，导致严重的资源争用。
优化版本使用 __ballot 和 __popc（计算掩码中1的位数）来减少原子操作调用次数：
// 优化代码：减少争用
unsigned int mask = __ballot(condition);
if ((threadIdx.x % 32) == 0) { // 每个warp选一个leader线程
int num_active = __popc(mask); // 计算满足条件的线程数
if (num_active > 0) {
atomicAdd(&counter, num_active * value); // Leader线程执行一次原子加
}
}
在这个优化版本中，无论warp中有多少线程满足条件，每个warp只执行一次原子操作，由选出的leader线程完成，从而显著降低了争用。
注意：投票函数的结果受线程活动状态影响。如果投票函数被放在条件语句（如 if(tid % 2 == 0)）内部，那么只有满足条件的“活动”线程会参与投票，不满足条件的线程被掩蔽（masked），不参与投票计算。
总结
本节课中我们一起学习了CUDA编程的三个高级主题。
-
动态并行：允许内核从设备端启动新的内核，适用于处理运行时才能确定工作量的嵌套并行任务，但需要注意其启动开销和内存访问规则。
-
多GPU编程：通过
cudaSetDevice管理多个GPU设备，利用异步操作实现真正的并发执行，并可以使用点对点访问优化GPU间的数据传输。 -
线程束投票：提供了
__all、__any和__ballot等函数，用于在warp内进行高效的集体决策和通信，常用于优化原子操作等高性能计算模式。
掌握这些高级特性，能够帮助你编写出更灵活、更高效、能够充分利用硬件资源的CUDA程序。
019：CUDA中的多维数组 🚀
在本节课中，我们将学习如何在CUDA中声明和使用多维数组。我们将通过三个具体的应用实例来深入理解这一概念：矩阵平方、二维最小值算法和模板计算。通过本课，你将学会如何更直观地处理多维数据，避免复杂的索引计算，并编写更易读的GPU代码。
概述 📋
在之前的课程中，我们使用一维数组来表示二维矩阵，并通过公式 i * matrix_size + j 来访问元素。本节课，我们将学习如何直接在CUDA中声明和访问二维（甚至更高维）数组，从而使代码更简洁、更易维护。
矩阵平方应用 🔲
上一节我们回顾了使用一维数组进行矩阵平方计算的方法。本节中，我们来看看如何使用真正的二维数组来实现相同的功能。
在之前的实现中，我们使用一维数组并通过公式计算索引来模拟二维访问。现在，我们将直接在GPU上声明二维数组。
CPU版本（使用一维数组）
void cpu_square(float* matrix, float* result, int matrix_size) {
for (int i = 0; i < matrix_size; i++) {
for (int j = 0; j < matrix_size; j++) {
for (int k = 0; k < matrix_size; k++) {
result[i * matrix_size + j] += matrix[i * matrix_size + k] * matrix[k * matrix_size + j];
}
}
}
}
GPU版本（使用二维数组）
以下是使用CUDA二维数组实现矩阵平方的关键步骤：
-
在GPU上声明二维数组：
__device__ float d_A[N][N]; __device__ float d_B[N][N]; -
主机端代码配置：
dim3 blocksPerGrid(1, 1); dim3 threadsPerBlock(N, N); // 使用 cudaMemcpyToSymbol 将数据从主机复制到设备 cudaMemcpyToSymbol(d_A, A, size); cudaMemcpyToSymbol(d_B, B, size); // 启动核函数 square_kernel<<<blocksPerGrid, threadsPerBlock>>>(N); // 将结果复制回主机 cudaMemcpyFromSymbol(C, d_B, size);
-
核函数实现：
__global__ void square_kernel(int n) { int i = threadIdx.y; int j = threadIdx.x; if (i < n && j < n) { float sum = 0.0f; for (int k = 0; k < n; k++) { sum += d_A[i][k] * d_A[k][j]; } d_B[i][j] = sum; } }
核心优势： 我们不再需要复杂的索引计算 i * n + j，可以直接使用 d_A[i][j] 这种直观的语法访问元素，代码可读性大大增强。
二维最小值算法 📉
上一节我们介绍了矩阵操作，本节中我们来看看一个图像处理中常见的算法——二维最小值算法。
该算法扫描输入矩阵中每个2x2的窗口，找出窗口内的最小值，并放入输出矩阵的对应位置。这类似于卷积操作。
算法定义
对于输出矩阵 O 中的元素 O[i][j]，其值为输入矩阵 A 中以下四个元素的最小值：
A[i-1][j], A[i-1][j+1], A[i][j], A[i][j+1]
边界情况（第一行和最后一列）保持不变。
CPU实现
void compute(float grid[N][N], float prev_grid[N][N]) {
for (int i = 1; i < N; i++) { // 忽略第一行
for (int j = 0; j < N-1; j++) { // 忽略最后一列
float min_val = fminf(prev_grid[i-1][j], prev_grid[i-1][j+1]);
min_val = fminf(min_val, prev_grid[i][j]);
min_val = fminf(min_val, prev_grid[i][j+1]);
grid[i][j] = min_val;
}
}
}
GPU实现
以下是GPU核函数的关键部分：
__global__ void min2d_kernel(int n) {
int i = blockIdx.y * blockDim.y + threadIdx.y;
int j = blockIdx.x * blockDim.x + threadIdx.x;
// 处理边界，避免第一行和最后一列
if (i > 0 && j < n-1) {
float min_val = fminf(d_prev_grid[i-1][j], d_prev_grid[i-1][j+1]);
min_val = fminf(min_val, d_prev_grid[i][j]);
min_val = fminf(min_val, d_prev_grid[i][j+1]);
d_grid[i][j] = min_val;
} else if (i == 0 || j == n-1) {
// 边界直接复制
d_grid[i][j] = d_prev_grid[i][j];
}
}
实现要点： 每个GPU线程负责计算一个输出元素。我们使用二维线程块来自然映射到矩阵的(i, j)坐标，并通过条件判断优雅地处理了边界情况。
模板计算 🧮
上一节我们学习了空间上的邻域操作，本节中我们来看看一个涉及多次迭代的模板计算。
模板计算中，每个单元格的新值由其上方两个邻居的当前值决定，公式为：
A[i][j] += A[i-1][j] + A[i-1][j+1]
该计算需要迭代多次（例如10次），且每次迭代后，新值成为下一次迭代的输入。
基础版本（使用额外空间）
最简单的实现是使用两个独立的数组，一个用于读取（上一轮结果），一个用于写入（本轮结果）。
GPU核函数示例：
__global__ void stencil_kernel(int n) {
int i = blockIdx.y * blockDim.y + threadIdx.y;
int j = blockIdx.x * blockDim.x + threadIdx.x;
if (i > 0 && j < n-1) {
d_grid_new[i][j] = d_grid_old[i][j] + d_grid_old[i-1][j] + d_grid_old[i-1][j+1];
}
}
// 每次迭代后，需要将 d_grid_new 的数据复制到 d_grid_old
进阶挑战：原地计算
现在，我们增加一个限制：不能使用额外的数组，必须在原输入矩阵上就地更新数据。
这就带来了数据竞争问题：当线程正在读取A[i-1][j]和A[i-1][j+1]来计算A[i][j]的新值时，另一个负责计算A[i-1][j]的线程可能正在写入它。
解决方案：按行顺序执行
为了消除数据竞争，我们可以将并行度限制在单行内。即，每次只计算一行：
-
从最后一行开始计算（因为它没有下方的依赖）。
-
等待该行所有线程计算完成后，再开始计算上一行。
-
依此类推，直到第一行。
主机端代码逻辑：
for (int iter = 0; iter < 10; iter++) {
for (int row = N-1; row >= 1; row--) { // 从底向上
// 1. 将当前行数据复制到设备
// 2. 启动核函数，只计算这一行
dim3 threadsPerBlock(N, 1); // 一行N个线程
stencil_inplace_kernel<<<1, threadsPerBlock>>>(N, row);
// 3. 将计算后的行复制回主机
}
}
核函数实现：
__global__ void stencil_inplace_kernel(int n, int current_row) {
int j = threadIdx.x; // 列索引
int i = current_row; // 行索引由主机传入
if (j < n-1) { // 忽略最后一列
d_grid[i][j] += d_grid[i-1][j] + d_grid[i-1][j+1];
}
}
性能权衡： 这种方法保证了正确性，但牺牲了并行度（从并行处理整个矩阵变为每次只并行处理一行）。另一种可能的方案是使用原子操作，但需要精心设计以避免将并行计算完全序列化。这可以作为课后思考题。
总结 🎯
本节课我们一起学习了CUDA中多维数组的核心用法：
-
声明与访问： 使用
__device__ float arr[N][M]在GPU上声明二维数组，并可直接用arr[i][j]语法访问，无需手动计算线性索引。 -
三个应用实例：
-
矩阵平方：展示了用二维数组简化代码逻辑。
-
二维最小值算法：演示了如何将图像处理类算法映射到二维线程网格，并处理边界条件。
-
模板计算：深入探讨了迭代计算中的依赖问题，并给出了通过限制执行顺序来实现原地计算的方案。
-
-
核心思想： 利用CUDA的二维/三维线程组织来自然匹配多维数据结构，使算法更清晰，并注意处理边界条件和迭代计算中的数据竞争问题。
通过掌握这些技巧，你可以更高效、更直观地在GPU上处理科学计算、图像处理等领域的多维数据问题。
浙公网安备 33010602011771号