多GPU多kernel变量共享

加速计算 CUDA
1

现在单节点有多块GPU卡，希望多线程利用多块卡，而且每块卡上会多次调用同一kernel。遇到两个疑问：

  1. 因为拷贝的数据很大，所以希望多个kernel能够共享一个变量。现在通过在host上cudamalloc，然后指针传到device，这样变量的lifetime是kernel还是application？看到帖子说变量加上__device__，如果使用能动态分配空间吗？该如何拷贝数据
  2. 因为变量很多，为了方便，使用了全局变量，四个线程都用到此变量。这时候通过cudasetdevice先选择GPU，再cudaMalloc会存在一个全局变量多次分配的问题吗？是否得使用数组？
2

多个kernel共享一个变量？是数组还是一个变量？如何共享？共享读取？写入？这个还请楼主先说明一下，最好举例。

第二个问题，你对一个全局变量依次再不同设备上malloc，那么出现以下问题:第二个分配后，第一个分配的空间就再找不到了…最后结果就是前面3个分配的空间都泄露了，你无法再访问到。

3

----------------a.c---------------------
variable *XX;
func1(){
for(…){
for(i = 0; i < 4; i++)
cutStartThread(func2,args)
}
}
---------------b.cu-------------------
func2(args){
cudaSetDevice(i);
xx = cudaMalloc();
mykernel<<<…>>>(xx);
}

一个简单的示例，XX是一个很大的数组，在myKernel中是只读的。希望第一次for循环将xx拷贝到GPU，之后for循环再启动mykernel不用再拷贝xx。该如何实现呢?

4

LZ您好，这里仅回复您3#的内容:
“一个简单的示例，XX是一个很大的数组，在myKernel中是只读的。希望第一次for循环将xx拷贝到GPU，之后for循环再启动mykernel不用再拷贝xx。该如何实现呢?”

1：因为XX在kernel中是只读的，那么只要XX的容量没有超过单个显卡的显存容量，那么您给每个显卡copy一分即可，无需考虑其他同步，是否被修改等的问题。

2：这样的话，您只需对每个GPU，先cudaMalloc，然后cudaMemcpy，然后各自使用即可。注意不要用同一个指针变量申请，而是每个卡配备自己的指针（也就是每个卡维护自己的副本）。

3：cudaMalloc到的数据，在cudaFree之前都是有效的，和启动过多少次kernel无关，所以您尽可在第一次copy完成之后，反复启动kernel计算。

大致如上，祝您编码顺利~

5

十分感谢版主！那就是现在将XX改为数组，每一个元素维护一个。还有一个小疑问，如果需要将XX绑定到texture，那么texture是不需要定义为数组的？

6

不明白您的意思。

“那就是现在将XX改为数组”，XX本来不就是数组么？您在3#里面就是这么说的。
“每一个元素维护一个”，每个什么元素维护一个什么？每个线程维护一个数组？每个线程维护一个数组元素？

具体实现你安排好即可。

如果您需要使用 texture，请参考programming guide上的详细用法。
就性能而言，在fermi上，如果您不是为了使用texture的插值特性等硬件特性，那么不推荐使用texture，直接使用一般的数组即可。
如果在SM 3.0的kepler上，可以使用texture来优化访存。

但这一切都需要建立在您的程序确实需要优化这方面的基础上，如果瓶颈在其他地方，那么优化这些也没有意义。

所以，建议您先使用一般的数组，保证合并访问，用NVVP评估程序性能和瓶颈，然后再考虑其他。

7

sorry，没说清楚。是指针数组，比如XX[4]，第k个线程维护XX[k]，XX[k]是一个大数组。现在是四块tesla c2075的卡，但是实际绑定到texture的数组比较小且频繁访问。不会使用到插值，只是简单的访问。单纯考虑这个的话性能提升不了多少？在Linux上开发的，目前只会cuda-gdb和command line profiler，调试是比较麻烦

8

LZ您好，如果您总的数据可以分割为4部分，并且每个显卡只负责处理其中一部分，显卡间无交互或者交互很少的话，多卡效果应该还是不错的。

关于那些比较小且频繁访问的数组，您可以分别用常规数组和texture实现一下，看看哪个效果好。
command line profiler估计可以看访存的情况，但是我不用这个，并不知道详细的情况。

祝您一切顺利~

9
  1. 因为拷贝的数据很大，所以希望多个kernel能够共享一个变量。现在通过在host上cudamalloc，然后指针传到device，这样变量的lifetime是kernel还是application？看到帖子说变量加上__device__，如果使用能动态分配空间吗？该如何拷贝数据

—这个分配空间的生存周期是cuda context, 也就是在您在cudaSetDevice(n)保持n为一个值的时候，cudaFree/cudaDeviceReset将对其释放有效。以及，无论n为哪个值，直接您的host process结束，它们均会被释放。

  1. 因为变量很多，为了方便，使用了全局变量，四个线程都用到此变量。这时候通过cudasetdevice先选择GPU，再cudaMalloc会存在一个全局变量多次分配的问题吗？是否得使用数组？
    –正因为上一段所说的，cudaMalloc出来的设备内存只能在一个cuda context里有效，而该context实际上只在1个gpu上，所以您说的没错，的确需要多次cudaSetDevice, 然后多次cudaMalloc。
    –是否将这些多次分配的结果放置在一个元素数目为您使用的gpu数目个的数组里，您可以自行根据是否方便决定。
10

从上文我的回复你实际上也看出了，“你的试图多个gpu上的kernel共享数据，从而节省空间/节省copy的方式”，实际上这样不能满足你的目的。

那么你怎么办？那么就不能多次cudaSetDevice + 多次cudaMalloc了。
我给你2个建议，符合你的最开始的假设和愿望：

（1）使用在多个context间portable的pinned host memory, 并依次映射到设备地址空间，这样对巨大的数据，只需要分配一次。但device访问host memory效率较低。
（这种方式您可以参考cudaSetDeviceFlags()和cudaHostAlloc()和cudaHostGetDevicePointer())

（2）你可以尝试多个GPU间共享数据，通过pci-e总线进行交换，但不经过host memory.这种方式，
需要你的卡支持p2p access。你可以参考手册中的例子，以及参考cudaDeviceEnablePeerAccess，以及，必须是64位程序才有效。

以及，如果只是为了规避“拷贝大数据的时间”，而不是因为“显存放不下”。那么我不建议你使用peer access之类的方式，而是直接进行拷贝。你的本能感觉“大数据拷贝不好”而“尽量共享内存/显存的访问”，实际上可能和你的想想是相反的。你直觉认为较慢的方式可能较快。我建议你尝试下，说不定性能实际上反而是最好的。

楼主加油。

11

以上2楼是我对您的1#的回复。

对于您的2楼的问题，您可以直接考虑第一次启动前单独cudaMemcpy一次，然后多次使用，最后释放即可。

因为2楼这个问题已经被ICE详细解答过（实际1#也是), 我就不赘述了。此楼表示我也重视过您的2#的问题。

12

谢谢版主，现在比较明晰了，就先在一个全局数组维护每个device的指针，然后每个device上cudaMalloc一份，由于内存较大，估计不是问题。再次感谢！！很有帮助

13

确实满足您说的情况，四个分块没有数据相关性。谢谢ice，回复的十分及时，我先调试试一试，遇到问题再请教！！

14

您客气了。以及ICE提前已经回复的差不多了。我只是稍微补充了下。

感谢您的莅临，
祝您生活愉快！

15

您客气了，横扫斑竹回复的十分详尽，我只说了3#问题的一小部分。

欢迎您继续回来反馈讨论此问题。

祝您编码顺利！