【发布时间】：2018-12-24 04:57:28
【问题描述】：
我在 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;
}
【问题讨论】：
-
您应该提供一个minimal reproducible example 我的猜测是由于取消引用
nodes内的nodes指针graph，但您没有提供您如何构建的描述graph。与其试图解决这一点，我建议您提供一个minimal reproducible example 这样的问题，请参阅第1 项here。 -
@RobertCrovella 你是完全正确的。在这里。
-
我没有看到在您的代码中使用统一（即托管）内存的任何证据。也许您对CUDA unified memory 是什么感到困惑。您还可以在
cuda标签上研究 CUDA UM 示例代码以及有关它的许多问题。撇开这一点不谈，您的代码试图取消引用设备代码中的主机指针（由malloc提供）。这是基本的 CUDA 禁忌，也是您观察到的错误的近端原因。 -
@RobertCrovella 啊，我可能误解了这个例子（这里）[devblogs.nvidia.com/unified-memory-in-cuda-6/].我见过的大多数示例将内存创建为 UM（如链接文档
cudaMallocManaged中所示），然后用东西填充它。我还没有找到一个示例，其中我已经有一个如图所示的主机结构（由另一个进程生成，在 C 上没有对其进行控制），我想将它传递给 CUDA。我不擅长谷歌或者我不能这样做？ -
in the docs (K.1.6)...阅读整个 K.1.6 部分。到最后一句话/行。将主机分配器（例如malloc）直接用作“托管”分配器取决于 ATS 支持。 K.1.6 节第一段的最后一句指出“应用程序可以通过检查新的pageableMemoryAccessUsesHostPageTables属性来查询设备是否支持通过 ATS 一致地访问可分页内存。”你检查过那个属性吗？ （提示：您的系统不支持此功能）。
标签： cuda unified-memory