CUDA学习（七）

设备内存：
正如异构编程中提到的那样，CUDA编程模型假设一个由主机和设备组成的系统，每个系统都有自己独立的内存。 内核在设备内存之外运行，因此runtime提供了分配，解除分配和复制设备内存的功能，以及在主机内存和设备内存之间传输数据。
设备内存可以分配为线性内存或CUDA数组。
CUDA数组是不透明的内存布局，针对纹理获取进行了优化。 它们在纹理和表面记忆中被描述。
线性内存在设备上存在于40位地址空间中，因此分开分配的实体可以通过指针相互引用，例如在二叉树。
线性内存通常使用cudaMalloc（）进行分配，并使用cudaFree（）进行释放，主机内存和设备内存之间的数据传输通常使用cudaMemcpy（）完成。 在内核的向量附加代码示例中，向量需要从主机内存复制到设备内存：

// 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
    ...
}

线性内存也可以通过cudaMallocPitch（）和cudaMalloc3D（）来分配。 建议将这些功能用于分配2维或者3维的数组，因为它确保分配适当地填充以满足器件存储器访问中描述的对齐要求，因此在访问行地址或在2D数组与其他区域之间执行复制时确保最佳性能 的设备内存（使用cudaMemcpy2D（）和cudaMemcpy3D（）函数）。返回的pitch（或stride）必须用于访问数组元素。以下代码示例分配宽x高二维浮点值数组，并显示如何遍历设备代码中的数组元素：

// Host code
int width = 64, height = 64;
float* devPtr;
size_t pitch;
cudaMallocPitch(&devPtr, &pitch,
    width * sizeof(float), height);
MyKernel << <100, 512 >> >(devPtr, pitch, width, height);
// Device code
__global__ void MyKernel(float* devPtr,
    size_t pitch, int width, int height)
{
    for (int r = 0; r < height; ++r) {
        float* row = (float*)((char*)devPtr + r * pitch);
        for (int c = 0; c < width; ++c) {
            float element = row[c];
        }
    }
}

以下代码示例分配宽度x高度x深度3D浮点数值数组，并显示如何遍历设备代码中的数组元素：

// Host code
int width = 64, height = 64, depth = 64;
cudaExtent extent = make_cudaExtent(width * sizeof(float),
    height, depth);
cudaPitchedPtr devPitchedPtr;
cudaMalloc3D(&devPitchedPtr, extent);
MyKernel << <100, 512 >> >(devPitchedPtr, width, height, depth);
// Device code
__global__ void MyKernel(cudaPitchedPtr devPitchedPtr,
    int width, int height, int depth)
{
    char* devPtr = devPitchedPtr.ptr;
    size_t pitch = devPitchedPtr.pitch;
    size_t slicePitch = pitch * height;
    for (int z = 0; z < depth; ++z) {
        char* slice = devPtr + z * slicePitch;
        for (int y = 0; y < height; ++y) {
            float* row = (float*)(slice + y * pitch);
            for (int x = 0; x < width; ++x) {
                float element = row[x];
            }
        }
    }
}

参考手册列出了用cudaMalloc（）分配的线性存储器，用cudaMallocPitch（）或cudaMalloc3D（）分配的线性存储器，CUDA数组和分配给在全局或常量存储空间中声明的变量的存储器之间复制存储器的所有函数。
以下代码示例说明了通过运行时API访问全局变量的各种方法：

__constant__ float constData[256];
float data[256];
cudaMemcpyToSymbol(constData, data, sizeof(data));
cudaMemcpyFromSymbol(data, constData, sizeof(data));
__device__ float devData;
float value = 3.14f;
cudaMemcpyToSymbol(devData, &value, sizeof(float));
__device__ float* devPointer;
float* ptr;
cudaMalloc(&ptr, 256 * sizeof(float));
cudaMemcpyToSymbol(devPointer, &ptr, sizeof(ptr));

cudaGetSymbolAddress（）用于检索指向为全局内存空间中声明的变量分配的内存的地址。 分配的内存大小是通 cudaGetSymbolSize（）获得的。