CUDA学习（九十三）

简化GPU编程：
存储空间的统一意味着主机和设备之间不再需要显式存储器传输。 在托管内存空间中创建的任何分配都会自动迁移到需要的位置。
程序以两种方式之一分配托管内存：通过cudaMallocManaged（）例程，它在语义上类似于cudaMalloc（）; 或者通过定义一个全局的__managed__变量，这个变量在语义上类似于__device__变量。 这些文件的精确定义见本文后面。
在具有计算能力6.x的设备的支持平台上，Unified Memory将使应用程序能够使用默认系统分配器分配和共享数据。 这允许GPU在不使用特殊分配器的情况下访问整个系统虚拟内存。
以下代码示例说明了如何使用托管内存可以更改写入主机代码的方式。 首先，一个没有统一内存利益的简单程序：

__global__ void AplusB(int *ret, int a, int b) {
    ret[threadIdx.x] = a + b + threadIdx.x;
}
int main() {
    int *ret;
    cudaMalloc(&ret, 1000 * sizeof(int));
    AplusB << < 1, 1000 >> >(ret, 10, 100);
    int *host_ret = (int *)malloc(1000 * sizeof(int));
    cudaMemcpy(host_ret, ret, 1000 * sizeof(int), cudaMemcpyDefault);
    for (int i = 0; i<1000; i++)
        printf("%d: A+B = %d\n", i, host_ret[i]);
    free(host_ret);
    cudaFree(ret);
    return 0;
}

第一个示例将GPU上的两个数字与每个线程ID结合在一起，并将数值返回到数组中。 如果没有托管内存，则需要用于返回值的主机和设备端存储（在本例中为host_ret和ret），因为两者之间使用cudaMemcpy（）进行显式拷贝。
将此与程序的统一内存版本进行比较，该版本允许从主机直接访问GPU数据。 注意cudaMallocManaged（）例程，它返回一个有效来自主机和设备代码的指针。 这允许在没有单独的host_ret副本的情况下使用ret，极大地简化和减小了程序的大小。

__global__ void AplusB(int *ret, int a, int b) {
    ret[threadIdx.x] = a + b + threadIdx.x;
}
int main() {
    int *ret;
    cudaMalloc(&ret, 1000 * sizeof(int));
    AplusB << < 1, 1000 >> >(ret, 10, 100);
    int *host_ret = (int *)malloc(1000 * sizeof(int));
    cudaMemcpy(host_ret, ret, 1000 * sizeof(int), cudaMemcpyDefault);
    for (int i = 0; i<1000; i++)
        printf("%d: A+B = %d\n", i, host_ret[i]);
    free(host_ret);
    cudaFree(ret);
    return 0;
}

最后，语言集成允许直接引用GPU声明的__managed__变量，并在使用全局变量时进一步简化程序。

__device__ __managed__ int ret[1000];
__global__ void AplusB(int a, int b) {
    ret[threadIdx.x] = a + b + threadIdx.x;
}
int main() {
    AplusB << < 1, 1000 >> >(10, 100);
    cudaDeviceSynchronize();
    for (int i = 0; i<1000; i++)
        printf("%d: A+B = %d\n", i, ret[i]);
    return 0;
}

请注意，缺少显式的cudaMemcpy（）命令以及返回数组ret在CPU和GPU上都可见的事实。
对主机和设备之间的同步值得评论。 请注意，在非托管示例中，同步cudaMemcpy（）例程用于同步内核（即等待它完成运行）并将数据传输到主机。 统一内存示例不会调用cudaMemcpy（），因此在主机程序可以安全地使用GPU输出之前需要明确的cudaDeviceSynchronize（）。
这里的另一种方法是设置环境变量CUDA_LAUNCH_BLOCKING = 1，确保所有内核的启动都是同步完成的。 这通过消除所有显式同步来简化代码，但显然对整个执行行为具有更广泛的影响。