从简单CUDA开始

从简单CUDA开始

我们将从一个简单的C++程序开始，这个程序将两个包含一百万个元素的数组相加。下面是完整的代码。

#include<iostream>#include<math.h>// 函数：将两个数组的元素相加voidadd(intn,float*x,float*y){for(inti=0;i<n;i++)y[i]=x[i]+y[i];}intmain(void){intN=1<<20;// 1M elementsfloat*x=newfloat[N];float*y=newfloat[N];// 在CPU上初始化数组for(inti=0;i<N;i++){x[i]=1.0f;y[i]=2.0f;}// 在CPU上运行add函数add(N,x,y);// 检查结果中是否有错误floatmaxError=0.0f;for(inti=0;i<N;i++)maxError=fmax(maxError,fabs(y[i]-3.0f));std::cout<<"Max error: "<<maxError<<std::endl;// 释放内存delete[]x;delete[]y;return0;}

这个程序在CPU上初始化两个一百万元素的数组x和y，然后调用add函数将它们相加，并将结果存回数组y。最后，它检查结果是否正确。这是一个典型的内存带宽受限的计算，因为每个加法操作都需要两次内存读取和一次内存写入。

转换为CUDA内核

为了将这个函数转换为可以在GPU上运行的CUDA内核，我们需要做的主要改动是添加__global__修饰符。这个修饰符告诉CUDA编译器，这是一个我们将在GPU上调用的函数。

__global__voidadd(intn,float*x,float*y){for(inti=0;i<n;i++)y[i]=x[i]+y[i];}

在CUDA中，在CPU上运行的代码称为主机代码(host code)，而在GPU上运行的代码称为设备代码(device code)。__global__表示一个函数可以在设备上执行，并且可以从主机全局调用。我们通过一种特殊的语法<<<...>>>来从主机调用这个内核。

add<<<1,1>>>(N,x,y);

现在，add函数将在GPU上执行。但我们还需要处理内存。

CUDA中的内存分配

为了让GPU能够访问数据，我们需要在GPU内存中分配数据，或者使用一种特殊的CUDA特性，称为统一内存(Unified Memory)。统一内存创建了一个托管的内存池，CPU和GPU都可以访问。要使用统一内存，我们用cudaMallocManaged()替换new，用cudaFree()替换delete。

#include<iostream>#include<math.h>// CUDA内核：在GPU上执行__global__voidadd(intn,float*x,float*y){for(inti=0;i<n;i++)y[i]=x[i]+y[i];}intmain(void){intN=1<<20;float*x,*y;// 使用统一内存分配x和ycudaMallocManaged(&x,N*sizeof(float));cudaMallocManaged(&y,N*sizeof(float));// 在CPU上初始化数组for(inti=0;i<N;i++){x[i]=1.0f;y[i]=2.0f;}// 在GPU上运行内核add<<<1,1>>>(N,x,y);// 等待GPU完成计算cudaDeviceSynchronize();// 检查结果floatmaxError=0.0f;for(inti=0;i<N;i++)maxError=fmax(maxError,fabs(y[i]-3.0f));std::cout<<"Max error: "<<maxError<<std::endl;// 释放内存cudaFree(x);cudaFree(y);return0;}

我们还需要在调用内核后添加cudaDeviceSynchronize()。这是因为CUDA内核启动是异步的。主机代码在启动内核后会立即继续执行，而不会等待内核完成。cudaDeviceSynchronize()会阻塞主机线程，直到所有先前发出的CUDA任务（包括内核）完成。这确保了我们在CPU上访问结果之前，GPU已经完成了计算。

现在，我们可以编译并运行这个CUDA程序了。

# 编译CUDA代码nvcc add.cu -o add_cuda# 运行程序./add_cuda# 输出: Max error: 0.0

程序成功运行，但速度非常慢。这是因为内核只在单个GPU线程上运行。

性能分析

为了衡量性能，我们可以使用NVIDIA Nsight Systems命令行工具nsys。我们可以创建一个简单的bash脚本nsys_easy来简化这个过程。

#!/bin/bashnsys profile -t cuda --stats=true"$@"

现在，我们可以用这个脚本来分析我们的程序。

# 分析单线程CUDA版本./nsys_easy ./add_cuda

在NVIDIA T4 GPU上，单线程内核的执行时间大约是91.8毫秒。这比CPU版本慢得多。为了进行比较，我们可以测量CPU版本的性能。在我的系统上，CPU版本大约需要2.5毫秒。GPU版本慢了36倍！这是因为我们的CUDA内核只使用了一个GPU线程。

引入线程

为了利用GPU的并行能力，我们需要使用多个线程。我们可以修改内核启动配置，使用256个线程。

// 启动256个线程add<<<1,256>>>(N,x,y);

现在，我们需要修改内核代码，让每个线程处理一部分数据。我们可以使用内置变量threadIdx.x，它提供了当前线程在块内的索引。

__global__voidadd(intn,float*x,float*y){intindex=threadIdx.x;intstride=blockDim.x;for(inti=index;i<n;i+=stride)y[i]=x[i]+y[i];}

在这个修改后的内核中，每个线程从threadIdx.x开始，以blockDim.x（块中的线程总数，这里是256）为步长，处理数组中的元素。这种循环方式被称为grid-stride loop，它有几个好处：

1. 可扩展性：无论我们用多少线程启动内核，它都能正确工作。
2. 高效性：它能很好地合并内存访问。

再次运行性能分析，我们看到执行时间降到了2.05毫秒，比单线程版本快了45倍！

引入块

我们已经使用了256个线程，但现代GPU可以同时运行成千上万个线程。为了进一步扩展，我们可以使用多个线程块(thread blocks)。我们可以修改内核启动配置，使用多个块。

intblockSize=256;intnumBlocks=(N+blockSize-1)/blockSize;add<<<numBlocks,blockSize>>>(N,x,y);

这里，我们计算了需要的块数，以确保每个元素都至少被一个线程访问。现在，我们需要修改内核，使用blockIdx.x（当前块在grid中的索引）和gridDim.x（grid中的块总数）来计算全局索引。

__global__voidadd(intn,float*x,float*y){intindex=blockIdx.x*blockDim.x+threadIdx.x;intstride=blockDim.x*gridDim.x;for(inti=index;i<n;i+=stride)y[i]=x[i]+y[i];}

再次运行性能分析，执行时间降到了47.5微秒，比单块版本快了43倍，比单线程版本快了1932倍！这个版本的内核实现了265 GB/s的内存带宽，达到了NVIDIA T4 GPU峰值带宽(320 GB/s)的80%以上。

统一内存预取

虽然我们的内核现在非常快，但Nsight Systems的分析显示，仍然存在一些性能瓶颈。这是因为统一内存的按需页面迁移(on-demand page migration)。当GPU内核首次访问尚未迁移到GPU内存的数据时，会发生页面错误(page fault)，导致数据从CPU内存迁移到GPU内存。这些迁移会增加内核的执行时间。

为了解决这个问题，我们可以使用cudaMemPrefetchAsync()在内核启动前将数据**预取(prefetch)**到GPU。

// 在内核启动前预取数据到GPUcudaMemPrefetchAsync(x,N*sizeof(float),0,0);cudaMemPrefetchAsync(y,N*sizeof(float),0,0);add<<<numBlocks,blockSize>>>(N,x,y);cudaDeviceSynchronize();

预取操作是异步的，并且可以与数据传输和内核执行重叠。通过预取，我们告诉CUDA运行时，我们即将在GPU上使用这些数据，从而避免了运行时的页面错误。

总结

通过这个简单的向量加法示例，我们学习了如何将一个CPU程序转换为一个高度并行的CUDA程序。我们使用了统一内存来简化内存管理，并使用grid-stride loop来实现可扩展的并行内核。最后，我们使用预取来优化数据传输。

以下是不同版本的性能对比：

下一步

这只是CUDA编程的开始。要继续学习，我们推荐以下资源：

• CUDA Toolkit文档: Quick Start Guide, Programming Guide, Best Practices Guide
• 后续教程系列: 探索更高级的CUDA主题，如共享内存、原子操作和多维索引。
• NVIDIA DLI课程: Getting Started with Accelerated Computing in Modern CUDA C++