多GPU多kernel变量共享

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

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

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

第二个问题,你对一个全局变量依次再不同设备上malloc,那么出现以下问题:第二个分配后,第一个分配的空间就再找不到了…最后结果就是前面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。该如何实现呢?

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计算。

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

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

不明白您的意思。

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

具体实现你安排好即可。

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

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

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

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

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

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

祝您一切顺利~

  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数目个的数组里,您可以自行根据是否方便决定。

从上文我的回复你实际上也看出了,“你的试图多个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之类的方式,而是直接进行拷贝。你的本能感觉“大数据拷贝不好”而“尽量共享内存/显存的访问”,实际上可能和你的想想是相反的。你直觉认为较慢的方式可能较快。我建议你尝试下,说不定性能实际上反而是最好的。

楼主加油。

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

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

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

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

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

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

感谢您的莅临,
祝您生活愉快!

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

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

祝您编码顺利!