jjzjj

c++ - 扭曲如何与原子操作一起工作？

coder 2024-02-12 原文

warp 中的线程在物理上并行运行，所以如果其中一个(称为线程 X)启动原子操作，其他线程会做什么？等待？这是否意味着，当线程 X 被推送到原子队列时，所有线程都将等待，获取访问权限(互斥锁)并使用内存做一些事情，内存受到该互斥锁的保护，然后才是真正的互斥锁？

有没有办法让其他线程进行某些工作，比如读取一些内存，这样原子操作会隐藏它的延迟？我的意思是，有 15 个空闲线程......我猜不太好。 Atomic 真的很慢，是吗？我怎样才能加速它？有什么模式可以使用它吗？

共享内存的原子操作是否锁定银行或整个内存？
例如(没有互斥量)，有 __shared__ float smem[256];

  • 线程 1 运行 atomicAdd(smem, 1);
  • 线程 2 运行 atomicAdd(smem + 1, 1);

    • 这些线程适用于不同的银行，但通常共享内存。他们是并行运行还是要排队？这个例子有什么区别，如果Thread1和Thread2来自分开的warp还是一般的warp？

    最佳答案

    我数了大约 10 个问题。这让人很难回答。建议您每个问题问一个问题。

    一般来说，warp 中的所有线程都在执行相同的指令流。所以有两种情况我们可以考虑:

  • 没有条件(例如 if...then...else)在这种情况下，所有线程都在执行相同的指令，这恰好是一条原子指令。然后所有 32 个线程都将执行一个原子，尽管不一定在同一位置。所有这些原子都将由 SM 处理，并且在某种程度上会序列化(如果它们更新相同的位置，它们将完全序列化)。
  • 带条件 例如，假设我们有 if (!threadIdx.x) AtomicAdd(*data, 1);然后线程 0 将执行原子，并且
    其他人不会。看起来我们可以让其他人做
    别的东西，但锁步扭曲执行不允许这样做。
    Warp 执行被序列化，以便所有线程都使用 if (true) path 将一起执行，并且所有执行该路径的线程if (false) path会一起执行，但是true和false
    路径将被序列化。再说一次，我们真的不能有不同的
    warp 中的线程同时执行不同的指令。

    • 它的网络是，在一个经线内，我们不能让一个线程执行原子操作，而其他线程同时执行其他操作。

    您的许多其他问题似乎期望内存事务在它们起源的指令周期结束时完成。事实并非如此。对于全局内存和共享内存，我们必须在代码中采取特殊步骤以确保之前的写入事务对其他线程可见(这可以作为事务完成的证据。)一种典型的方法是使用屏障说明，例如 __syncthreads()或 __threadfence()但是如果没有这些屏障指令，线程就不会“等待”写入完成。一次(依赖于 a 的操作)读取可以停止线程。写入通常不能停止线程。

    现在让我们看看你的问题:

    so if one of them start an atomic operation, what other will do? Wait?



    不，他们不等。原子操作被分派(dispatch)到处理原子的 SM 上的一个功能单元，并且所有线程一起以锁步继续。由于原子通常意味着读取，是的，读取可以阻止扭曲。但是线程不会等到原子操作完成(即写入)。但是，对该位置的后续读取可能会再次阻止扭曲，等待原子(写入)完成。在保证更新全局内存的全局原子的情况下，它将使原始 SM 中的 L1(如果启用)和 L2 无效，如果它们包含该位置作为条目。

    Is there any way to take other threads for some work, like reads some memory, so the atomic operation will hide it's latency?



    不是真的，因为我在开头说的原因。

    Atomic is really slow, does it? How can I accelerate it? Is there any pattern to work with it?



    是的，如果原子支配事件(例如朴素的减少或朴素的直方图)，原子可以使程序运行得更慢。一般来说，加速原子操作的方法是不使用它们，或者谨慎地使用它们，以这样的方式不支配程序事件。例如，朴素的归约将使用原子将每个元素添加到全局总和中。对于线程块中完成的工作，智能并行缩减将根本不使用原子。在线程块减少结束时，可以使用单个原子将线程块部分总和更新为全局总和。这意味着我可以对任意数量的元素进行快速并行减少，大约 32 个原子添加或更少。这种对原子的 Thrift 使用在整个程序执行中基本上不会被注意到，除了它可以在单个内核调用而不是 2 个内核调用中完成并行减少。

    Shared memory: Does they run parralel or they will be queued?



    他们会排队。这样做的原因是，可以在共享内存上处理原子操作的功能单元数量有限，不足以在单个周期内为来自 warp 的所有请求提供服务。

    我避免尝试回答与原子操作吞吐量相关的问题，因为文档 AFAIK 中没有很好地指定这些数据。可能是，如果您同时发出足够多的或几乎同时发生的原子操作，由于提供原子功能单元的队列已满，某些扭曲将在原子指令上停滞。我不知道这是真的，我无法回答有关它的问题。

    关于c++ - 扭曲如何与原子操作一起工作？，我们在Stack Overflow上找到一个类似的问题： https://stackoverflow.com/questions/20726299/

    amp43brcodeblockquotec++cperformancecudaatomic

    有关c++ - 扭曲如何与原子操作一起工作？的更多相关文章

    1. ruby - 如何使用 Nokogiri 的 xpath 和 at_xpath 方法 - 2

      我正在学习如何使用Nokogiri，根据这段代码我遇到了一些问题:require'rubygems'require'mechanize'post_agent=WWW::Mechanize.newpost_page=post_agent.get('http://www.vbulletin.org/forum/showthread.php?t=230708')puts"\nabsolutepathwithtbodygivesnil"putspost_page.parser.xpath('/html/body/div/div/div/div/div/table/tbody/tr/td/div

    2. ruby - 如何从 ruby​​ 中的字符串运行任意对象方法？ - 2

      总的来说，我对ruby​​还比较陌生，我正在为我正在创建的对象编写一些rspec测试用例。许多测试用例都非常基础，我只是想确保正确填充和返回值。我想知道是否有办法使用循环结构来执行此操作。不必为我要测试的每个方法都设置一个assertEquals。例如:describeitem,"TestingtheItem"doit"willhaveanullvaluetostart"doitem=Item.new#HereIcoulddotheitem.name.shouldbe_nil#thenIcoulddoitem.category.shouldbe_nilendend但我想要一些方法来使用

    3. python - 如何使用 Ruby 或 Python 创建一系列高音调和低音调的蜂鸣声？ - 2

      关闭。这个问题是opinion-based.它目前不接受答案。想要改进这个问题？更新问题，以便editingthispost可以用事实和引用来回答它.关闭4年前。Improvethisquestion我想在固定时间创建一系列低音和高音调的哔哔声。例如:在150毫秒时发出高音调的蜂鸣声在151毫秒时发出低音调的蜂鸣声200毫秒时发出低音调的蜂鸣声250毫秒的高音调蜂鸣声有没有办法在Ruby或Python中做到这一点？我真的不在乎输出编码是什么(.wav、.mp3、.ogg等等)，但我确实想创建一个输出文件。

    4. ruby-on-rails - 由于 "wkhtmltopdf"，PDFKIT 显然无法正常工作 - 2

      我在从html页面生成PDF时遇到问题。我正在使用PDFkit。在安装它的过程中，我注意到我需要wkhtmltopdf。所以我也安装了它。我做了PDFkit的文档所说的一切......现在我在尝试加载PDF时遇到了这个错误。这里是错误:commandfailed:"/usr/local/bin/wkhtmltopdf""--margin-right""0.75in""--page-size""Letter""--margin-top""0.75in""--margin-bottom""0.75in""--encoding""UTF-8""--margin-left""0.75in""-

    5. ruby-on-rails - 如何验证 update_all 是否实际在 Rails 中更新 - 2

      给定这段代码defcreate@upgrades=User.update_all(["role=?","upgraded"],:id=>params[:upgrade])redirect_toadmin_upgrades_path,:notice=>"Successfullyupgradeduser."end我如何在该操作中实际验证它们是否已保存或未重定向到适当的页面和消息？ 最佳答案 在Rails3中，update_all不返回任何有意义的信息，除了已更新的记录数(这可能取决于您的DBMS是否返回该信息)。http://ar.ru

    6. ruby-on-rails - 'compass watch' 是如何工作的/它是如何与 rails 一起使用的 - 2

      我在我的项目目录中完成了compasscreate.和compassinitrails。几个问题:我已将我的.sass文件放在public/stylesheets中。这是放置它们的正确位置吗？当我运行compasswatch时，它不会自动编译这些.sass文件。我必须手动指定文件:compasswatchpublic/stylesheets/myfile.sass等。如何让它自动运行？文件ie.css、print.css和screen.css已放在stylesheets/compiled。如何在编译后不让它们重新出现的情况下删除它们？我自己编译的.sass文件编译成compiled/t

    7. ruby - 如何将脚本文件的末尾读取为数据文件(Perl 或任何其他语言) - 2

      我正在寻找执行以下操作的正确语法(在Perl、Shell或Ruby中):#variabletoaccessthedatalinesappendedasafileEND_OF_SCRIPT_MARKERrawdatastartshereanditcontinues. 最佳答案 Perl用__DATA__做这个:#!/usr/bin/perlusestrict;usewarnings;while(){print;}__DATA__Texttoprintgoeshere 关于ruby-如何将脚

    8. ruby - 如何指定 Rack 处理程序 - 2

      Rackup通过Rack的默认处理程序成功运行任何Rack应用程序。例如:classRackAppdefcall(environment)['200',{'Content-Type'=>'text/html'},["Helloworld"]]endendrunRackApp.new但是当最后一行更改为使用Rack的内置CGI处理程序时，rackup给出“NoMethodErrorat/undefinedmethod`call'fornil:NilClass”:Rack::Handler::CGI.runRackApp.newRack的其他内置处理程序也提出了同样的反对意见。例如Rack

    9. ruby - 如何每月在 Heroku 运行一次 Scheduler 插件？ - 2

      在选择我想要运行操作的频率时，唯一的选项是“每天”、“每小时”和“每10分钟”。谢谢!我想为我的Rails3.1应用程序运行调度程序。 最佳答案 这不是一个优雅的解决方案，但您可以安排它每天运行，并在实际开始工作之前检查日期是否为当月的第一天。 关于ruby-如何每月在Heroku运行一次Scheduler插件？，我们在StackOverflow上找到一个类似的问题： https://stackoverflow.com/questions/8692687/

    10. ruby-on-rails - 如何从 format.xml 中删除 <hash></hash> - 2

      我有一个对象has_many应呈现为xml的子对象。这不是问题。我的问题是我创建了一个Hash包含此数据，就像解析器需要它一样。但是rails自动将整个文件包含在.........我需要摆脱type="array"和我该如何处理？我没有在文档中找到任何内容。 最佳答案 我遇到了同样的问题；这是我的XML:我在用这个:entries.to_xml将散列数据转换为XML，但这会将条目的数据包装到中所以我修改了:entries.to_xml(root:"Contacts")但这仍然将转换后的XML包装在“联系人”中，将我的XML代码修改为

    随机推荐