同步多个Cuda流
对于我目前正在开发的应用程序，我希望有一个长内核(即一个相对于其他内核需要很长时间才能完成的内核)与多个同时运行的较短内核序列并发执行。然而，更复杂的是，四个较短的内核都需要在完成后进行同步，以便执行另一个短内核，该内核收集和处理其他短内核输出的数据。
以下是我所想到的一个示意图，编号的绿色条形代表不同的内核：
为了实现这一点，我编写了如下代码：
// definitions of kernels 1-6
class Calc
{
Calc()
{
// ...
cudaStream_t stream[5];
for(int i=0; i<5; i++) cudaStreamCreate(&stream[i]);
// ...
}
~Calc()
{
// ...
for(int i=0; i<5; i++) cudaStreamDestroy(stream[i]);
// ...
}
void compute()
{
kernel1<<<32, 32, 0, stream[0]>>>(...);
for(int i=0; i<20; i++) // this 20 is a constant throughout the program
{
kernel2<<<1, 32, 0, stream[1]>>>(...);
kernel3<<<1, 32, 0, stream[2]>>>(...);
kernel4<<<1, 32, 0, stream[3]>>>(...);
kernel5<<<1, 32, 0, stream[4]>>>(...);
// ?? synchronisation ??
kernel6<<<1, 32, 0, stream[1]>>>(...);
}
}
}
int main()
{
// preparation
Calc C;
// run compute-heavy function as many times as needed
for(int i=0; i<100; i++)
{
C.compute();
}
// ...
return 0;
}注意:块、线程和共享内存的数量只是任意数字。
现在，我将如何正确地同步每次迭代2-5内核？首先，我不知道哪个内核需要最长的时间才能完成，因为这可能取决于用户的输入。此外，我尝试过使用cudaDeviceSynchronize()和cudaStreamSynchronize()，但这使总执行时间增加了两倍多。
库达事件也许是要走的路吗？如果是这样的话，我应该如何应用它们呢？如果不这样做，怎样才是正确的方法？
非常感谢。
回答 1
Stack Overflow用户
发布于 2016-07-18 21:26:51
有两项意见需要首先提出。
- 启动小内核(一个块)通常不是从GPU中获得良好性能的方法。同样，每个块(32)具有少量线程的内核通常会施加占用限制，这将阻止GPU的全面性能。启动多个并发内核并不会减轻第二个考虑。我不会再花更多的时间在这里，因为你已经说过这些数字是任意的(但是请看下面的下一个评论)。
- 见证实际的内核并发是很困难的。我们需要执行时间相对较长但对GPU资源的需求相对较低的内核。
<<<32,32>>>的内核可能会填充正在运行的GPU的，从而防止并发内核中的块运行。
您的问题似乎归结为“如何防止kernel6在kernel2-5完成之前启动。
可以使用事件来实现这一目的。基本上，您将在内核2-5启动之后，将记录事件放入每个流中，并在启动kernel6之前为4个事件中的每个事件设置一个cudaStreamWaitEvent调用。
就像这样：
kernel2<<<1, 32, 0, stream[1]>>>(...);
cudaEventRecord(event1, stream[1]);
kernel3<<<1, 32, 0, stream[2]>>>(...);
cudaEventRecord(event2, stream[2]);
kernel4<<<1, 32, 0, stream[3]>>>(...);
cudaEventRecord(event3, stream[3]);
kernel5<<<1, 32, 0, stream[4]>>>(...);
cudaEventRecord(event4, stream[4]);
// ?? synchronisation ??
cudaStreamWaitEvent(stream[1], event1);
cudaStreamWaitEvent(stream[1], event2);
cudaStreamWaitEvent(stream[1], event3);
cudaStreamWaitEvent(stream[1], event4);
kernel6<<<1, 32, 0, stream[1]>>>(...);请注意，上述所有调用都是异步的。它们中的任何一个都不应超过几微秒来处理，而且它们都不会阻止CPU线程继续进行，这与您使用cudaDeviceSynchronize()或cudaStreamSynchronize()不同，后者通常会阻塞CPU线程。
因此，您可能希望在循环中执行上述序列(例如cudaStreamSynchronize(stream[1]);)之后进行某种同步，或者所有这些的异步性质都会变得毛茸茸的(另外，根据您的原理图，您可能不希望迭代i+1的内核2-5在迭代I的kernel6完成之前开始？)请注意，我忽略了事件创建和其他样板，我假设您可以找出或引用使用事件的任何示例代码，或者参考文档。
即使您实现了所有这些基础设施，您见证(或不)实际内核并发的能力也将取决于内核本身，而不是我在这个答案中建议的任何内容。所以，如果你回来说“我做到了，但我的内核不是同时运行的”，这实际上是一个与你提出的不同的问题，在这里，我建议你首先参考我上面的第2条评论。
https://stackoverflow.com/questions/38445282
复制相似问题
腾讯云开发者
