CUDA 统一内存可以用作固定内存（统一虚拟内存）吗？

Question

据我所知，我们可以在内核内存中分配一个固定内存区域。 （来自KGPU）

然后，在 Pinned 内存中分配 linux 内核数据并传输到 GPU。

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

今天，一个案例就是一棵树。

我尝试将其从固定内存传递到 GPU。

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

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

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

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

Answer 1

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

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

还是说统一内存只能使用cudaMallocManaged？

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

The programming guide 有更多信息。

针对下面的 cmets，这是一个使用固定内存创建单链表并在设备代码中遍历该列表的简单示例：

$ 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。