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

Question

我最近一直在玩CUDA，并希望尝试统一内存模型。我试着玩样例代码，奇怪的是，当启动内核时，似乎没有值在更新。修改主机上的统一数据工作得很好，但是启动的内核根本不会修改统一数据。

我的卡是GTX 770与4GB的内存。我正在运行Arch Linux，内核3.14-2，使用GCC 4.8来编译我的示例。我将计算arch设置为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("%sn", error);
    x[0] = 0;
    x[1] = 0;
    kernel<<<1, 2>>>(x);
    cudaDeviceSynchronize();
    printf("result = %dn", 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: %dn", 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测试是这样)，并且整个对象的数据都不可访问。在第一次内核启动后，我仍然可以看到的边缘和映射指针的唯一部分是它们各自的位置。

任何帮助都将非常感激!非常感谢!

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帮我缩小了问题的范围!

Answer 2

我做了一个与Matthew Daiter提到的类似的过程。我删除了IOMMU选项，但从BIOS中执行。这个工作很完美!!