CUDA C 的编程接口主要包括一个对C语言的小扩展集以及一个runtime库。
核心的语言扩展在于编程模型，就是kernel,block,grid,完整的扩展参阅 C Language Extensions，任何源文件包括了扩展内容就必须用nvcc编译，如果没有使用扩展语言，可以在正常的.cpp文件中调用CUDA函数
在编译流程中引入runtime，它让host端通过C函数来控制device端内存分配和释放，两者内存之间的数据转换，管理多device系统等。
runtime建立在底层的CUDA驱动程序API之上，驱动程序API也提供了一些接口可以由应用程序访问，例如CUDA contexts和CUDA module。但是大多数应用程序不使用驱动程序API，因为它们不需要这种额外的控制，并且在使用runtime时，上下文和模块管理是隐式的，代码可以更简洁。
基础
- 文件后缀为.cu 和 .cuh，CUDA 源文件和头文件，其他的和正常的.cpp、.hpp一致。
- host端支持所有的C++语法，但是device端支持部分C++语法
函数类型限定符Function Type Qualifiers
__host__
函数的调用和执行都在host上
__device__
函数的调用和执行都在device上
__global__
函数在host上调用，在device上执行，对于Compute Capability 3.2及以上也可以在device上调用，特性如下
- 对应函数返回值必定为为void
- 调用
__global__函数必须有执行配置（execution configuration），即<<< Dg, Db, Ns, S >>>参数 - 对应修饰的函数为异步
变量类型修饰符Variable Type Qualifiers
指明在device中变量的内存位置，在device端的代码中没有__device__, __shared__ and __constant__修饰符，则为自动变量，在寄存器register中。
__device__
变量在device中，常常和其他类型修饰符一起使用，如果只有该修饰符
- 在global内存空间中
- 拥有整个应用的生命周期
- 可以被grid中的所有线程访问，同时还host端还可以通过runtime库访问
(cudaGetSymbolAddress() / cudaGetSymbolSize() / cudaMemcpyToSymbol() / cudaMemcpyFromSymbol()).
__constant__
可以和__device__修饰符一同使用
- 在constant内存空间中
- 拥有整个应用的生命周期
- 可以被grid中的所有线程访问，同时还host端还可以通过runtime库访问
__shared__
可以和__device__修饰符一同使用
- 在block中的共享内存空间中
- 拥有block的声明周期
- 只能被block中的线程访问
__managed__
可以和__device__修饰符一同使用
- device和host都可以访问
- 拥有整个应用的生命周期
GPU 编程的步骤
- 在host上设定输入数据
- 在host上分配内存作为输出 (malloc)
- 在GPU上分配输入/输出内存 (cudaMalloc)
- 把host上的输入复制到GPU(cudaMemcpy)
- 运行GPU kernel
- 把输出从GPU复制到host (cudaMemcpy)
device对应的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下对应的malloc/free
float *h_A = (float *)malloc(size);
free(h_A);
数据传输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)
错误类型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错误处理
#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);
}
}
调用Kernel
<<<blocksPerGrid, threadsPerBlock>>>设定执行次数
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完成之前所有的任务
cudaDeviceReset() 销毁分配的资源，重置状态，最后调用
#include "cuda_runtime.h"
