cuda unified memory简介 和 cuda深拷贝问题

个人博客中有本文原文https://whatghost.github.io/2020/03/12/cuda-unified-memory/
如果喜欢可以访问个人博客WhatGhost.github.io，会有更多分享内容

cuda unified memory简介 和 cuda深拷贝问题

unified memory

unified memory 是cuda6.0以后的版本增加的一个编程模型
统一内存创建了一个受管内存池，该池在CPU和GPU之间共享，弥合了CPU与GPU的鸿沟。使得CPU端和GPU端都能通过同一个指针访问数据，从而避免了程序员编写代码时显示的执行GPU和CPU端的内存拷贝命令

一个使用unified memory的例子如下

void sortfile(FILE *fp, int N){
    char *data; 
    cudaMallocManaged(data, N);

    fread(data, 1, N, fp);

    qsort<<<...>>>(data, N, 1, compare);
    cudaDeviceSynchronize()
    usedata(data);
    free(data);
}

观察上面的例子我们看到，以前需要的 cudaMemcpy 不再需要了
如果不用Unified Memory，我们如果需要在GPU上处理数据，需要显示的调用cudaMemcpy函数到GPU端，需要在CPU端处理的时候还需要在cudaMemcpy回来,会带来很大的麻烦。而unified memory则可以解决这个问题（拷贝依然会执行，但是是系统帮程序员执行，不需要显示调用cudamemcpy）

使用unified memry 需要使用 cudaMallocManaged 函数来进行分配。
这样带来的好处主要是可以减少代码量，使得程序员不用费尽心思去进行CPU和GPU端数据的拷贝，尤其是对于一些嵌套结构体，链表之类的拷贝起来比较麻烦的问题。

深拷贝问题

使用unified memory可以很好地将解决这个问题，使得我们不用再考虑深拷贝的问题，而专心的解决程序算法的编写，这个是显而易见的，就不在啰嗦

但是很多时候尤其是在优化现有项目的时候，我们想向GPU端传一个已经初始化好的结构体，如何进行深拷贝
有以下的解决方案

加入我们有以下结构体

struct dataElem {
  int prop1;
  int prop2;
  char *name;
}

里面包含一个指针，这样单纯的cudaMemcpy(dst,src,sizeof(dataElem),cudaMemcpyHostToDevice)是不行的，因为这样的只是浅拷贝，只会拷贝指针的值就是地址，但是GPU端是没法使用CPU端的指针的，所以我们需要一下的深拷贝方法

void launch(dataElem *elem) {
  dataElem *d_elem;
  char *d_name;

  int namelen = strlen(elem->name) + 1;

  //为elem和name分别分配空间
  cudaMalloc(&d_elem, sizeof(dataElem));
  cudaMalloc(&d_name, namelen);

  // 把elem和里面name的值分别拷贝到GPU端
  cudaMemcpy(d_elem, elem, sizeof(dataElem), cudaMemcpyHostToDevice);
  cudaMemcpy(d_name, elem->name, namelen, cudaMemcpyHostToDevice);
  //将d_elem里name的指针重新赋值为GPU端的指针
  cudaMemcpy(&(d_elem->name), &d_name, sizeof(char*), cudaMemcpyHostToDevice);

  // 运行核函数
  Kernel<<< ... >>>(d_elem);
}

但是如果是嵌套非常深的结构体（最近被这种问题搞得很难受），上述的方法显然麻烦得很，那就需要把原来的结构体展开成一个没有多层嵌套的结构体，就相当于把这个很深结构拉开展平，在进行传参。

最后，写cuda时不要定义很深的结构体，如果有还是要用unified memory