3.2 CUDA C
CUDA C为熟悉C语言的用户提供了一个简单途径，让他们能够轻易的写出能够在设备上执行的程序。
CUDA C包含了一个C语言的最小扩展集和一个运行时库。语言核心扩展在第二章已经介绍了。本节继续介绍运行时。所有扩展的完整的描述可在附录B找到，CUDA运行时的完整描述可在CUDA参考手册中找到。
cudart动态库是运行时的实现，它所有的入口点前缀都是cuda。
运行时没有显式的初始化函数;在初次调用运行时函数(更精确地，不在参考手册中设备和版本管理节中的任何函数)时初始化。在计算运行时函数调用时间和解析初次调用运行时产生的错误码时必须牢记这点。
一旦运行时在主机线程中初始化，在主机线程中通过一些运行时函数调用分配的任何资源(存储器，流，事件等)只在当前主机线程的上下文中有效。因此只有在这个主机线程中调用的运行时函数(存储器拷贝，内核发射等)才能操作这些资源。这是因为CUDA上下文(参见3.3.1节)作为初始化的一部分建立且成为主机线程的当前上下文，且不能成为其它主机线程的当前上下文。
在多设备的系统中，内核默认在0号设备上执行，详见3.2.3节。
3.2.1 设备存储器
正如2.4节所提到的，CUDA编程模型假定系统包含主机和设备，它们各有自己独立的存储器。内核不能操作设备存储器，所以运行时提供了分配，释放，拷贝设备存储器和在设备和主机间传输数据的函数。
设备存储器可被分配为线性存储器或CUDA数组。
CUDA数组是不透明的存储器层次，为纹理获取做了优化。它们的细节在3.2.4节。
计算能力1.x的设备，其线性存储器存在于32位地址空间内，计算能力2.0的设备，其线性存储器存在于40位地址空间内，所以独立分配的实体能够通过指针引用，如，二叉树。
典型地，线性存储器使用cudaMalloc()分配，通过cudaFree()释放，使用cudaMemcpy()在设备和主机间传输。在2.1节的向量加法代码中，向量要从主机存储器复制到设备存储器：
__global__ void VecAdd(float* A, float* B, float* C, int N) {
int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i < N)
C[i] = A[i] + B[i];
}
// Host code
int main(){
int N = ...;
size_t size = N * sizeof(float);
// Allocate input vectors h_A and h_B in host memory
float* h_A = (float*)malloc(size);
float* h_B = (float*)malloc(size);
// Initialize input vectors
...
// Allocate vectors in device memory
float* d_A;
cudaMalloc((void**)&d_A, size);
float* d_B;
cudaMalloc((void**)&d_B, size);
float* d_C;
cudaMalloc((void**)&d_C, size);
// Copy vectors from host memory to device memory
cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
// Invoke kernel
int threadsPerBlock = 256;
