超级会员免费看
1.2 高性能CUDA编程范式
引言
掌握了GPU的底层硬件架构后，下一步是学习如何通过CUDA编程模型来高效地驾驭它。本节将聚焦于三个核心的高性能编程范式，它们是区分CUDA新手与专家的关键分水岭。我们将学习如何利用流（Stream）和事件（Event）来掩盖数据传输延迟，实现计算与I/O的并行流水线；我们将重温并深化共享内存（Shared Memory）的应用，将其作为用户管理的高速缓存来最大化数据复用；最后，我们将介绍CUDA Graphs这一终极武器，以消除CPU启动开销，将应用的延迟推向极致。
1.2.1 CUDA流（Stream）与事件（Event）异步执行
原理简介
在典型的异构计算中，工作流包含三个阶段：1) 将数据从CPU主机拷贝到GPU设备 (H2D)，2) 在GPU上执行计算核函数，3) 将结果从GPU设备拷贝回CPU主机 (D2H)。默认情况下，这些操作在所谓的“默认流”中同步或隐式同步执行，即一个操作必须等待前一个完成后才能开始。这导致GPU在数据拷贝时处于空闲状态，CPU在等待GPU计算时也处于空闲状态，造成了严重的资源浪费。
CUDA流（Stream） 是一个GPU操作的有序队列。提交到同一个流的操作按顺序执行，但不同流中的操作可以并行执行（只要硬件资源允许）。通过创建多个流，我们可以构建一个流水线（Pipeline），将数据分割成块（Chunks），让第 N 块数据的计算与第 N+1 块数据的H2D拷贝以及第 N-1 块数据的D2H拷贝同时进行，从而有效隐藏数据传输的延迟。
CUDA事件（Event） 是CUDA流中的标记，可用于精确计时和流间同步。
技术手册：可执行代码
以下程序将通过实验对比三种执行模式的性能：
-
同步模式：完全串行执行。
-
单流异步模式：展示仅使用异步API但未构建流水线的情况。
-
多流异步流水线模式：通过重叠计算和数据传输实现最高性能。
C++
// =================================================================
// 编译指令:
// nvcc -o streams_pipeline streams_pipeline.cu
// =================================================================
#include <iostream>
#include <vector>
#include <cuda_runtime.h>
#define CUDA_CHECK(call) do { \
cudaError_t err = call; \
if (err != cudaSuccess) { \
fprintf(stderr, "CUDA Error: %s:%d, ", __FILE__, __LINE__); \
fprintf(stderr, "code: %d, reason: %s\n", err, cudaGetErrorString(err)); \
exit(1); \
} \
} while(0)
// 一个简单的计算密集型核函数
__global__ void simple_kernel(float* data, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
float val = data[idx];
for(int i = 0; i < 50; ++i) {
val = sinf(val) * cosf(val);
}
data[idx] = val;
}
}
int main() {
int n = 1 << 25; // approx 33 million floats
size_t bytes = n * sizeof(float);
// 1. 分配主机内存。使用锁页内存(Pinned Memory)以实现真正的异步H2D/D2H拷贝
float *h_data;
CUDA_CHECK(cudaMallocHost(&h_data, bytes));
for (int i = 0; i < n; ++i) {
h_data[i] = static_cast<float>(i);
}
// 分配设备内存
float* d_data;
CUDA_CHECK(cudaMalloc(&d_data, bytes));
// 计时事件
cudaEvent_t start, stop;
CUDA_CHECK(cudaEventCreate(&start));
CUDA_CHECK(cudaEventCreate(&stop));
float ms;
int threadsPerBlock = 256;
int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
// --- 模式1: 同步执行 ---
CUDA_CHECK(cudaEventRecord(start));
CUDA_CHECK(cudaMemcpy(d_data, h_data, bytes, cudaMemcpyHostToDevice));
simple_kernel<<<blocksPerGrid, threadsPerBlock>>>(d_data, n);
CUDA_CHECK(cudaMemcpy(h_data, d_data, bytes, cudaMemcpyDeviceToHost));
CUDA_CHECK(cudaEventRecord(stop));
CUDA_CHECK(cudaEventSynchronize(stop));
CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));
std::cout << "1. Synchronous Execution Time: " << ms << " ms" << std::endl;
// --- 模式2: 单流异步 ---
// 虽然API是异步的，但由于在同一个流中，操作仍然是串行的
cudaStream_t stream_single;
CUDA_CHECK(cudaStreamCreate(&stream_single));
CUDA_CHECK(cudaEventRecord(start));
CUDA_CHECK(cudaMemcpyAsync(d_data, h_data, bytes, cudaMemcpyHostToDevice, stream_single));
simple_kernel<<<blocksPerGrid, threadsPerBlock, 0, stream_single>>>(d_data, n);
CUDA_CHECK(cudaMemcpyAsync(h_data, d_data, bytes, cudaMemcpyDeviceToHost, stream_single));
// 等待流中所有操作完成
CUDA_CHECK(cudaStreamSynchronize(stream_single));
CUDA_CHECK(cudaEventRecord(stop));
CUDA_CHECK(cudaEventSynchronize(stop));
CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));
std::cout << "2. Single Stream Async Time: " << ms << " ms" << std::endl;
CUDA_CHECK(cudaStreamDestroy(stream_single));
// --- 模式3: 多流异步流水线 ---
int n_streams = 4;
cudaStream_t streams[n_streams];
float* d_chunks[n_streams];
int chunk_size = n / n_streams;
size_t chunk_bytes = chunk_size * sizeof(float);
int blocksPerGrid_chunk = (chunk_size + threadsPerBlock - 1) / threadsPerBlock;
for (int i = 0; i < n_streams; ++i) {
CUDA_CHECK(cudaStreamCreate(&streams[i]));
CUDA_CHECK(cudaMalloc(&d_chunks[i], chunk_bytes));
}
CUDA_CHECK(cudaEventRecord(start));
// 启动流水线
for (int i = 0; i < n_streams; ++i) {
int offset = i * chunk_size;
// 在第 i 个流上: H2D -> Kernel -> D2H
CUDA_CHECK(cudaMemcpyAsync(d_chunks[i], h_data + offset, chunk_bytes, cudaMemcpyHostToDevice, streams[i]));
simple_kernel<<<blocksPerGrid_chunk, threadsPerBlock, 0, streams[i]>>>(d_chunks[i], chunk_size);
CUDA_CHECK(cudaMemcpyAsync(h_data + offset, d_chunks[i], chunk_bytes, cudaMemcpyDeviceToHost, streams[i]));
}
// 同步所有流
for (int i = 0; i < n_streams; ++i) {
CUDA_CHECK(cudaStreamSynchronize(streams[i]));
}
CUDA_CHECK(cudaEventRecord(stop));
CUDA_CHECK(cudaEventSynchronize(stop));
CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));
std::cout << "3. Multi-Stream Pipeline Time: " << ms << " ms" << std::endl;
// 清理
for (int i = 0; i < n_streams; ++i) {
CUDA_CHECK(cudaStreamDestroy(streams[i]));
CUDA_CHECK(cudaFree(d_chunks[i]));
}
CUDA_CHECK(cudaFree(d_data));
CUDA_CHECK(cudaFreeHost(h_data));
CUDA_CHECK(cudaEventDestroy(start));
CUDA_CHECK(cudaEventDestroy(stop));
return 0;
}
分析与预期结果：
-
模式1和模式2的执行时间约等于 H2D时间 + Kernel时间 + D2H时间 的总和。
-
模式3的执行时间理论上约等于
max(单次H2D时间, 单次Kernel时间, 单次D2H时间) + 流水线建立和排空时间。您将观察到模式3的时间显著少于前两者，因为它成功地用计算时间“隐藏”了大部分的数据传输时间。注意：使用cudaMallocHost分配锁页内存是实现计算和拷贝重叠的关键。
1.2.2 共享内存（Shared Memory）与访存优化
原理简介
共享内存是位于SM芯片上的一小块、由程序员显式管理的高速读写内存。其访问延迟远低于全局内存，与L1缓存相当。当一个算法中存在数据复用时（即同一数据被多个线程多次读取），共享内存就成为性能优化的关键。典型的例子是卷积或模板（Stencil）计算，计算一个点的输出需要其周围邻域的输入。朴素实现会导致每个线程都从全局内存中读取邻域数据，造成大量冗余的全局内存访问。
优化策略是：让一个线程块（Block）协作，先将计算所需的一个数据瓦片（Tile）（包含邻域所需的“光环”或“幽灵”区域）一次性、以合并的方式从全局内存加载到共享内存中。然后，块内所有线程都从极速的共享内存中读取数据进行计算。这样，每个全局内存地址在一个块中只被读取一次，极大地降低了对内存带宽的压力。
技术手册：可执行代码
以下程序实现了一个简单的一维平均模板（每个输出是对应输入及其左右邻居的平均值），并对比了朴素实现与使用共享内存优化的性能。
C++
// =================================================================
// 编译指令:
// nvcc -o shared_memory_stencil shared_memory_stencil.cu
// =================================================================
#include <iostream>
#include <vector>
#include <cuda_runtime.h>
#include <numeric>
#define CUDA_CHECK(call) do { \
cudaError_t err = call; \
if (err != cudaSuccess) { \
fprintf(stderr, "CUDA Error: %s:%d, ", __FILE__, __LINE__); \
fprintf(stderr, "code: %d, reason: %s\n", err, cudaGetErrorString(err)); \
exit(1); \
} \
} while(0)
// 内核1: 朴素实现，从全局内存重复读取
__global__ void naive_stencil(float* out, const float* in, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx > 0 && idx < n - 1) {
// 每个线程从全局内存读取3次
float left = in[idx - 1];
float center = in[idx];
float right = in[idx + 1];
out[idx] = (left + center + right) / 3.0f;
}
}
// 内核2: 使用共享内存优化
// 每个块处理 TILE_SIZE 个元素, 但需要加载 TILE_SIZE + 2 个元素 (左右各一个光环)
#define TILE_SIZE 256
#define BLOCK_SIZE (TILE_SIZE)
__global__ void shared_mem_stencil(float* out, const float* in, int n) {
// 声明共享内存，大小为块尺寸加上左右光环
__shared__ float tile[BLOCK_SIZE + 2];
int g_idx = blockIdx.x * BLOCK_SIZE + threadIdx.x; // 全局索引
int l_idx = threadIdx.x; // 块内索引
// 协作加载数据到共享内存
// 每个线程负责加载一个元素
if (g_idx < n) {
tile[l_idx + 1] = in[g_idx];
}
// 块内的第一个线程加载左侧光环
if (l_idx == 0 && g_idx > 0) {
tile[0] = in[g_idx - 1];
}
// 块内的最后一个线程加载右侧光环
if (l_idx == BLOCK_SIZE - 1 && g_idx < n - 1) {
tile[BLOCK_SIZE + 1] = in[g_idx + 1];
}
__syncthreads(); // 同步，确保所有数据已加载
// 从共享内存中计算
if (g_idx > 0 && g_idx < n - 1) {
float left = tile[l_idx];
float center = tile[l_idx + 1];
float right = tile[l_idx + 2];
out[g_idx] = (left + center + right) / 3.0f;
}
}
// CPU版本用于验证结果
void cpu_stencil(std::vector<float>& out, const std::vector<float>& in) {
for (size_t i = 1; i < in.size() - 1; ++i) {
out[i] = (in[i - 1] + in[i] + in[i + 1]) / 3.0f;
}
}
int main() {
int n = 1 << 24;
size_t bytes = n * sizeof(float);
std::vector<float> h_in(n);
for (int i = 0; i < n; ++i) h_in[i] = (float)i;
std::vector<float> h_out_naive(n, 0);
std::vector<float> h_out_shared(n, 0);
float *d_in, *d_out;
CUDA_CHECK(cudaMalloc(&d_in, bytes));
CUDA_CHECK(cudaMalloc(&d_out, bytes));
CUDA_CHECK(cudaMemcpy(d_in, h_in.data(), bytes, cudaMemcpyHostToDevice));
cudaEvent_t start, stop;
CUDA_CHECK(cudaEventCreate(&start));
CUDA_CHECK(cudaEventCreate(&stop));
float ms;
// --- 测试 Naive Kernel ---
int threadsPerBlock = 256;
int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
CUDA_CHECK(cudaEventRecord(start));
naive_stencil<<<blocksPerGrid, threadsPerBlock>>>(d_out, d_in, n);
CUDA_CHECK(cudaEventRecord(stop));
CUDA_CHECK(cudaEventSynchronize(stop));
CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));
std::cout << "Naive Stencil Kernel Time: " << ms << " ms" << std::endl;
CUDA_CHECK(cudaMemcpy(h_out_naive.data(), d_out, bytes, cudaMemcpyDeviceToHost));
// --- 测试 Shared Memory Kernel ---
// 注意块大小现在固定为BLOCK_SIZE
blocksPerGrid = (n + BLOCK_SIZE - 1) / BLOCK_SIZE;
CUDA_CHECK(cudaEventRecord(start));
shared_mem_stencil<<<blocksPerGrid, BLOCK_SIZE>>>(d_out, d_in, n);
CUDA_CHECK(cudaEventRecord(stop));
CUDA_CHECK(cudaEventSynchronize(stop));
CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));
std::cout << "Shared Memory Stencil Kernel Time: " << ms << " ms" << std::endl;
CUDA_CHECK(cudaMemcpy(h_out_shared.data(), d_out, bytes, cudaMemcpyDeviceToHost));
// --- 验证结果 ---
std::vector<float> h_out_cpu(n, 0);
cpu_stencil(h_out_cpu, h_in);
bool success = true;
for(int i=1; i < n-1; ++i) {
if (fabs(h_out_cpu[i] - h_out_shared[i]) > 1e-5) {
success = false;
break;
}
}
std::cout << "Verification: " << (success ? "SUCCESS" : "FAILED") << std::endl;
// 清理
cudaFree(d_in);
cudaFree(d_out);
cudaEventDestroy(start);
cudaEventDestroy(stop);
return 0;
}
分析与预期结果： 您将观察到，Shared Memory Stencil Kernel 的执行时间远少于 Naive Stencil Kernel。在朴素版本中，对于 n 个元素的计算，大约发生了 3*n 次全局内存读取。在优化版本中，每个元素平均只被从全局内存读取约 (TILE_SIZE + 2) / TILE_SIZE 次，约等于1次。这种对全局内存访问量的巨大削减，直接转化为性能的大幅提升。
1.2.3 CUDA Graphs：消除启动开销
原理简介
对于由一系列计算量小、执行速度快的核函数组成的工作流，CPU通过CUDA API逐个启动这些核函数的**启动开销（Launch Overhead）**本身可能会成为性能瓶颈。这个开销包括驱动程序调用、参数验证、将启动指令推送到GPU等步骤，通常需要几微秒（μs）的时间。如果一个核函数在GPU上只执行1μs，那么整个流程的绝大部分时间都耗费在了CPU的启动上。
CUDA Graphs 提供了一种“一次定义，多次运行”的机制来解决这个问题。它允许我们将一系列CUDA操作（核函数启动、内存拷贝等）捕获（Capture）到一个图中。一旦图被定义和实例化（Instantiate），我们就可以用一个API调用来**启动（Launch）**整个图的执行。这绕过了逐个核函数的启动路径，将CPU开销降至几乎为零，对于需要极低延迟的循环工作流至关重要。
技术手册：可执行代码
以下程序定义了一个包含三个微小内核的简单工作流，并对比了在循环中传统启动与使用CUDA Graph启动的CPU端总耗时。
C++
// =================================================================
// 编译指令:
// nvcc -o cuda_graphs cuda_graphs.cu
// =================================================================
#include <iostream>
#include <vector>
#include <chrono>
#include <cuda_runtime.h>
#define CUDA_CHECK(call) do { \
cudaError_t err = call; \
if (err != cudaSuccess) { \
fprintf(stderr, "CUDA Error: %s:%d, ", __FILE__, __LINE__); \
fprintf(stderr, "code: %d, reason: %s\n", err, cudaGetErrorString(err)); \
exit(1); \
} \
} while(0)
// 三个非常快的核函数
__global__ void add_one(float* data, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) data[idx] += 1.0f;
}
__global__ void mul_two(float* data, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) data[idx] *= 2.0f;
}
__global__ void sub_three(float* data, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) data[idx] -= 3.0f;
}
int main() {
int n = 1 << 10; // 使用小数据量以凸显启动开销
size_t bytes = n * sizeof(float);
int iterations = 10000;
std::vector<float> h_data(n, 1.0f);
float *d_data;
CUDA_CHECK(cudaMalloc(&d_data, bytes));
int threadsPerBlock = 128;
int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
// --- 模式1: 传统循环启动 ---
CUDA_CHECK(cudaMemcpy(d_data, h_data.data(), bytes, cudaMemcpyHostToDevice));
auto start_cpu = std::chrono::high_resolution_clock::now();
for (int i = 0; i < iterations; ++i) {
add_one<<<blocksPerGrid, threadsPerBlock>>>(d_data, n);
mul_two<<<blocksPerGrid, threadsPerBlock>>>(d_data, n);
sub_three<<<blocksPerGrid, threadsPerBlock>>>(d_data, n);
}
// 需要同步以确保所有GPU工作完成，才能准确计时
CUDA_CHECK(cudaDeviceSynchronize());
auto stop_cpu = std::chrono::high_resolution_clock::now();
auto duration_ms = std::chrono::duration_cast<std::chrono::milliseconds>(stop_cpu - start_cpu).count();
std::cout << "1. Traditional Loop CPU Time: " << duration_ms << " ms" << std::endl;
// --- 模式2: CUDA Graph ---
CUDA_CHECK(cudaMemcpy(d_data, h_data.data(), bytes, cudaMemcpyHostToDevice));
cudaGraph_t graph;
cudaGraphExec_t instance;
cudaStream_t stream;
CUDA_CHECK(cudaStreamCreate(&stream));
// -- 捕获Graph --
CUDA_CHECK(cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal));
// 将工作流定义一次
add_one<<<blocksPerGrid, threadsPerBlock, 0, stream>>>(d_data, n);
mul_two<<<blocksPerGrid, threadsPerBlock, 0, stream>>>(d_data, n);
sub_three<<<blocksPerGrid, threadsPerBlock, 0, stream>>>(d_data, n);
CUDA_CHECK(cudaStreamEndCapture(stream, &graph));
// -- 实例化Graph --
CUDA_CHECK(cudaGraphInstantiate(&instance, graph, NULL, NULL, 0));
// -- 循环启动Graph --
start_cpu = std::chrono::high_resolution_clock::now();
for (int i = 0; i < iterations; ++i) {
CUDA_CHECK(cudaGraphLaunch(instance, stream));
}
// 等待Graph流完成
CUDA_CHECK(cudaStreamSynchronize(stream));
stop_cpu = std::chrono::high_resolution_clock::now();
duration_ms = std::chrono::duration_cast<std::chrono::milliseconds>(stop_cpu - start_cpu).count();
std::cout << "2. CUDA Graph Loop CPU Time: " << duration_ms << " ms" << std::endl;
// 验证结果
float final_result;
CUDA_CHECK(cudaMemcpy(&final_result, d_data, sizeof(float), cudaMemcpyDeviceToHost));
// 预期结果: (1.0f + 1) * 2 - 3 = 1.0f
std::cout << "Verification: final value is " << final_result << " (expected 1.0)" << std::endl;
// --- 清理 ---
CUDA_CHECK(cudaGraphExecDestroy(instance));
CUDA_CHECK(cudaGraphDestroy(graph));
CUDA_CHECK(cudaStreamDestroy(stream));
CUDA_CHECK(cudaFree(d_data));
return 0;
}
分析与预期结果： 您将看到一个惊人的性能差异：CUDA Graph Loop CPU Time 可能只有几毫秒，而 Traditional Loop CPU Time 可能需要数百甚至数千毫秒。这是因为在传统循环中，CPU在 iterations * 3 次核函数启动上花费了大量时间。而在CUDA Graph模式下，CPU只在循环中调用了iterations次cudaGraphLaunch，这是一个开销极低的API。这证明了CUDA Graphs在消除重复工作流的启动开销方面的绝对优势。
扫一扫
49
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
