CUDA统一内存可以用作固定内存(统一虚拟内存)吗？

Question

据我所知，我们可以在内核内存中分配一个 Pinned 内存区域。 (来自KGPU)

然后，在Pinned memory中分配linux kernel数据，并传输给GPU。

但问题是linux内核数据应该排列成数组。

今天的案例就是一棵树。

我试过将它从固定内存传递到 GPU。

但是当一个节点访问下一个节点时，内存访问错误。

我想知道是否可以将统一内存分配为内核内存中的固定内存区域？

因此树可以在统一内存区域中构建并由 GPU 使用，而无需其他运行时 API，如 cudaMallocMaganed。

还是统一内存只能使用cudaMallocMaganed？

Answer 1

But when a node access next node, memory access error occurred.

这只是意味着您的代码中存在错误。

Or is that Unified memory must only use cudaMallocManaged?

目前，访问统一内存功能的唯一方法是使用托管分配器。对于动态分配，即 cudaMallocManaged()。对于静态分配，它是通过 __managed__ 关键字。

The programming guide有额外的信息。

为了回应下面的评论，这里有一个简单的工作示例，它使用固定内存创建单链表，并在设备代码中遍历该列表:

$ cat t1115.cu
#include <stdio.h>
#define NUM_ELE 5

struct ListElem{

   int id;
   bool last;
   ListElem *next;
};

__global__ void test_kernel(ListElem *list){

  int count = 0;
  while (!(list->last)){
    printf("List element %d has id %d\n", count++, list->id);
    list = list->next;}
  printf("List element %d is the last item in the list\n", count);
}

int main(){
  ListElem *h_list, *my_list;
  cudaHostAlloc(&h_list, sizeof(ListElem), cudaHostAllocDefault);
  my_list = h_list;
  for (int i = 0; i < NUM_ELE-1; i++){
    my_list->id = i+101;
    my_list->last = false;
    cudaHostAlloc(&(my_list->next), sizeof(ListElem), cudaHostAllocDefault);
    my_list = my_list->next;}
  my_list->last = true;
  test_kernel<<<1,1>>>(h_list);
  cudaDeviceSynchronize();
}

$ nvcc -o t1115 t1115.cu
$ cuda-memcheck ./t1115
========= CUDA-MEMCHECK
List element 0 has id 101
List element 1 has id 102
List element 2 has id 103
List element 3 has id 104
List element 4 is the last item in the list
========= ERROR SUMMARY: 0 errors
$

请注意，为了简洁起见，我省略了 proper CUDA error checking在此示例中(尽管使用 cuda-memcheck 运行代码表明没有 CUDA 运行时错误)，但我建议您在使用 CUDA 代码时遇到问题时使用它。另请注意，此示例假定 proper UVA environment .