host多线程分别处理不同gpu上的内存问题

  • 3 replies
  • 138 views
host多线程分别处理不同gpu上的内存问题
« 于: 十二月 13, 2019, 12:52:22 pm »
是这样的,我现在有三个相机,每个相机的数据接收buffer是分别在三个gpu上申请的,并且在三个gpu上分别创建了一个stream;
我开了三个线程,每个线程进行接收一个相机的数据,并调用cuda核函数去处理图像,cuda核函数是调用了当前gpu上的stream进行异步处理;

我现在是每个线程每次运行时,都会调用cudaSetDevice()来切换到当前线程对应的gpu,以便我能调用该gpu对应的stream,并能访问该gpu上的内存。但是由于每个线程的核函数都是异步运行,最后就必须调用流同步函数来同步该流。由于流同步函数是阻塞的,那么在阻塞时,cpu会切换到其他线程继续运行吗?
另外就是有没有其他线程绑定方法来解决我这个每次线程运行都要调用cudaSetDevice()切换对应gpu?

第三个问题是,我在gpu上申请的内存是用cudaMallocMange来申请的,那么比如当前某个线程调用了cudaSetDevice(0),那么另外一个线程可以在不调用cudaSetDevice(1)就能访问gpu 1上我申请的内存吗?
多谢指导!

Re: host多线程分别处理不同gpu上的内存问题
« 回复 #1 于: 十二月 13, 2019, 02:53:54 pm »
是这样的,我现在有三个相机,每个相机的数据接收buffer是分别在三个gpu上申请的,并且在三个gpu上分别创建了一个stream;
我开了三个线程,每个线程进行接收一个相机的数据,并调用cuda核函数去处理图像,cuda核函数是调用了当前gpu上的stream进行异步处理;

我现在是每个线程每次运行时,都会调用cudaSetDevice()来切换到当前线程对应的gpu,以便我能调用该gpu对应的stream,并能访问该gpu上的内存。但是由于每个线程的核函数都是异步运行,最后就必须调用流同步函数来同步该流。由于流同步函数是阻塞的,那么在阻塞时,cpu会切换到其他线程继续运行吗?
另外就是有没有其他线程绑定方法来解决我这个每次线程运行都要调用cudaSetDevice()切换对应gpu?

第三个问题是,我在gpu上申请的内存是用cudaMallocMange来申请的,那么比如当前某个线程调用了cudaSetDevice(0),那么另外一个线程可以在不调用cudaSetDevice(1)就能访问gpu 1上我申请的内存吗?
多谢指导!

你的中文基本上没看懂。我大致猜测了一下你的意思:

(1)当一个CPU线程被阻塞后,其他的线程有机会继续使用CPU的。这个没什么问题。这个是基本的OS的调度能力,和CUDA其实并没有关系。

你可以想象一下你以前用单核的机器,Windows XP的时代,你可能同时运行一堆应用程序,存在大量线程,难道某个应用程序的进程里面的1个线程被阻塞了,整个系统上的所有线程都被卡住么?显然不会的,他们依然会得到CPU调度执行机会的。我建议你了解一下基本的操作系统相关概念。

(2)使用CUDA Runtime API的时候,必须通过cudaSetDevice()来设定本Host线程所使用的当前CPU的。不能越过。这个函数的切换代价非常低的,你过虑了。

再说了,你完全可以保持一个Host线程设定device后,长期执行。这样连极地的切换代价也木有了(除了第一次切换)。

(3)我没有看懂你的中文。Managed Memory是自动管理的,你是想强制要求某段managed memory必须驻留在某个GPU的显存里?然后却又强制性的从其他GPU访问该GPU上的驻留的managed memory?

如果是这样想法,那就很奇葩了。。你越过了自动的数据临近性存放位置管理。如果你真的确定你这种想法是需要的,你可以考虑设定一下该段manged memory的preferred location,然后手工执行一次async的prefetch,尽量将数据迁移到某GPU中,然后同时设定禁用read mostly(自动只读副本),同时在确保另外的一个GPU能访问某GPU中的直接显存内容,这样还有可能。

但即使如此,在某些情况下,managed memory的自动管理机制依然可能会将数据迁移到其他临近访问的GPU中的。

大部分试图自己做的人,可能往往都没有自动做的好。如果不是非常必须,我不建议你这样。

Re: host多线程分别处理不同gpu上的内存问题
« 回复 #2 于: 十二月 13, 2019, 03:46:16 pm »
您好,谢谢您的详细回答。
1,关于线程调度的问题,我的担忧是:我调用cudaStreamSync()时,这个函数是cuda函数,我主要是不了解它,不知道它内部阻塞时是释放了cpu使用权,还是像一个死循环一样一直阻塞?如果它内部释放了cpu,倒还好,如果不是,我担心其他线程是在时间片到了之后才有机会执行自己,如果时间片是几ms,那么对于我们的实时性也是浪费的。所以跟您确认一下它内部的机制;
2,我并不是一定是在用的gpu上申请内存,我那样做是因为我不确定我在不调用cudaSetDevice()时,直接用cudaMallocMange()函数申请的内存,是不是默认所有的gpu都可以对它访问?因为我对cudaMallocMange()函数申请的内存了解不透彻,只知道当cpu和gpu互相拷贝数据时,它可以不用手动调用cudaMemcpy()函数。
3,为了降低某个gpu的负载,我想让每个相机的图像处理工作分别放在3个gpu上,所以在程序开始初始化时,我先调用cudaSetDevice(),然后创建stream,这相当于在3个gpu上分别创建了3个stream。我用了3个线程,每个线程里写了一个核函数进行图像处理,每个线程分别用了一个stream,那么是不是每次线程执行时,都要先调用cudaSetDevice(),然后才能用对应的stream?
上面三个问题再跟您确认一下,谢谢!

Re: host多线程分别处理不同gpu上的内存问题
« 回复 #3 于: 十二月 13, 2019, 06:16:18 pm »
您好,谢谢您的详细回答。
1,关于线程调度的问题,我的担忧是:我调用cudaStreamSync()时,这个函数是cuda函数,我主要是不了解它,不知道它内部阻塞时是释放了cpu使用权,还是像一个死循环一样一直阻塞?如果它内部释放了cpu,倒还好,如果不是,我担心其他线程是在时间片到了之后才有机会执行自己,如果时间片是几ms,那么对于我们的实时性也是浪费的。所以跟您确认一下它内部的机制;
2,我并不是一定是在用的gpu上申请内存,我那样做是因为我不确定我在不调用cudaSetDevice()时,直接用cudaMallocMange()函数申请的内存,是不是默认所有的gpu都可以对它访问?因为我对cudaMallocMange()函数申请的内存了解不透彻,只知道当cpu和gpu互相拷贝数据时,它可以不用手动调用cudaMemcpy()函数。
3,为了降低某个gpu的负载,我想让每个相机的图像处理工作分别放在3个gpu上,所以在程序开始初始化时,我先调用cudaSetDevice(),然后创建stream,这相当于在3个gpu上分别创建了3个stream。我用了3个线程,每个线程里写了一个核函数进行图像处理,每个线程分别用了一个stream,那么是不是每次线程执行时,都要先调用cudaSetDevice(),然后才能用对应的stream?
上面三个问题再跟您确认一下,谢谢!

(1)阻塞在同步对象上(例如某些平台上可能是个内核态的mutex),和持续spin都是有可能的。这个是自动的。如果你需要手工控制,可以使用手册的cudaSetDeviceFlags()函数,指定spin或者blocking sync的标志。

详情可以参考本论坛Sisiy妹子的《阅读CUDA手册100天》。

(2)新人建议总是使用明确的显存分配,明确的内存(包括page-locked的)分配,和明确的cudaMemcpy*()。

如果想知道unified memory的详情,请参考本论坛的sisiy妹子的《阅读CUDA手册100天》

(3)没错,只有先确定了device后,后续的一切操作才能有效(包括你想要的启动kernel,在某个stream中;也包括存储器分配等等)