目录
一、CUDA编程
二、第一个CUDA程序
三、CUDA关键字
四、device管理
4.1 初始化
4.2 Runtime API查询GPU信息
4.3 决定最佳GPU
CUDA C++ 编程指南CUDA C++在线文档:CUDA C++ 编程指南
CUDA是并行计算的平台和类C编程模型,能很容易的实现并行算法。只需配备NVIDIA GPU,就可以在许多设备上运行并行程序
一、CUDA编程
CUDA编程允许程序执行在异构系统上,即CUP和GPU,二者有各自的存储空间,并由PCI-Express 总线区分开。注意二者术语上的区分:
- Host:CPU and itsmemory (host memory)
- Device: GPU and its memory (device memory)
device 可以独立于 host 进行大部分操作。当一个 kernel 启动后,控制权会立刻返还给 CPU 来执行其他额外的任务。所以CUDA编程是异步的。一个典型的CUDA程序包含由并行代码补足的串行代码,串行代码由host执行,并行代码在device中执行
host 端代码是标准C,device 是CUDA C代码。可以把所有代码放到一个单独的源文件,也可以使用多个文件或库。NVIDIA C编译器(nvcc)可以编译 host 和 device 端代码生成可执行程序
一个典型的CUDA程序结构包含五个主要步骤:
- 分配GPU空间
- 将数据从CPU端复制到GPU端
- 调用CUDA kernel来执行计算
- 计算完成后将数据从GPU拷贝回CPU
- 清理GPU内存空间
二、第一个CUDA程序
若是第一次使用CUDA,在Linux下可以使用下面的命令来检查CUDA编译器是否安装正确:
还需检查下机器上的GPU
以上输出显示仅有一个GPU显卡安装在机器上
CUDA 为许多常用编程语言提供扩展,如 C、C++、Python 和 Fortran 等语言。CUDA 加速程序的文件扩展名是.cu
下面包含两个函数,第一个函数将在 CPU 上运行,第二个将在 GPU 上运行
void CPUFunction()
{
printf("This function is defined to run on the CPU.\n");
}
__global__ void GPUFunction()
{
printf("This function is defined to run on the GPU.\n");
}
int main()
{
CPUFunction();
GPUFunction<<<1, 1>>>();
cudaDeviceSynchronize();
return 0;
}
- __global__ void GPUFunction()
__global__ 关键字表明以下函数将在 GPU 上运行并可全局调用
将在 CPU 上执行的代码称为主机代码,而将在 GPU 上运行的代码称为设备代码
注意返回类型为 void,使用 __global__ 关键字定义的函数要求返回 void 类型
- GPUFunction<<<1, 1>>>();
当调用要在 GPU 上运行的函数时,将此种函数称为已启动的核函数
启动核函数时,必须提供执行配置,即在向核函数传递任何预期参数之前使用 <<< … >>> 语法完成的配置。在宏观层面,程序员可通过执行配置为核函数启动指定线程层次结构,从而定义线程组(称为线程块)的数量,以及要在每个线程块中执行的线程数量
- cudaDeviceSynchronize();
与许多 C/C++ 代码不同,核函数启动方式为异步:CPU 代码将继续执行而无需等待核函数完成启动。调用 CUDA 运行时提供的函数 cudaDeviceSynchronize 将导致主机 (CPU) 代码暂作等待,直至设备 (GPU) 代码执行完成,才能在 CPU 上恢复执行
三、CUDA关键字
_global__关键字
__global__执行空间说明符将函数声明为内核。 其功能是:
- 在设备上执行
- 可从主机调用,可在计算能力为 3.2或更高的设备调用
- __global__ 函数必须具有 void 返回类型,并且不能是类的成员函数
- 对 global 函数的任何调用都必须指定其执行配置
- 对 global 函数的调用是异步的,这意味着其在设备完成执行之前返回
__device__关键字
- 在设备上执行
- 只能从设备调用
- __global__ 和 __device__ 执行空间说明符不能一起使用
__host__关键字
- 在主机上执行
- 只能从主机调用
- __global__ 和 __host__ 执行空间说明符不能一起使用
- __device__ 和 __host__ 执行空间说明符可以一起使用,此时该函数是为主机和设备编译的
四、device管理
4.1 初始化
当第一次调用任何CUDA运行时API(如cudaMalloc、cudaMemcpy等)时,CUDA Runtime会被初始化。这个初始化过程包括设置必要的内部数据结构、分配资源等,以便CUDA运行时能够管理后续的CUDA操作
每个CUDA设备都有一个与之关联的主上下文。主上下文是设备上的默认上下文,当没有显式创建任何上下文时,所有的CUDA运行时API调用都会在该主上下文中执行。主上下文包含了设备上的全局资源,如内存、纹理、表面等
开发者可以在程序启动时显式地指定哪个GPU成为"默认"设备。这个变化通常通过设置环境变量CUDA_VISIBLE_DEVICES或在程序中使用CUDA API(如cudaSetDevice)显式选择设备来实现。一旦选择了设备,随后的CUDA运行时初始化就会在这个指定的设备上创建主上下文
在没有显式指定设备的情况下,CUDA程序会默认在编号为0的设备(通常是第一个检测到的GPU)上执行操作
可以设置环境变量CUDA_VISIBLE_DEVICES-2来屏蔽其他GPU,这样只有GPU2能被使用。也可以使用CUDA_VISIBLE_DEVICES-2,3来设置多个GPU,其 device ID 分别为0和1
cudaDeviceReset
其作用是重置当前线程所关联的CUDA设备的状态,并释放该设备上所有已分配并未释放的资源
使用场景:
- 在程序结束时,调用该函数可以确保所有已分配的GPU资源都被正确释放,避免内存泄漏
- 若在程序的执行过程中遇到错误或需要中途退出,可释放已分配的资源,确保设备状态正确
- 在某些情况下,若设备状态出错(如由于之前的错误操作导致设备进入不可预测的状态),调用该函数可以尝试恢复设备到一个可用的状态
注意:
- 在调用该函数前,应确保所有已分配的设备内存和其他资源都已被正确地处理(如过cudaFree释放内存)。尽管其会释放这些资源,但最好还是在代码中显式地进行释放,以提高代码的可读性和可维护性
- 调用该函数后,当前线程与设备的关联关系可能会被重置。若需要继续使用设备,可能需要重新调用cudaSetDevice来设置当前线程要使用的设备
4.2 Runtime API查询GPU信息
cudaError_t cudaGetDeviceProperties(cudaDeviceProp *prop, int device);
GPU的信息被存放在cudaDeviceProp结构体中
#include <cuda_runtime_api.h>
#include <iostream>
#include <cmath>
using namespace std;
int main()
{
// 获取GPU数量
int deviceCount = 0;
cudaError_t errorId = cudaGetDeviceCount(&deviceCount);
if (errorId != cudaSuccess) {
printf("cudaGetDeviceCount returned %d\n-> %s\n", static_cast<int>(errorId), cudaGetErrorString(errorId));
printf("Result = FAIL\n");
exit(EXIT_FAILURE);
}
if (deviceCount == 0) {
printf("There are no available device(s) that support CUDA\n");
} else {
printf("Detected %d CUDA Capable device(s)\n", deviceCount);
}
// 指定第一个GPU
int device = 0;
cudaSetDevice(device);
// 获取GPU信息
cudaDeviceProp deviceProp;
cudaGetDeviceProperties(&deviceProp, device);
int driverVersion = 0, runtimeVersion = 0;
cudaDriverGetVersion(&driverVersion);
cudaRuntimeGetVersion(&runtimeVersion);
// 打印信息
printf(" Device %d: \"%s\"\n", device, deviceProp.name);
printf(" CUDA Driver Version / Runtime Version %d.%d / %d.%d\n", driverVersion/1000, (driverVersion%100)/10,runtimeVersion/1000, (runtimeVersion%100) / 10);
printf(" CUDA Capability Major/Minor version number: %d.%d\n", deviceProp.major, deviceProp.minor);
printf(" 全局内存总量: %.2f MBytes (%llu bytes)\n", (float)deviceProp.totalGlobalMem/(pow(1024.0,3)), static_cast<unsigned long long>(deviceProp.totalGlobalMem));
printf(" GPU Clock rate: %.0f MHz (%0.2f GHz)\n", deviceProp.clockRate * 1e-3f, deviceProp.clockRate * 1e-6f);
printf(" Memory Clock rate: %.0f Mhz\n", deviceProp.memoryClockRate * 1e-3f);
printf(" Memory Bus Width: %d-bit\n", deviceProp.memoryBusWidth);
if (deviceProp.l2CacheSize) {
printf(" L2 Cache Size: %d bytes\n",
deviceProp.l2CacheSize);
}
printf(" Max Texture Dimension Size (x,y,z) 1D=(%d), 2D=(%d,%d), 3D=(%d,%d,%d)\n",
deviceProp.maxTexture1D , deviceProp.maxTexture2D[0],
deviceProp.maxTexture2D[1],
deviceProp.maxTexture3D[0], deviceProp.maxTexture3D[1],
deviceProp.maxTexture3D[2]);
printf(" Max Layered Texture Size (dim) x layers 1D=(%d) x %d, 2D=(%d,%d) x %d\n",
deviceProp.maxTexture1DLayered[0], deviceProp.maxTexture1DLayered[1],
deviceProp.maxTexture2DLayered[0], deviceProp.maxTexture2DLayered[1],
deviceProp.maxTexture2DLayered[2]);
printf(" 常量内存总量: %lu bytes\n",deviceProp.totalConstMem);
printf(" 每个块的共享内存总量: %lu bytes\n",deviceProp.sharedMemPerBlock);
printf(" 每个块可用的寄存器总数: %d\n",deviceProp.regsPerBlock);
printf(" Warp size: %d\n", deviceProp.warpSize);
printf(" 每个多处理器的最大线程数: %d\n",deviceProp.maxThreadsPerMultiProcessor);
printf(" 每个块的最大线程数: %d\n",deviceProp.maxThreadsPerBlock);
printf(" 块各维度的最大尺寸: %d x %d x %d\n", deviceProp.maxThreadsDim[0], deviceProp.maxThreadsDim[1], deviceProp.maxThreadsDim[2]);
printf(" 网格每个维度的最大尺寸: %d x %d x %d\n", deviceProp.maxGridSize[0], deviceProp.maxGridSize[1], deviceProp.maxGridSize[2]);
printf(" Maximum memory pitch: %lu bytes\n", deviceProp.memPitch);
return 0;
}
4.3 决定最佳GPU
对于支持多GPU的系统,需从中选择一个来作为device,抉择出最佳计算性能GPU的一种方法就是由其拥有的处理器数量决定
int main()
{
int numDevices = 0;
cudaGetDeviceCount(&numDevices);
if (numDevices > 1)
{
int maxMultiprocessors = 0, maxDevice = 0;
for (int device=0; device < numDevices; ++device)
{
cudaDeviceProp props;
cudaGetDeviceProperties(&props, device);
if (maxMultiprocessors < props.multiProcessorCount) {
maxMultiprocessors = props.multiProcessorCount;
maxDevice = device;
}
}
cudaSetDevice(maxDevice);
}
return 0;
}