CUDA 统一内存工作（具体来说，cudaMallocManaged();）

Posted 2023-04-15

技术标签:

【中文标题】CUDA 统一内存工作（具体来说，cudaMallocManaged();）【英文标题】：CUDA Unified Memory Working (in specific, cudaMallocManaged();) 【发布时间】：2014-07-01 23:55:36 【问题描述】：

我最近一直在玩 CUDA，并希望尝试统一内存模型。我尝试使用示例代码，但奇怪的是，在启动内核时，似乎没有更新任何值。从主机修改统一数据可以正常工作，但启动的内核根本不会修改统一数据。

我的卡是具有 4GB 内存的 GTX 770。我正在运行 Arch Linux，内核 3.14-2，使用 GCC 4.8 来编译我的示例。我将计算架构设置为 sm_30，并激活 -m64 标志

这是我正在玩的一个示例。 X[0] 和 X[1] 始终计算为 0，即使在内核启动时也是如此。

#include<stdio.h>
#include <cuda.h>

__global__ void kernel(int* x)
    x[threadIdx.x] = 2;


int main()
    int* x;
    cudaMallocManaged(&x, sizeof(int) * 2);
    cudaError_t error = cudaGetLastError();
    printf("%s\n", error);
    x[0] = 0;
    x[1] = 0;

    kernel<<<1, 2>>>(x);
    cudaDeviceSynchronize();

    printf("result = %d\n", x[1]);

    cudaFree(x);
    return 0;

另一个示例是这样的：

__global__ void adjacency_map_init_gpu(adjacency_map_t* map)
    int row = threadIdx.y + blockIdx.y * blockDim.y;
    int col = threadIdx.x + blockIdx.x * blockDim.x;

    int i = row * map->width + col;

    max(i, 0);
    min(i, map->width * map->height);

    map->connections[i] = 0;


__global__ void adjacency_map_connect_gpu(edge_t* edges, int num_edges, adjacency_map_t* map)

    int i = threadIdx.x + (((gridDim.x * blockIdx.y) + blockIdx.x)*blockDim.x);

    max(i, 0);
    min(i, num_edges);

    int n_start = edges[i].n_start;
    int n_end = edges[i].n_end;

    int map_index = n_start * map->width + n_end;
    map->connections[map_index] = 1;
    printf("%d new value: %d\n", map_index, map->connections[map_index]);


adjacency_map_t* adjacency_map_init(int num_nodes, edge_t* edges, int num_edges)
    adjacency_map_t *map;// = (adjacency_map_t*)malloc(sizeof(adjacency_map_t));
    cudaMallocManaged(&map, sizeof(adjacency_map_t));
    cudaMallocManaged(&(map->connections), num_nodes * num_nodes * sizeof(int));
    //map->connections = (int*)malloc(sizeof(int) * num_nodes * num_nodes);

    map->width = num_nodes;
    map->height = num_nodes;

    map->stride = 0;

    //GPU stuff
//    adjacency_map_t *d_map;
//    int* d_connections;

//    cudaMalloc((void**) &d_map, sizeof(adjacency_map_t));
//    cudaMalloc((void**) &d_connections, num_nodes * num_nodes * sizeof(int));

//    cudaMemcpy(d_map, map, sizeof(adjacency_map_t), cudaMemcpyHostToDevice);
//    cudaMemcpy(d_connections, map->connections, num_nodes * num_nodes, cudaMemcpyHostToDevice);
//cudaMemcpy(&(d_map->connections), &d_connections, sizeof(int*), cudaMemcpyHostToDevice);

//    edge_t* d_edges;
//    cudaMalloc((void**) &d_edges, num_edges * sizeof(edge_t));
//    cudaMemcpy(d_edges, edges, num_edges * sizeof(edge_t), cudaMemcpyHostToDevice);

adjacency_map_init_gpu<<<1, 3>>>(map);
cudaDeviceSynchronize();
//adjacency_map_connect_gpu<<<1, 3>>>(edges, num_edges, map);

cudaDeviceSynchronize();

//    cudaMemcpy(map, d_map, sizeof(adjacency_map_t), cudaMemcpyDeviceToHost);
//Synchronize everything
//    cudaFree(map);
//    cudaFree(edges);

return map;

基本上，我可以访问主机上原始结构中的所有元素，以获取第二个 sn-p 代码。但是，一旦我尝试启动内核函数，指针就会变得不可访问（至少，从 gdb 测试），并且整个对象的数据都不可访问。在第一次内核启动后，我仍然可以看到的唯一部分边缘和映射指针是它们各自的位置。

任何帮助将不胜感激！非常感谢！

【问题讨论】：

您的第一个代码适用于我在 CUDA 6/RHEL 6.2 上。也就是说，我得到了result = 2 printed。您应该将 proper cuda error checking 添加到所有 cuda API 调用和内核调用中。 Arch Linux/gcc4.8 看起来不像a supported distro combination。首先测试没有gdb 的代码，然后使用cuda-memcheck。还提供您的确切编译命令行。重新编译 gcc 到 4.8（Arch Linux 的版本现在是 4.9）。使用本地版本的 GCC 进行编译。这是我使用的命令： /opt/cuda/bin/nvcc -ccbin /opt/gcc/usr/local/bin/gcc mem.cu -arch=compute_30 -code=sm_30 -g -m64 （mem.cu 是名称第一个 sn-p 的文件。）添加了您链接到的错误检查。包装 cudaMallocManaged()、cudaPeekAtLastError() 和 cudaDeviceSynchronize() 时没有打印任何内容。我不知道是什么问题。我能提供的唯一建议是切换到受支持的发行版。正如我已经指出的那样，当我在受支持的发行版/配置上运行时，您的代码示例（至少是第一个）对我来说很好。我在第一条评论中提到了我的 linux 版本 (RHEL 6.2) 驱动程序版本是 CUDA 6 linux .run 安装程序 (331.62) 附带的驱动程序。 xorg.conf 与此问题/问题无关。我使用的大多数系统都没有运行 X。 【参考方案1】：

知道了！

原来是启用了 IOMMU 内核选项的问题。我的主板 GIGABYTE 990-FXAUD3 似乎在 GPU 和 CPU 之间的 IOMMU 出现错误。

检测：每当您在控制台中启动统一内存访问代码（不带 X）时，都会出现类似以下的错误消息：

AMD-Vi：记录事件 [IO_PAGE_FAULT 设备=01:00.0 域=0x0017 地址=0x00000002d80d5000 标志=0x0010]

向下滚动页面。屏幕右上角也可能有一些变色（至少对我来说是这样）。

这是解决方案（假设您使用 GRUB）：

打开 /etc/default/grub，并为 GRUB_CMDLINE_LINUX_DEFAULT="" 行在引号内添加选项 iommu=soft。

希望这对人们有所帮助！非常感谢 Robert Crovella 帮助我缩小问题范围！

【讨论】：

【参考方案2】：

我做了一个与 Matthew Daiter 提到的类似的程序。我删除了 IOMMU 选项，但从 Bios 中删除了。这工作完美！

【讨论】：