草庐IT

c++ - 使用CUDA对两个数组求和

coder 2024-02-05 原文

我在学习this guide的同时正在学习CUDA。

我还没有完成,但是我决定尝试一下到目前为止所看到的。

我试图重写第一个使用256个线程的示例。我想这样做,以便每个线程都在数组的连续切片上进行操作。

目标是将2个数组与1,048,576个项相加。

为了进行比较,这是原始代码,其中根据跨步访问每个数组项:

__global__
void add(int n, float *x, float *y)
{
  int index = threadIdx.x;
  int stride = blockDim.x;
  for (int i = index; i < n; i += stride)
      y[i] = x[i] + y[i];
}


这是我的功能:
__global__
void add2(int n, float* x, float* y) {
    int sliceSize = n / blockDim.x;
    int lower = threadIdx.x * sliceSize;
    int upper = lower + sliceSize;
    for (int i = lower; i < upper; i++) {
        y[i] = x[i] + y[i];
    }
}

事实证明,最后一个摘要的执行速度比上一个摘要慢了近7倍(22毫秒对3毫秒)。我以为,通过在连续的片上访问它们,它将执行相同或更快的操作。

我正在使用add<<<1, threads>>>(n, x, y)add<<<1, threads>>>(n, x, y)(256个线程)调用该函数。
sliceSize的值始终是4096。在这种情况下,应该发生的是:
  • threadIdx.x = 0从0到4095
  • threadIdx.x = 1从4096变为8191
  • ...
  • threadIdx.x = 255从1044480变为1048576

  • 我打开了NVidia Visual Profiler,然后了解到我的内存访问模式效率不高(全局内存加载/存储效率低)。第一个代码段中不存在此警告。为什么会这样呢?

    我以为第一个被删除的对象会在整个数组中跳转,从而造成错误的访问模式。实际上,这似乎很好。

    我已经阅读了可视分析器随附的一些有关内存优化的文档,但是我不太明白为什么它这么慢。

    最佳答案

    您正在探索合并和未合并的内存访问之间的区别。或者我们可以简单地说“最有效”和“效率较低”的内存访问。

    在GPU上,所有指令都在整个warp范围内执行。因此,当warp中的一个线程正在读取内存中的位置时,warp中的所有线程都正在从内存中读取。粗略地说,最佳模式是经线中的所有线程都从相邻位置读取时。这导致以下情况:GPU内存 Controller 在检查特定时间段内warp中每个线程请求的内存地址后,可以将合并地址在一起,从而导致需要从内存请求的行数最少。高速缓存(或要从DRAM请求的最小段数)。

    在幻灯片36(或37)here上以图形方式描绘了这种情况。

    在您的第一个代码段中表示了100%合并的大小写。从全局内存读取的示例在这里:

      y[i] = x[i] + y[i];
             ^
             reading from the vector x in global memory
    

    让我们考虑循环的第一遍,并考虑第一次扭曲的情况(即线程块中的前32个线程)。在这种情况下,ithreadIdx.x给出。因此,线程0的索引为0,线程1的索引为1,依此类推。因此,每个线程都在读取全局内存中的相邻位置。假设我们错过了所有缓存,这将转化为DRAM读取请求,并且存储器 Controller 可以针对DRAM中的段(或等效地,针对缓存中的行)生成最少数量的请求(更精确地说:事务)。从“总线带宽利用率”为100%的角度来看,这是最佳的。在该读取周期中,请求的每个字节实际上都被扭曲中的线程使用。

    “不公开”访问通常可以指任何不符合以上描述的情况。转化为上述更细粒度的“总线带宽利用率”数字,根据具体情况和不同情况,非强制访问的程度可能有所不同,从略低于100%的最佳情况到最坏的情况是12.5%或3.125%。 GPU。

    在幻灯片44(或45)here中给出了根据此描述的最坏情况的不分时段访问模式示例。这并不能完全描述您的最坏情况的代码段,但是对于足够大的sliceSize来说,它是等效的。代码行是相同的。考虑到相同的读取请求(对于x,在循环的第一次迭代中,对于warp 0,为warp 0),唯一的区别在于i在整个warp中采用的值:
    int sliceSize = n / blockDim.x;
    int lower = threadIdx.x * sliceSize;
    ...
    for (int i = lower; i < upper; i++) {
        y[i] = x[i] + y[i];
    

    因此,ilower开始,即threadIdx.x * sliceSize。让我们假设sliceSize大于1。然后,第一个线程将读取位置0。第二个线程将读取位置sliceSize。第三个线程将读取位置2*sliceSize等。这些位置由sliceSize距离分隔。即使sliceSize只有2,该模式的效率仍然较低,因为内存 Controller 现在必须请求两倍的行数或段数才能满足在warp 0上的特定读取周期。如果sliceSize足够大,则内存 Controller 必须请求a每个线程的唯一行或段,这是最坏的情况。

    作为最后的说明/要点,可以对“快速分析”做出有益的观察:
  • 在大多数情况下可实现最佳的合并访问,我们要确保内存索引的计算不会使而不是涉及到threadIdx.x乘以除1以外的任何数量。
  • 相反,
  • ,如果我们可以证明在任何给定的索引计算中threadIdx.x都乘以不等于1的某个数字,那么不管其他考虑如何,这几乎都是普遍的指示,即所生成的访问模式将是非最佳的。

  • 为了清楚起见,请重复此操作:
    index = any_constant_across_the_warp + threadIdx.x;
    

    通常将是最佳的访问模式。
    index = any_constant_across_the_warp + C*threadIdx.x;
    

    通常不会是最佳的访问模式。注意any_constant_across_the_warp可以由数量上的任意算术组成,例如:循环索引,blockIdx.?blockDim.?gridDim.?和任何其他常量。必须考虑2D或3D线程块模式,在这种情况下将考虑使用threadIdx.y,但通常不难将这种理解扩展到2D情况。对于典型的线程块形状,为了进行快速分析,通常不希望在threadIdx.xthreadIdx.y上使用常数乘数。

    整个讨论适用于全局内存读/写。共享内存还具有用于最佳访问的规则,这些规则在某种程度上类似于上述说明,但在某些方面却大不相同。但是,对于全局内存,完全最佳的100%合并模式也将是共享内存读/写的最佳模式,这通常是正确的。换句话说,扭曲中的相邻访问通常对于共享内存也是最佳的(但它不是共享内存的唯一可能的最佳模式)。

    已经链接到here的演示文稿将对该主题进行更全面的处理,网络上的许多其他演示文稿和处理也将提供该处理。

    关于c++ - 使用CUDA对两个数组求和,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/51027354/

    有关c++ - 使用CUDA对两个数组求和的更多相关文章

    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 - 使用 RubyZip 生成 ZIP 文件时设置压缩级别 - 2

      我有一个Ruby程序,它使用rubyzip压缩XML文件的目录树。gem。我的问题是文件开始变得很重,我想提高压缩级别,因为压缩时间不是问题。我在rubyzipdocumentation中找不到一种为创建的ZIP文件指定压缩级别的方法。有人知道如何更改此设置吗?是否有另一个允许指定压缩级别的Ruby库? 最佳答案 这是我通过查看ruby​​zip内部创建的代码。level=Zlib::BEST_COMPRESSIONZip::ZipOutputStream.open(zip_file)do|zip|Dir.glob("**/*")d

    3. ruby - 为什么我可以在 Ruby 中使用 Object#send 访问私有(private)/ protected 方法? - 2

      类classAprivatedeffooputs:fooendpublicdefbarputs:barendprivatedefzimputs:zimendprotecteddefdibputs:dibendendA的实例a=A.new测试a.foorescueputs:faila.barrescueputs:faila.zimrescueputs:faila.dibrescueputs:faila.gazrescueputs:fail测试输出failbarfailfailfail.发送测试[:foo,:bar,:zim,:dib,:gaz].each{|m|a.send(m)resc

    4. ruby-on-rails - 使用 Ruby on Rails 进行自动化测试 - 最佳实践 - 2

      很好奇,就使用ruby​​onrails自动化单元测试而言,你们正在做什么?您是否创建了一个脚本来在cron中运行rake作业并将结果邮寄给您?git中的预提交Hook?只是手动调用?我完全理解测试,但想知道在错误发生之前捕获错误的最佳实践是什么。让我们理所当然地认为测试本身是完美无缺的,并且可以正常工作。下一步是什么以确保他们在正确的时间将可能有害的结果传达给您? 最佳答案 不确定您到底想听什么,但是有几个级别的自动代码库控制:在处理某项功能时,您可以使用类似autotest的内容获得关于哪些有效,哪些无效的即时反馈。要确保您的提

    5. ruby - 在 Ruby 中使用匿名模块 - 2

      假设我做了一个模块如下:m=Module.newdoclassCendend三个问题:除了对m的引用之外,还有什么方法可以访问C和m中的其他内容?我可以在创建匿名模块后为其命名吗(就像我输入“module...”一样)?如何在使用完匿名模块后将其删除,使其定义的常量不再存在? 最佳答案 三个答案:是的,使用ObjectSpace.此代码使c引用你的类(class)C不引用m:c=nilObjectSpace.each_object{|obj|c=objif(Class===objandobj.name=~/::C$/)}当然这取决于

    6. ruby - 使用 ruby​​ 和 savon 的 SOAP 服务 - 2

      我正在尝试使用ruby​​和Savon来使用网络服务。测试服务为http://www.webservicex.net/WS/WSDetails.aspx?WSID=9&CATID=2require'rubygems'require'savon'client=Savon::Client.new"http://www.webservicex.net/stockquote.asmx?WSDL"client.get_quotedo|soap|soap.body={:symbol=>"AAPL"}end返回SOAP异常。检查soap信封,在我看来soap请求没有正确的命名空间。任何人都可以建议我

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

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

    8. 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

    9. ruby - 使用 ruby​​ 将 HTML 转换为纯文本并维护结构/格式 - 2

      我想将html转换为纯文本。不过,我不想只删除标签,我想智能地保留尽可能多的格式。为插入换行符标签,检测段落并格式化它们等。输入非常简单,通常是格式良好的html(不是整个文档,只是一堆内容,通常没有anchor或图像)。我可以将几个正则表达式放在一起,让我达到80%,但我认为可能有一些现有的解决方案更智能。 最佳答案 首先,不要尝试为此使用正则表达式。很有可能你会想出一个脆弱/脆弱的解决方案,它会随着HTML的变化而崩溃,或者很难管理和维护。您可以使用Nokogiri快速解析HTML并提取文本:require'nokogiri'h

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

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

    随机推荐