统一内存和结构与数组

Question

我在 CUDA 上有一个很大的结构数组结构，它对我的应用程序来说是常量和只读的。一个非常简单的例子是

struct Graph{
    Node * nodes;
    int nNode;
}
struct Node{
   int* pos;
   int nPos;
}

我的内核需要导航此图并查询它。如您所知，使用 cudaMalloc 和 cudaMemcpy 将此结构复制到 GPU 内存只是大量代码，统一内存应该消除对它的需求。

在我的代码中，我在 CPU 中生成了图表，然后为了测试，我设计了以下内核

__global__ void testKernel(const Graph graph,int * d_res){
    d_res[0]=graph.nNode;


};

被称为：

// using malloc for testing to make sure I know what I am doing
int * d_res,* h_res;
cudaMalloc((void **)&d_res,sizeof(int));
h_res=(int*)malloc(sizeof(int));

testKernel<<<1,1>>>(graph,d_res);

gpuErrchk( cudaPeekAtLastError() );
gpuErrchk(cudaMemcpy(h_res,d_res,sizeof(int),cudaMemcpyDeviceToHost));

使用错误检查 from here。

当我如图所示使用 testKernel 时，它工作正常，但如果我将内核更改为：

__global__ void testKernel(const Graph graph,int * d_res){
    d_res[0]=graph.nodes[0].nPos;

};

我遇到非法内存访问错误。

是不是统一内存没有正确处理这类数据？有没有办法确保我可以避免将所有显式副本写入 GPU 内存？

完整 MCVE：

#include <algorithm>
#include <cuda_runtime_api.h>
#include <cuda.h>
typedef struct node{
    int* pos;
    int nPos;
}Node;
typedef struct Graph{
    Node * nodes;
    int nNode;
}Graph;


#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort = true)
{
    if (code != cudaSuccess)
    {
        fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
        if (abort) exit(code);
    }
}

__global__ void testKernel(const Graph graph, int * d_res){
    d_res[0] = graph.nNode;
    // d_res[0]=graph.nodes[0].nPos; // Not working

};



int main(void){

    // fake data, this comes from another process
     Graph graph;
    graph.nodes = (Node*)malloc(2*sizeof(Node));
    graph.nNode = 2;
    for (int i = 0; i < 2; i++){


    // They can have different sizes in the original code
    graph.nodes[i].pos = (int*)malloc(3 * sizeof(int));
    graph.nodes[i].pos[0] = 0;
    graph.nodes[i].pos[1] = 1;
    graph.nodes[i].pos[2] = 2;

    graph.nodes[i].nPos = 3;

}



printf("%d\n", graph.nNode); // Change to the kernel variable for comparison
int * d_res, *h_res;
cudaMalloc((void **)&d_res, sizeof(int));
h_res = (int*)malloc(sizeof(int));
testKernel << <1, 1 >> >(graph, d_res);
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaMemcpy(h_res, d_res, sizeof(int), cudaMemcpyDeviceToHost));

printf("%d", h_res[0]);
return 0;
}

Answer 1

您的代码未使用 CUDA 统一内存。 UM 无论如何都不是 "automatic"。它需要特定的编程步骤才能利用它，并且具有特定的系统要求。

所有这些都包含在 UM section of the programming guide 中。

Is there a way to make sure I can avoid writing all the explicit copies to GPU memory?

正确使用 UM 应该可以做到这一点。这是一个完整的例子。我所做的唯一一件事就是将主机代码中的 malloc 操作机械地转换为等效的 cudaMallocManaged 操作。

$ cat t1389.cu
#include <algorithm>
#include <stdio.h>

typedef struct node{
    int* pos;
    int nPos;
}Node;
typedef struct Graph{
    Node * nodes;
    int nNode;
}Graph;


#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort = true)
{
    if (code != cudaSuccess)
    {
        fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
        if (abort) exit(code);
    }
}

__global__ void testKernel(const Graph graph, int * d_res){
    d_res[0] = graph.nNode;
     d_res[0]=graph.nodes[0].nPos; // Not working

};



int main(void){

    // fake data, this comes from another process
     Graph graph;
    cudaMallocManaged(&(graph.nodes), 2*sizeof(Node));
    graph.nNode = 2;
    for (int i = 0; i < 2; i++){


    // They can have different sizes in the original code
    cudaMallocManaged(&(graph.nodes[i].pos), 3 * sizeof(int));
    graph.nodes[i].pos[0] = 0;
    graph.nodes[i].pos[1] = 1;
    graph.nodes[i].pos[2] = 2;

    graph.nodes[i].nPos = 3;

}



printf("%d\n", graph.nNode); // Change to the kernel variable for comparison
int * d_res, *h_res;
cudaMalloc((void **)&d_res, sizeof(int));
h_res = (int*)malloc(sizeof(int));
testKernel << <1, 1 >> >(graph, d_res);
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaMemcpy(h_res, d_res, sizeof(int), cudaMemcpyDeviceToHost));

printf("%d", h_res[0]);
return 0;
}
$ nvcc t1389.cu -o t1389
$ cuda-memcheck ./t1389
========= CUDA-MEMCHECK
2
3========= ERROR SUMMARY: 0 errors
$

UM 有许多记录在案的系统要求。我不打算在这里一一列举。首先你需要一个 cc3.0 或更高版本的 GPU。您的 MCVE 不包含任何标准错误检查，我也没有尝试添加它。但是如果你仍然对这段代码有问题，一定要使用正确的 CUDA 错误检查和运行它与 cuda-memcheck.

如果您的整个数据结构（包括嵌入式指针）是使用普通主机分配器分配的，并且您无法控制它，那么您将无法在 UM 机制中直接使用它，而无需进行某种排序涉及复制。如上面链接的编程指南部分的 K.1.6 节所述，这里的例外情况是在 IBM Power9 系统上。

在尝试将主机分配器（例如 malloc）与 UM 一起使用之前，您应该首先测试 pageableMemoryAccessUsesHostPageTables 属性，如该部分所述。

属性当前不会在任何系统上设置，除非正确配置的 IBM Power9 系统。目前没有 x86 系统有这个属性 set/available.

统一内存和结构与数组

Unified memory and struct with arrays

cuda

unified-memory