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

Question

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

我的显卡是 GTX 770，内存为 4GB。我正在运行 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;

}

基本上，对于第二段代码，我可以访问主机上原始结构中的所有元素。然而，一旦我尝试启动内核函数，指针就变得不可访问(至少，从 gdb 测试)，并且整个对象的数据不可访问。在第一次内核启动后，我仍然能看到的边缘和 map 指针的唯一部分是它们各自的位置。

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

Answer 1

知道了!

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

检测:每当您在控制台(没有 X)中启动统一内存访问代码时，应该会出现类似如下的错误消息:

AMD-Vi:事件记录 [IO_PAGE_FAULT device=01:00.0 domain=0x0017 address=0x00000002d80d5000 flags=0x0010]

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

这是解决方案(假设您使用 GRUB):

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

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