草庐IT

c++ - 同步多个 Cuda 流

coder 2024-02-23 原文

对于我目前正在开发的应用程序,我希望有一个长内核(即,相对于其他内核需要很长时间才能完成的内核)与一系列同时运行的多个较短内核同时执行.然而,使这变得更复杂的是,四个较短的内核在完成后都需要同步,以便执行另一个短内核收集和处理其他短内核输出的数据。

下面是我的想法的示意图,带编号的绿色条代表不同的内核:

为了实现这一点,我编写了如下所示的代码:

// 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;
}

注意: block 、线程和共享内存的数量只是任意数字。

现在,我将如何在每次迭代中正确同步内核 2-5?首先,我不知道哪个内核需要最长的时间才能完成,因为这可能取决于用户输入。此外,我尝试过使用 cudaDeviceSynchronize()cudaStreamSynchronize(),但总执行时间增加了两倍多。

Cuda 事件是否可行?如果是这样,我应该如何应用它们?如果不是,执行此操作的正确方法是什么?

非常感谢。

最佳答案

有两条评论需要先说

  1. 启动小内核(一个 block )通常不是获得 GPU 良好性能的方法。同样,每个 block (32) 具有少量线程的内核通常会施加占用限制,这将阻止 GPU 的全部性能。启动多个并发内核并不能减轻第二个考虑因素。我不会在这里花更多时间,因为你已经说过数字是任意的(但请参阅下面的下一条评论)。

  2. 见证实际的内核并发是很困难的。我们需要执行时间相对较长但对 GPU 资源要求相对较低的内核。 <<<32,32>>> 的内核可能会填充您正在运行的 GPU,从而阻止并发内核中的 block 运行的任何能力。

您的问题似乎可以归结为“我如何防止 kernel6kernel2-5 完成之前开始。

为此可以使用事件。基本上,你会 record an event到每个流中,在 kernel2-5 启动后,你会放一个 cudaStreamWaitEvent kernel6 启动之前调用,为 4 个事件中的每一个调用.

像这样:

        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 的 kernel2-5 在迭代 i 的 kernel6 完成之前开始?)请注意,我为此省略了事件创建和其他样板,我'我假设您可以弄清楚或引用任何使用事件的示例代码,或引用文档。

即使您实现了所有这些基础设施,您见证(或不见证)实际内核并发的能力也将由您的内核本身决定,而不是我在此答案中提出的任何建议。因此,如果您回过头来说“我这样做了,但我的内核没有同时运行”,这实际上与您在这里提出的问题不同,我建议您首先引用我上面的评论 #2。

关于c++ - 同步多个 Cuda 流,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/38445282/

有关c++ - 同步多个 Cuda 流的更多相关文章

  1. ruby-on-rails - Rails 3 中的多个路由文件 - 2

    Rails2.3可以选择随时使用RouteSet#add_configuration_file添加更多路由。是否可以在Rails3项目中做同样的事情? 最佳答案 在config/application.rb中:config.paths.config.routes在Rails3.2(也可能是Rails3.1)中,使用:config.paths["config/routes"] 关于ruby-on-rails-Rails3中的多个路由文件,我们在StackOverflow上找到一个类似的问题

  2. ruby-on-rails - 在 Ruby 中循环遍历多个数组 - 2

    我有多个ActiveRecord子类Item的实例数组,我需要根据最早的事件循环打印。在这种情况下,我需要打印付款和维护日期,如下所示:ItemAmaintenancerequiredin5daysItemBpaymentrequiredin6daysItemApaymentrequiredin7daysItemBmaintenancerequiredin8days我目前有两个查询,用于查找maintenance和payment项目(非排他性查询),并输出如下内容:paymentrequiredin...maintenancerequiredin...有什么方法可以改善上述(丑陋的)代

  3. ruby-on-rails - Rails - 一个 View 中的多个模型 - 2

    我需要从一个View访问多个模型。以前,我的links_controller仅用于提供以不同方式排序的链接资源。现在我想包括一个部分(我假设)显示按分数排序的顶级用户(@users=User.all.sort_by(&:score))我知道我可以将此代码插入每个链接操作并从View访问它,但这似乎不是“ruby方式”,我将需要在不久的将来访问更多模型。这可能会变得很脏,是否有针对这种情况的任何技术?注意事项:我认为我的应用程序正朝着单一格式和动态页面内容的方向发展,本质上是一个典型的网络应用程序。我知道before_filter但考虑到我希望应用程序进入的方向,这似乎很麻烦。最终从任何

  4. ruby-on-rails - 如何优雅地重启 thin + nginx? - 2

    我的瘦服务器配置了nginx,我的ROR应用程序正在它们上运行。在我发布代码更新时运行thinrestart会给我的应用程序带来一些停机时间。我试图弄清楚如何优雅地重启正在运行的Thin实例,但找不到好的解决方案。有没有人能做到这一点? 最佳答案 #Restartjustthethinserverdescribedbythatconfigsudothin-C/etc/thin/mysite.ymlrestartNginx将继续运行并代理请求。如果您将Nginx设置为使用多个上游服务器,例如server{listen80;server

  5. ruby - 多个属性的 update_column 方法 - 2

    我有一个具有一些属性的模型:attr1、attr2和attr3。我需要在不执行回调和验证的情况下更新此属性。我找到了update_column方法,但我想同时更新三个属性。我需要这样的东西:update_columns({attr1:val1,attr2:val2,attr3:val3})代替update_column(attr1,val1)update_column(attr2,val2)update_column(attr3,val3) 最佳答案 您可以使用update_columns(attr1:val1,attr2:val2

  6. ruby-on-rails - 在 ruby​​ .gemspec 文件中,如何指定依赖项的多个版本? - 2

    我正在尝试修改当前依赖于定义为activeresource的gem:s.add_dependency"activeresource","~>3.0"为了让gem与Rails4一起工作,我需要扩展依赖关系以与activeresource的版本3或4一起工作。我不想简单地添加以下内容,因为它可能会在以后引起问题:s.add_dependency"activeresource",">=3.0"有没有办法指定可接受版本的列表?~>3.0还是~>4.0? 最佳答案 根据thedocumentation,如果你想要3到4之间的所有版本,你可以这

  7. ruby - 使用多个数组创建计数 - 2

    我正在尝试按0-9和a-z的顺序创建数字和字母列表。我有一组值value_array=['0','1','2','3','4','5','6','7','8','9','a','b','光盘','e','f','g','h','i','j','k','l','m','n','o','p','q','r','s','','u','v','w','x','y','z']和一个组合列表的数组,按顺序,这些数字可以产生x个字符,比方说三个list_array=[]和一个当前字母和数字组合的数组(在将它插入列表数组之前我会把它变成一个字符串,]current_combo['0','0','0']

  8. ruby-on-rails - before_filter 运行多个方法 - 2

    是否有可能:before_filter:authenticate_user!||:authenticate_admin! 最佳答案 before_filter:do_authenticationdefdo_authenticationauthenticate_user!||authenticate_admin!end 关于ruby-on-rails-before_filter运行多个方法,我们在StackOverflow上找到一个类似的问题: https://

  9. ruby-on-rails - Rails 3.1 中具有相同形式的多个模型? - 2

    我正在使用Rails3.1并在一个论坛上工作。我有一个名为Topic的模型,每个模型都有许多Post。当用户创建新主题时,他们也应该创建第一个Post。但是,我不确定如何以相同的形式执行此操作。这是我的代码:classTopic:destroyaccepts_nested_attributes_for:postsvalidates_presence_of:titleendclassPost...但这似乎不起作用。有什么想法吗?谢谢! 最佳答案 @Pablo的回答似乎有你需要的一切。但更具体地说...首先改变你View中的这一行对此#

  10. ruby-on-rails - 使用 ruby​​ 将多个实例变量转换为散列的更好方法? - 2

    我收到格式为的回复#我需要将其转换为哈希值(针对活跃商家)。目前我正在遍历变量并执行此操作:response.instance_variables.eachdo|r|my_hash.merge!(r.to_s.delete("@").intern=>response.instance_eval(r.to_s.delete("@")))end这有效,它将生成{:first="charlie",:last=>"kelly"},但它似乎有点hacky和不稳定。有更好的方法吗?编辑:我刚刚意识到我可以使用instance_variable_get作为该等式的第二部分,但这仍然是主要问题。

随机推荐