目錄

一、CUDA編程

二、第一個CUDA程序

三、CUDA關鍵字

四、device管理

4.1 初始化

4.2?Runtime API查詢GPU信息

4.3?決定最佳GPU

---

CUDA C++ 編程指南CUDA C++在線文檔：CUDA C++ 編程指南

CUDA是并行計算的平臺和類C編程模型，能很容易的實現(xiàn)并行算法。只需配備NVIDIA GPU，就可以在許多設備上運行并行程序

一、CUDA編程

CUDA編程允許程序執(zhí)行在異構系統(tǒng)上，即CUP和GPU，二者有各自的存儲空間，并由PCI-Express 總線區(qū)分開。注意二者術語上的區(qū)分：

• Host：CPU and itsmemory (host memory)
• Device: GPU and its memory (device memory)

device 可以獨立于 host 進行大部分操作。當一個 kernel 啟動后，控制權會立刻返還給 CPU 來執(zhí)行其他額外的任務。所以CUDA編程是異步的。一個典型的CUDA程序包含由并行代碼補足的串行代碼，串行代碼由host執(zhí)行，并行代碼在device中執(zhí)行

host 端代碼是標準C，device 是CUDA C代碼?？梢园阉写a放到一個單獨的源文件，也可以使用多個文件或庫。NVIDIA C編譯器（nvcc）可以編譯 host 和 device 端代碼生成可執(zhí)行程序

一個典型的CUDA程序結構包含五個主要步驟：

1. 分配GPU空間
2. 將數(shù)據(jù)從CPU端復制到GPU端
3. 調用CUDA kernel來執(zhí)行計算
4. 計算完成后將數(shù)據(jù)從GPU拷貝回CPU
5. 清理GPU內存空間

二、第一個CUDA程序

若是第一次使用CUDA，在Linux下可以使用下面的命令來檢查CUDA編譯器是否安裝正確：

還需檢查下機器上的GPU

以上輸出顯示僅有一個GPU顯卡安裝在機器上

CUDA 為許多常用編程語言提供擴展，如 C、C++、Python 和 Fortran 等語言。CUDA 加速程序的文件擴展名是.cu

下面包含兩個函數(shù)，第一個函數(shù)將在 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__ 關鍵字表明以下函數(shù)將在 GPU 上運行并可全局調用
將在 CPU 上執(zhí)行的代碼稱為主機代碼，而將在 GPU 上運行的代碼稱為設備代碼
注意返回類型為 void，使用 __global__ 關鍵字定義的函數(shù)要求返回 void 類型

• GPUFunction<<<1, 1>>>();

當調用要在 GPU 上運行的函數(shù)時，將此種函數(shù)稱為已啟動的核函數(shù)
啟動核函數(shù)時，必須提供執(zhí)行配置，即在向核函數(shù)傳遞任何預期參數(shù)之前使用 <<< … >>> 語法完成的配置。在宏觀層面，程序員可通過執(zhí)行配置為核函數(shù)啟動指定線程層次結構，從而定義線程組（稱為線程塊）的數(shù)量，以及要在每個線程塊中執(zhí)行的線程數(shù)量

• cudaDeviceSynchronize();

與許多 C/C++ 代碼不同，核函數(shù)啟動方式為異步：CPU 代碼將繼續(xù)執(zhí)行而無需等待核函數(shù)完成啟動。調用 CUDA 運行時提供的函數(shù) cudaDeviceSynchronize 將導致主機 (CPU) 代碼暫作等待，直至設備 (GPU) 代碼執(zhí)行完成，才能在 CPU 上恢復執(zhí)行

三、CUDA關鍵字

_global__關鍵字

__global__執(zhí)行空間說明符將函數(shù)聲明為內核。 其功能是：

• 在設備上執(zhí)行
• 可從主機調用，可在計算能力為 3.2或更高的設備調用
• __global__ 函數(shù)必須具有 void 返回類型，并且不能是類的成員函數(shù)
• 對 global 函數(shù)的任何調用都必須指定其執(zhí)行配置
• 對 global 函數(shù)的調用是異步的，這意味著其在設備完成執(zhí)行之前返回

__device__關鍵字

• 在設備上執(zhí)行
• 只能從設備調用
• __global__ 和 __device__ 執(zhí)行空間說明符不能一起使用

__host__關鍵字

• 在主機上執(zhí)行
• 只能從主機調用
• __global__ 和 __host__ 執(zhí)行空間說明符不能一起使用
• __device__ 和 __host__ 執(zhí)行空間說明符可以一起使用，此時該函數(shù)是為主機和設備編譯的

四、device管理

4.1 初始化

當?shù)谝淮握{用任何CUDA運行時API（如cudaMalloc、cudaMemcpy等）時，CUDA Runtime會被初始化。這個初始化過程包括設置必要的內部數(shù)據(jù)結構、分配資源等，以便CUDA運行時能夠管理后續(xù)的CUDA操作

每個CUDA設備都有一個與之關聯(lián)的主上下文。主上下文是設備上的默認上下文，當沒有顯式創(chuàng)建任何上下文時，所有的CUDA運行時API調用都會在該主上下文中執(zhí)行。主上下文包含了設備上的全局資源，如內存、紋理、表面等

開發(fā)者可以在程序啟動時顯式地指定哪個GPU成為"默認"設備。這個變化通常通過設置環(huán)境變量CUDA_VISIBLE_DEVICES或在程序中使用CUDA API（如cudaSetDevice）顯式選擇設備來實現(xiàn)。一旦選擇了設備，隨后的CUDA運行時初始化就會在這個指定的設備上創(chuàng)建主上下文

在沒有顯式指定設備的情況下，CUDA程序會默認在編號為0的設備（通常是第一個檢測到的GPU）上執(zhí)行操作

可以設置環(huán)境變量CUDA_VISIBLE_DEVICES-2來屏蔽其他GPU，這樣只有GPU2能被使用。也可以使用CUDA_VISIBLE_DEVICES-2,3來設置多個GPU，其 device ID 分別為0和1

cudaDeviceReset

其作用是重置當前線程所關聯(lián)的CUDA設備的狀態(tài)，并釋放該設備上所有已分配并未釋放的資源

使用場景：

1. 在程序結束時，調用該函數(shù)可以確保所有已分配的GPU資源都被正確釋放，避免內存泄漏
2. 若在程序的執(zhí)行過程中遇到錯誤或需要中途退出，可釋放已分配的資源，確保設備狀態(tài)正確
3. 在某些情況下，若設備狀態(tài)出錯（如由于之前的錯誤操作導致設備進入不可預測的狀態(tài)），調用該函數(shù)可以嘗試恢復設備到一個可用的狀態(tài)

注意：

1. 在調用該函數(shù)前，應確保所有已分配的設備內存和其他資源都已被正確地處理（如過cudaFree釋放內存）。盡管其會釋放這些資源，但最好還是在代碼中顯式地進行釋放，以提高代碼的可讀性和可維護性
2. 調用該函數(shù)后，當前線程與設備的關聯(lián)關系可能會被重置。若需要繼續(xù)使用設備，可能需要重新調用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數(shù)量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);}// 指定第一個GPUint 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(" 每個塊可用的寄存器總數(shù): %d\n",deviceProp.regsPerBlock);printf(" Warp size: %d\n", deviceProp.warpSize);printf(" 每個多處理器的最大線程數(shù): %d\n",deviceProp.maxThreadsPerMultiProcessor);printf(" 每個塊的最大線程數(shù): %d\n",deviceProp.maxThreadsPerBlock);printf(" 塊各維度的最大尺寸: %d x %d x %d\n", deviceProp.maxThreadsDim[0], deviceProp.maxThreadsDim[1], deviceProp.maxThreadsDim[2]);printf(" 網(wǎng)格每個維度的最大尺寸: %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的系統(tǒng)，需從中選擇一個來作為device，抉擇出最佳計算性能GPU的一種方法就是由其擁有的處理器數(shù)量決定

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;
}