CUDA C 的編程接口主要包括一個(gè)對(duì)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í)，上下文和模塊管理是隱式的，代碼可以更簡(jiǎn)潔。

基礎(chǔ)

1. 文件后綴為.cu 和 .cuh，CUDA 源文件和頭文件，其他的和正常的.cpp、.hpp一致。
2. 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í)行，對(duì)于Compute Capability 3.2及以上也可以在device上調(diào)用，特性如下

1. 對(duì)應(yīng)函數(shù)返回值必定為為void
2. 調(diào)用__global__函數(shù)必須有執(zhí)行配置（execution configuration），即<<< Dg, Db, Ns, S >>>參數(shù)
3. 對(duì)應(yīng)修飾的函數(shù)為異步

變量類型修飾符Variable Type Qualifiers

指明在device中變量的內(nèi)存位置，在device端的代碼中沒有__device__, __shared__ and __constant__修飾符，則為自動(dòng)變量，在寄存器register中。

__device__

變量在device中，常常和其他類型修飾符一起使用，如果只有該修飾符

1. 在global內(nèi)存空間中
2. 擁有整個(gè)應(yīng)用的生命周期
3. 可以被grid中的所有線程訪問，同時(shí)還host端還可以通過runtime庫訪問(cudaGetSymbolAddress() / cudaGetSymbolSize() / cudaMemcpyToSymbol() / cudaMemcpyFromSymbol()).

__constant__

可以和__device__修飾符一同使用

1. 在constant內(nèi)存空間中
2. 擁有整個(gè)應(yīng)用的生命周期
3. 可以被grid中的所有線程訪問，同時(shí)還host端還可以通過runtime庫訪問

__shared__

可以和__device__修飾符一同使用

1. 在block中的共享內(nèi)存空間中
2. 擁有block的聲明周期
3. 只能被block中的線程訪問

__managed__

可以和__device__修飾符一同使用

1. device和host都可以訪問
2. 擁有整個(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對(duì)應(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下對(duì)應(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"