CUDA C 的編程接口主要包括一個(gè)對C語言的小擴(kuò)展集以及一個(gè)runtime庫寺酪。
核心的語言擴(kuò)展在于編程模型撩笆，就是kernel,block,grid,完整的擴(kuò)展參閱 C Language Extensions，任何源文件包括了擴(kuò)展內(nèi)容就必須用nvcc編譯，如果沒有使用擴(kuò)展語言沮协，可以在正常的.cpp文件中調(diào)用CUDA函數(shù)
在編譯流程中引入runtime，它讓host端通過C函數(shù)來控制device端內(nèi)存分配和釋放偷溺，兩者內(nèi)存之間的數(shù)據(jù)轉(zhuǎn)換哄孤，管理多device系統(tǒng)等。
runtime建立在底層的CUDA驅(qū)動(dòng)程序API之上胸私，驅(qū)動(dòng)程序API也提供了一些接口可以由應(yīng)用程序訪問厌处，例如CUDA contexts和CUDA module。但是大多數(shù)應(yīng)用程序不使用驅(qū)動(dòng)程序API岁疼，因?yàn)樗鼈儾恍枰@種額外的控制阔涉，并且在使用runtime時(shí)，上下文和模塊管理是隱式的捷绒，代碼可以更簡潔瑰排。
基礎(chǔ)
- 文件后綴為.cu 和 .cuh，CUDA 源文件和頭文件疙驾，其他的和正常的.cpp凶伙、.hpp一致。
- host端支持所有的C++語法它碎，但是device端支持部分C++語法
函數(shù)類型限定符Function Type Qualifiers
__host__
函數(shù)的調(diào)用和執(zhí)行都在host上
__device__
函數(shù)的調(diào)用和執(zhí)行都在device上
__global__
函數(shù)在host上調(diào)用函荣，在device上執(zhí)行显押，對于Compute Capability 3.2及以上也可以在device上調(diào)用，特性如下
- 對應(yīng)函數(shù)返回值必定為為void
- 調(diào)用
__global__函數(shù)必須有執(zhí)行配置（execution configuration）傻挂，即<<< Dg, Db, Ns, S >>>參數(shù) - 對應(yīng)修飾的函數(shù)為異步
變量類型修飾符Variable Type Qualifiers
指明在device中變量的內(nèi)存位置乘碑，在device端的代碼中沒有__device__, __shared__ and __constant__修飾符，則為自動(dòng)變量金拒，在寄存器register中兽肤。
__device__
變量在device中，常常和其他類型修飾符一起使用绪抛，如果只有該修飾符
- 在global內(nèi)存空間中
- 擁有整個(gè)應(yīng)用的生命周期
- 可以被grid中的所有線程訪問资铡，同時(shí)還host端還可以通過runtime庫訪問
(cudaGetSymbolAddress() / cudaGetSymbolSize() / cudaMemcpyToSymbol() / cudaMemcpyFromSymbol()).
__constant__
可以和__device__修飾符一同使用
- 在constant內(nèi)存空間中
- 擁有整個(gè)應(yīng)用的生命周期
- 可以被grid中的所有線程訪問，同時(shí)還host端還可以通過runtime庫訪問
__shared__
可以和__device__修飾符一同使用
- 在block中的共享內(nèi)存空間中
- 擁有block的聲明周期
- 只能被block中的線程訪問
__managed__
可以和__device__修飾符一同使用
- device和host都可以訪問
- 擁有整個(gè)應(yīng)用的生命周期
GPU 編程的步驟
- 在host上設(shè)定輸入數(shù)據(jù)
- 在host上分配內(nèi)存作為輸出 (malloc)
- 在GPU上分配輸入/輸出內(nèi)存 (cudaMalloc)
- 把host上的輸入復(fù)制到GPU(cudaMemcpy)
- 運(yùn)行GPU kernel
- 把輸出從GPU復(fù)制到host (cudaMemcpy)
device對應(yīng)的cudaMemcpy/cudaFree
cudaError_t cudaMalloc(void **devPtr, size_t size)
Returns
cudaSuccess, cudaErrorInvalidDevicePointer, cudaErrorInitializationError
cudaError_t cudaFree (void *devPtr)
Returns
cudaSuccess, cudaErrorInvalidDevicePointer, cudaErrorInitializationError
例子：
cudaError_t err = cudaMalloc((void **)&d_A, size);
err = cudaFree(d_A);
host下對應(yīng)的malloc/free
float *h_A = (float *)malloc(size);
free(h_A);
數(shù)據(jù)傳輸cudaMemcpy
cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind);
enum __device_builtin__ cudaMemcpyKind
{
cudaMemcpyHostToHost = 0, /**< Host -> Host */
cudaMemcpyHostToDevice = 1, /**< Host -> Device */
cudaMemcpyDeviceToHost = 2, /**< Device -> Host */
cudaMemcpyDeviceToDevice = 3, /**< Device -> Device */
cudaMemcpyDefault = 4 /**< Direction of the transfer is inferred from the pointer values. Requires unified virtual addressing */
};
cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice)
錯(cuò)誤類型cudaError_t
typedef __device_builtin__ enum cudaError cudaError_t;
enum __device_builtin__ cudaError
{
cudaSuccess = 0,
cudaErrorMissingConfiguration = 1,
cudaErrorMemoryAllocation = 2,
......
}幢码；
char* cudaGetErrorName(cudaError_t error)
char* cudaGetErrorString(cudaError_t error)
cudaError_t cudaGetLastError (void)
例子
if (err != cudaSuccess)
{
fprintf(stderr, "Failed to allocate device vector A (error code %s)!\n", cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
通用的CUDA Runtime API錯(cuò)誤處理
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
if (code != cudaSuccess)
{
fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
if (abort) exit(code);
}
}
調(diào)用Kernel
<<<blocksPerGrid, threadsPerBlock>>>設(shè)定執(zhí)行次數(shù)
int threadsPerBlock = 256;
int blocksPerGrid =(numElements + threadsPerBlock - 1) / threadsPerBlock;
printf("CUDA kernel launch with %d blocks of %d threads\n", blocksPerGrid, threadsPerBlock);
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, numElements);
err = cudaGetLastError();
__global__ void
vectorAdd(const float *A, const float *B, float *C, int numElements)
{
int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i < numElements)
{
C[i] = A[i] + B[i];
}
}
其他
cudaDeviceSynchronize() 阻塞直到device完成之前所有的任務(wù)
cudaDeviceReset() 銷毀分配的資源笤休，重置狀態(tài)，最后調(diào)用
#include "cuda_runtime.h"
