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

Question

如我所知，我们可以在内核内存中分配固定内存区域。(来自KGPU)

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

但是问题是linux内核数据应该被安排为数组。

今天，一个案件是一棵树。

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

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

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

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

还是统一内存必须只使用cudaMallocMaganed

Answer 1

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

这只意味着您的代码中有一个bug。

还是统一内存必须只使用cudaMallocManaged？

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

程序设计指南有更多的信息。

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

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

请注意，为了简洁的表示，我在本例中不再使用正确的CUDA错误检查 (尽管使用cuda-memcheck运行代码说明没有CUDA运行时错误)，但我建议您在任何时候遇到CUDA代码问题时都可以使用它。还请注意，此示例假定为适当的UVA环境。