从零到多GPU：NVIDIA开源教程仓库教你系统掌握CUDA编程，容器化部署一键启动

GPU编程长久以来都被视作一道高耸的技术门槛——初学者不仅要掌握线程层级架构、显存空间模型这类抽象概念，还得与编译工具链及运行时环境反复周旋。

NVIDIA 推出的accelerated-computing-hub代码仓库，正试图弥合这道鸿沟：它并非一份零散的示例代码集合，而是一套经过细致打磨、从“Hello GPU”一路延伸至多 GPU 分布式运算的完整教学体系。尤其值得关注的是，其代码体现了一种“教学驱动的软件工程”理念——每个源文件都瞄准特定的认知目标，代码复杂度严格递进，并借助 Docker 容器实现零配置运行环境。

• NVIDIA 精选的通用 GPU 编程相关教育资源合集
• https://github.com/NVIDIA/accelerated-computing-hub
• 4000 字，阅读 20 分钟，播客 16 分钟

该仓库提供了一份教程大纲，汇总了 NVIDIA 加速计算中心所有可在 NVIDIA Brev、Google Colab 上运行的交互式 GPU 编程教程。

下表清晰列出了教程内容主题、配套的Docker Compose 配置文件、适配的Brev 硬件实例型号，以及支持部署运行的Brev 服务提供商和系统环境要求，便于开发者快速匹配硬件与平台开展学习实操。其中，NVIDIA Brev 是 NVIDIA 面向人工智能原型开发与落地交付的开发者云平台。

本文将深入剖析这个仓库，聚焦其核心教学代码的设计细节，以及它是如何通过 Standard Parallelism 这条“零 CUDA 代码”路径，为传统 HPC 开发者开启 GPU 加速的大门。

unsetunset本文目录unsetunset

• 快速上手
• 一、仓库整体架构与课程编排逻辑
• 1.1 三条并行的学习路径
• 1.2 CUDA C++课程的三阶递进
• 二、从单线程到多线程：Kernel 编程的渐进式代码设计
• 2.1 起点：一个单线程的 GPU kernel
• 2.2 引入线程并行：stride loop 模式
• 2.3 kernel 启动语法的最小化演示
• 三、共享内存与线程同步：从原理到优化的完整闭环
• 3.1 共享内存的最小可运行示例
• 3.2 实战：基于共享内存的直方图优化
• 3.3 CUB 协作算法：站在巨人肩膀上
• 四、Standard Parallelism：零 CUDA 代码的 GPU 加速之路
• 4.1 用标准 C++算法直接跑 GPU
• 4.2 二维热传导方程：stdpar 的 HPC 实战
• 五、容器化基础设施：教学环境的工程保障
• 5.1 Docker Compose 的多服务编排
• 总结

unsetunset快速上手unsetunset

克隆仓库并通过 Docker Compose 一键启动 CUDA C++教程的 JupyterLab 环境：

git clone https://github.com/NVIDIA/accelerated-computing-hub.git
cd accelerated-computing-hub/tutorials/cuda-cpp/brev
docker compose up

启动后访问 http://127.0.0.1:8888 即可进入交互式 Notebook。也可直接通过 Google Colab 在线运行，无需本地 GPU。更多部署方式（NVIDIA Brev 云实例、GPU 型号选择等），参考 docs/brev.md。

unsetunset一、仓库整体架构与课程编排逻辑unsetunset

1.1 三条并行的学习路径

整个仓库围绕三大教学模块展开，对应三类目标受众：

三条路径彼此独立，却又共享相同的基础设施（Docker 镜像、Brev 配置、测试框架）。这种设计使得不同背景的开发者都能找到最契合自身知识结构的切入点。

1.2 CUDA C++课程的三阶递进

CUDA C++教程被拆分为三个精心设计的阶段：

第一阶段（01.xx）：并行算法入门——无需手写 Kernel，借助 thrust/CCCL 等高层抽象来理解执行空间与内存空间的概念。

第二阶段（02.xx）：异步与流——深入掌握 GPU 的异步执行本质、CUDA Stream 以及 Pinned Memory 的使用。

第三阶段（03.xx）：自定义 Kernel——从最基础的单线程 Kernel 起步，逐步演进到共享内存与协作算法的优化。

这种“先学会使用、再理解原理、最后深入优化”的教学理念，在代码层面得到了淋漓尽致的体现。

二、从单线程到多线程：Kernel 编程的渐进式代码设计

2.1 起点：一个单线程的 GPU kernel

教程的 Kernel 部分从一个刻意“低效”的实现开始，旨在让学习者先掌握基本语法：

// 来源：tutorials/cuda-cpp/notebooks/03.02-Kernels/Sources/simple-kernel.cpp  
#include "ach.h"  

__global__ void single_thread_kernel(ach::temperature_grid_f in, float *out) {  
for (int id = 0; id < in.size(); id++) {  
out[id] = ach::compute(id, in);  
}  
}  

void simulate(ach::temperature_grid_f temp_in, float *temp_out,  
cudaStream_t stream) {  
single_thread_kernel<<<1, 1, 0, stream>>>(temp_in, temp_out);  
}  

注意这里的 <<<1, 1>>>——它仅启动一个 Block 内的一个线程。这段代码的教学意图非常明确：让学习者首先理解 Kernel 的语法结构和启动方式，暂时不要考虑并行性的问题。

• ach::temperature_grid_f 是教程封装好的温度网格数据类型。
• ach::compute 封装了 Stencil 计算逻辑，让学习者能够专注于并行化改造本身。

2.2 引入线程并行：stride loop 模式

随后，代码进化到多线程版本，展示了经典的跨步循环模式：

// 来源：tutorials/cuda-cpp/notebooks/03.02-Kernels/Sources/block-256-kernel.cpp  
#include "ach.h"  

const int number_of_threads = 256;  

__global__ void block_kernel(ach::temperature_grid_f in, float *out) {  
int thread_index = threadIdx.x;  

for (int id = thread_index; id < in.size(); id += number_of_threads) {  
out[id] = ach::compute(id, in);  
}  
}  

void simulate(ach::temperature_grid_f temp_in, float *temp_out,  
cudaStream_t stream) {  
block_kernel<<<1, number_of_threads, 0, stream>>>(temp_in, temp_out);  
}  

这里展示了经典的 stride loop（跨步循环）模式：每个线程从自己的 threadIdx.x 出发，以线程总数为步长遍历所有数据。教程还特意提供了一个 number_of_threads = 2048 的“错误版本”（failed-block-kernel.cpp），让学习者亲身体验超过硬件线程块上限（1024）时引发的运行时错误。这种“先犯错、再纠正”的教学手法贯穿于整个仓库。

2.3 kernel 启动语法的最小化演示

为了单独解释 <<<blocks, threads, shmem, stream>>> 这个 CUDA 特有的语法，教程提供了一个极简的“Hello World” Kernel：

// 来源：tutorials/cuda-cpp/notebooks/03.02-Kernels/Sources/kernel-launch.cpp  
#include <cstdio>  

__global__ void kernel(int value) {  
std::printf("value on device = %dn", value);  
}  

int main() {  
int blocks_in_grid = 1;  
int threads_in_block = 1;  
cudaStream_t stream = 0;  
kernel<<<blocks_in_grid, threads_in_block, 0, stream>>>(42);  
cudaStreamSynchronize(stream);  
}  

每个启动参数都被赋予了具体的变量名，自文档化程度极高。这比直接展示 kernel<<<1,1>>>(42) 要清晰得多——变量名本身就是最好的注释。

三、共享内存与线程同步：从原理到优化的完整闭环

3.1 共享内存的最小可运行示例

共享内存是 CUDA 性能优化的核心武器之一。教程用一个仅有 4 个线程的极简案例来引入这一概念：

// 来源：tutorials/cuda-cpp/notebooks/03.05-Shared-Memory/Sources/simple-shmem.cpp  
#include <cstdio>  

__global__ void kernel()  
{  
__shared__ int shared[4];  
shared[threadIdx.x] = threadIdx.x;  
__syncthreads();  

if (threadIdx.x == 0)  
{  
for (int i = 0; i < 4; i++) {  
std::printf("shared[%d] = %dn", i, shared[i]);  
}  
}  
}  

int main() {  
kernel<<<1, 4>>>();  
cudaDeviceSynchronize();  
return 0;  
}  

这段代码精准地展示了三个关键点：

• __shared__ 声明将数组放入片上共享内存；
• __syncthreads() 确保所有线程完成写入后再读取；
• 仅由 0 号线程执行打印操作，避免输出混乱。

学习者可以立即编译并运行这段代码，观察到 shared[0]=0, shared[1]=1, ... 的输出，从而建立对共享内存可见性的直观理解。

3.2 实战：基于共享内存的直方图优化

在打好基础之后，教程直接切入实战环节——教你如何通过共享内存来优化直方图计算的 kernel：

// 来源：tutorials/cuda-cpp/notebooks/03.05-Shared-Memory/Solutions/shmem.cpp  
__global__ void histogram_kernel(cuda::std::span<float> temperatures,  
cuda::std::span<int> histogram) {  
__shared__ int block_histogram[10];  

if (threadIdx.x < 10) {  
block_histogram[threadIdx.x] = 0;  
}  
__syncthreads();  

int cell = blockIdx.x * blockDim.x + threadIdx.x;  
int bin = static_cast<int>(temperatures[cell] / 10);  

cuda::atomic_ref<int, cuda::thread_scope_block> block_ref(  
block_histogram[bin]);  
block_ref.fetch_add(1);  
__syncthreads();  

if (threadIdx.x < 10) {  
cuda::atomic_ref<int, cuda::thread_scope_device> ref(  
histogram[threadIdx.x]);  
ref.fetch_add(block_histogram[threadIdx.x]);  
}  
}

这段代码在教学层面蕴含着极高的价值：

• 两级归约机制：首先在 block 级别的共享内存中完成累加（借助 thread_scope_block 的原子操作），随后将各个 block 的局部结果汇总到全局内存中（通过 thread_scope_device 的原子操作实现）。
• 现代 CUDA C++ 风格：使用 cuda::std::span 替代了传统的裸指针，用 cuda::atomic_ref 取代了老式的 atomicAdd 函数，充分展示了 libcu++ 库提供的现代、类型安全的接口。
• 性能意义：在 block 内部对共享内存执行原子操作，其速度比直接操作全局内存快出整整一个数量级以上，这主要是因为它完全避免了跨流式多处理器（SM）的互联争用。

3.3 CUB 协作算法：站在巨人的肩膀上

教程的最后一站，引入了 CUB 库中的 block-level 原语，向学习者展示了工业级的协作归约实现：

// 来源：tutorials/cuda-cpp/notebooks/03.06-Cooperative-Algorithms/Sources/block-sum.cpp  
#include <cub/block/block_reduce.cuh>  

constexpr int block_threads = 128;  

__global__ void block_sum()  
{  
using block_reduce_t = cub::BlockReduce<int, block_threads>;  
using storage_t = block_reduce_t::TempStorage;  

__shared__ storage_t storage;  

int block_sum = block_reduce_t(storage).Sum(threadIdx.x);  

if (threadIdx.x == 0)  
{  
printf("block sum = %dn", block_sum);  
}  
}  

int main() {  
block_sum<<<1, block_threads>>>();  
cudaDeviceSynchronize();  
return 0;  
}

CUB 库提供的 BlockReduce 在底层实现了 warp-level shuffle 指令以及多轮归约树，其性能远超手写的共享内存归约版本。整个教程的编排逻辑非常清晰：先让你亲手编写共享内存代码来理解底层原理，然后再告诉你“在实际工程中应当使用 CUB”——这种从原理到最佳实践的过渡堪称教科书级别。

四、Standard Parallelism：零 CUDA 代码的 GPU 加速之路

4.1 用标准 C++ 算法直接跑 GPU

stdpar 教程为开发者展示了一条截然不同的 GPU 编程路径——完全借助 ISO C++ 标准并行算法来实现 GPU 加速：

// 来源：tutorials/stdpar/notebooks/cpp/lab1_daxpy/solutions/exercise5.cpp  
#include <algorithm>  
#include <execution>  
#include <ranges>  
#include <vector>  

void initialize(std::vector<double> &x, std::vector<double> &y) {  
std::for_each_n(std::execution::par, std::views::iota(0).begin(), x.size(),  
[x = x.data()](int i) { x[i] = (double)i; });  
std::fill_n(std::execution::par, y.data(), y.size(), 2.);  
}  

void daxpy(double a, std::vector<double> const &x, std::vector<double> &y) {  
std::for_each_n(std::execution::par,  
std::views::iota(0).begin(), x.size(),  
[a, x = x.data(), y = y.data()](int i) {  
y[i] += a * x[i];  
});  
}

这段代码中看不到任何 CUDA 专属的关键字——没有 __global__，没有 <<<>>>，也没有 cudaMalloc。开发者只需要传入 std::execution::par 执行策略，配合 nvc++ -stdpar=gpu 编译选项，这些标准算法就能自动在 GPU 上运行。 更妙的是，教程还通过对比四种编译器（g++、clang++、nvc++ CPU 后端、nvc++ GPU 后端）的运行结果，让学习者直观地看到同一份代码在不同执行后端上的性能差异。

4.2 二维热传导方程：stdpar 的 HPC 实战

教程进一步演示了如何通过标准算法解决实际的科学计算问题——针对二维热传导方程的并行 stencil 计算：

// 来源：tutorials/stdpar/notebooks/cpp/lab2_heat/exercise2.cpp  
double apply_stencil(double* u_new, double* u_old, grid g, parameters p) {
auto xs = std::views::iota(g.x_begin, g.x_end);
auto ys = std::views::iota(g.y_begin, g.y_end);
auto ids = std::views::cartesian_product(xs, ys);
return std::transform_reduce(
std::execution::par, ids.begin(), ids.end(),
        0., std::plus{}, [u_new, u_old, p](auto idx) {
auto [x, y] = idx;
return stencil(u_new, u_old, x, y, p);
});
}

借助 std::views::cartesian_product 生成二维索引空间，再通过 std::transform_reduce 在并行执行 stencil 的同时完成能量求和归约——这正是 map-reduce 模式在标准 C++ 中的简洁体现。当结合 MPI 实现多 GPU 通信时，这段代码充分展示了标准并行算法在工业级 HPC 应用中的完整能力。

五、容器化基础设施：教学环境的工程保障

5.1 Docker Compose 的多服务编排

该教程采用精心设计的 Docker Compose 配置，实现了“一键部署”：

# 来源：tutorials/cuda-cpp/brev/docker-compose.yml  
services:
base:
<<: [*gpu-config, *common-service]
image: *image
entrypoint: ["/accelerated-computing-hub/brev/entrypoint.bash", "base"]
jupyter:
<<: [*gpu-config, *common-service, *persistent-service]
image: *image
entrypoint: ["/accelerated-computing-hub/brev/entrypoint.bash", "jupyter"]
ports:
      - "127.0.0.1:8888:8888"
nsys:
image: nvcr.io/nvidia/devtools/nsight-streamer-nsys:2026.1.1
ports:
      - "127.0.0.1:8080:8080"

只需执行一次 docker compose up，即可同时启动 JupyterLab、Nsight Systems 性能分析器和 Nsight Compute 分析器这三个服务。学习者在 Notebook 中写完 kernel 后，能够直接在浏览器中查看 GPU 时间线和指令级分析——这种“编码-分析”的即时反馈循环，是传统教学方式难以提供的。

总结

NVIDIA 的 accelerated-computing-hub 不仅仅是一个代码仓库，它代表了一种教学驱动的软件工程范式：

1. 认知负载递进：每个代码文件只引入一个新概念，从 <<<1,1>>> 到 cub::BlockReduce 的跨度被拆解为十余个易于消化的步骤
2. 多入口设计：三条路径（CUDA C++/stdpar/Python）让不同背景的开发者都能找到合适的起点
3. 环境即代码：Docker Compose 与预构建镜像消除了“环境配置”这一最大的学习障碍
4. 现代 C++ 优先：全面采用 cuda::std::span、cuda::atomic_ref、std::mdspan 等现代接口，而非传统的 C 风格 CUDA API

对于希望系统学习 GPU 编程的开发者而言，该仓库提供了当前最具结构化、最可复现的开源学习路径。而对于教学设计者来说，它展示了如何通过代码的渐进式组织来实现复杂知识的有效传递。

参考资料

[1] CUDA Tile 教程: https://github.com/NVIDIA/accelerated-computing-hub/blob/main/tutorials/cuda-tile/README.md

[2] CUDA C++ 教程: https://github.com/NVIDIA/accelerated-computing-hub/blob/main/tutorials/cuda-cpp/README.md

[3] 标准并行性教程: https://github.com/NVIDIA/accelerated-computing-hub/blob/main/tutorials/stdpar/README.md

[4] 加速 Python 教程: https://github.com/NVIDIA/accelerated-computing-hub/blob/main/tutorials/accelerated-python/README.md

[5] nvmath-python 教程: https://github.com/NVIDIA/accelerated-computing-hub/blob/main/tutorials/nvmath-python/README.md

[6] CUDA Python 教程 – CuPy、cuDF、CCCL 及内核（8 小时版）: https://github.com/NVIDIA/accelerated-computing-hub/blob/main/tutorials/accelerated-python/notebooks/syllabi/cuda_python__cupy_cudf_cccl_kernels__8_hours.ipynb

[7] CUDA Python 教程 – cuda.core 与 CCCL（2 小时版）: https://github.com/NVIDIA/accelerated-computing-hub/blob/main/tutorials/accelerated-python/notebooks/syllabi/cuda_python__cuda_core_cccl__2_hours.ipynb

[8] PyHPC 教程 – NumPy、CuPy 与 mpi4py（4 小时版）: https://github.com/NVIDIA/accelerated-computing-hub/blob/main/tutorials/accelerated-python/notebooks/syllabi/pyhpc__numpy_cupy_mpi4py__4_hours.ipynb

GPU 部署教程：https://github.com/NVIDIA/accelerated-computing-hub/blob/main/tutorials/gpu-deployment/gpu-deployment-from-scratch.md

延伸阅读

• NVIDIA CCCL (CUDA Core Compute Libraries) 核心架构深度分析
• 暴涨 45 个 Star！NVIDIA CUDA C++ 核心库 Lead 开源 AutoCUDA：永不停机的 GPU 内核优化智能体框架
• 暴涨 45 个 Star！ptx-isa-markdown 语料金矿：借助 Claude Code 突破 NVIDIA 文档的检索壁垒