CUDA内存使用

CUDA线程可以在运行过程中从多中内存空间访问数据，分为三个层次：

1，局部内存：每一个线程有其私有的局部内存。

2，共享内存：每一个线程块（thread block）有一个共享内存，可以被该线程块中的所有线程访问。

3，全局内存：所有的线程都能访问。

此外还有两个能被所有线程都访问的只读内存：constant 和texture内存，如下图所示：



CUDA编程模型假设系统由一个主机（CPU）和一个设备（GPU）组成，它们各自有自己的内存。在核的执行过程中，需要提供函数用于分配内存，释放内存，复制设备内存以及数据在主机内存和设备内存之间的传输。

设备内存可以以线性内存或者CUDA数组的方式进行分配，CUDA数组主要用于纹理获取。

线性内存通常使用 cudaMalloc()进行分配，用cudaFree()对其释放，主机内存和设备内存之间的数据传输通过cudaMemcpy()实现。

cudaMallocPitch()和cudaMalloc3D()分别用于二维和三维数组的内存分配。

cudaMemcpy2D()和cudaMemcpy3D()分别用于二维和三维数组的内存复制。

以向量的加法为例，在调用核之前，向量的值需要从主机内存复制到设备内存，计算结束后，再将计算结果从设备内存复制到主机内存。

// Device code
__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(&d_A, size);
    float* d_B;
    cudaMalloc(&d_B, size);
    float* d_C;
    cudaMalloc(&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;
    int blocksPerGrid =
            (N + threadsPerBlock - 1) / threadsPerBlock;
    VecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

    // Copy result from device memory to host memory
    // h_C contains the result in host memory
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

    // Free device memory
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
            
    // Free host memory
    ...
}