草庐IT

memory - CUDA 合并访问全局内存

coder 2023-06-04 原文

我已阅读 CUDA 编程指南，但我错过了一件事。假设我在全局内存中有 32 位 int 数组，我想通过合并访问将它复制到共享内存。 全局数组的索引从 0 到 1024，假设我有 4 个 block ，每个 block 有 256 个线程。

__shared__ int sData[256];

何时执行合并访问？

1.

sData[threadIdx.x] = gData[threadIdx.x * blockIdx.x+gridDim.x*blockIdx.y];

全局内存中的地址从 0 复制到 255，每个被 32 个线程在 warp 中复制，这样就可以了？

2.

sData[threadIdx.x] = gData[threadIdx.x * blockIdx.x+gridDim.x*blockIdx.y + someIndex];

如果 someIndex 不是 32 的倍数，它不会合并？地址错位？对吗？

最佳答案

您最终想要什么取决于您的输入数据是一维数组还是二维数组，以及您的网格和 block 是一维还是二维。最简单的情况都是一维的:

shmem[threadIdx.x] = gmem[blockDim.x * blockIdx.x + threadIdx.x];

这是合并的。我使用的经验法则是将变化最快的坐标(threadIdx)作为偏移量添加到 block 偏移量(blockDim * blockIdx)上。最终结果是 block 中线程之间的索引步长为 1。如果步长变大，那么您将失去合并。

简单的规则(在 Fermi 和更高版本的 GPU 上)是，如果一个 warp 中所有线程的地址落入相同对齐的 128 字节范围内，则将产生单个内存事务(假设为加载启用了缓存，这是默认值)。如果它们落入两个对齐的 128 字节范围内，则会产生两个内存事务，等等。

在 GT2xx 和更早的 GPU 上，它变得更加复杂。但是您可以在编程指南中找到详细信息。

其他示例:

未合并:

shmem[threadIdx.x] = gmem[blockDim.x + blockIdx.x * threadIdx.x];

没有合并，但在 GT200 及更高版本上还不错:

stride = 2;
shmem[threadIdx.x] = gmem[blockDim.x * blockIdx.x + stride * threadIdx.x];

根本没有合并:

stride = 32;
shmem[threadIdx.x] = gmem[blockDim.x * blockIdx.x + stride * threadIdx.x];

合并的、2D 网格、1D block :

int elementPitch = blockDim.x * gridDim.x;
shmem[threadIdx.x] = gmem[blockIdx.y * elementPitch + 
                          blockIdx.x * blockDim.x + threadIdx.x];

合并的二维网格和 block :

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int elementPitch = blockDim.x * gridDim.x;
shmem[threadIdx.y * blockDim.x + threadIdx.x] = gmem[y * elementPitch + x];

关于memory - CUDA 合并访问全局内存，我们在Stack Overflow上找到一个类似的问题： https://stackoverflow.com/questions/10325244/

memoryCUDAthreadIdxblockIdxcodecopycoalescing

有关memory - CUDA 合并访问全局内存的更多相关文章

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

  2. ruby-on-rails - Ruby net/ldap 模块中的内存泄漏 - 2

    作为我的Rails应用程序的一部分，我编写了一个小导入程序，它从我们的LDAP系统中吸取数据并将其塞入一个用户表中。不幸的是，与LDAP相关的代码在遍历我们的32K用户时泄漏了大量内存，我一直无法弄清楚如何解决这个问题。这个问题似乎在某种程度上与LDAP库有关，因为当我删除对LDAP内容的调用时，内存使用情况会很好地稳定下来。此外，不断增加的对象是Net::BER::BerIdentifiedString和Net::BER::BerIdentifiedArray，它们都是LDAP库的一部分。当我运行导入时，内存使用量最终达到超过1GB的峰值。如果问题存在，我需要找到一些方法来更正我的代

  3. ruby-on-rails - 在混合/模块中覆盖模型的属性访问器 - 2

    我有一个包含模块的模型。我想在模块中覆盖模型的访问器方法。例如:classBlah这显然行不通。有什么想法可以实现吗？ 最佳答案 您的代码看起来是正确的。我们正在毫无困难地使用这个确切的模式。如果我没记错的话，Rails使用#method_missing作为属性setter，因此您的模块将优先，阻止ActiveRecord的setter。如果您正在使用ActiveSupport::Concern(参见thisblogpost)，那么您的实例方法需要进入一个特殊的模块:classBlah

  4. ruby - 续集在添加关联时访问many_to_many连接表 - 2

    我正在使用Sequel构建一个愿望list系统。我有一个wishlists和itemstable和一个items_wishlists连接表(该名称是续集选择的名称)。items_wishlists表还有一个用于facebookid的额外列(因此我可以存储opengraph操作)，这是一个NOTNULL列。我还有Wishlist和Item具有续集many_to_many关联的模型已建立。Wishlist类也有:selectmany_to_many关联的选项设置为select:[:items.*,:items_wishlists__facebook_action_id].有没有一种方法可以

  5. ruby - 如果指定键的值在数组中相同，如何合并哈希 - 2

    我有一个这样的哈希数组:[{:foo=>2,:date=>Sat,01Sep2014},{:foo2=>2,:date=>Sat,02Sep2014},{:foo3=>3,:date=>Sat,01Sep2014},{:foo4=>4,:date=>Sat,03Sep2014},{:foo5=>5,:date=>Sat,02Sep2014}]如果:date相同，我想合并哈希值。我对上面数组的期望是:[{:foo=>2,:foo3=>3,:date=>Sat,01Sep2014},{:foo2=>2,:foo5=>5:date=>Sat,02Sep2014},{:foo4=>4,:dat

  6. ruby-on-rails - Ruby 中的内存模型 - 2

    ruby如何管理内存。例如:如果我们在执行过程中采用C程序，则以下是内存模型。类似于这个ruby如何处理内存。C:__________________|||stack|||------------------||||------------------|||||Heap|||||__________________|||data|__________________|text|__________________Ruby:? 最佳答案 Ruby中没有“内存”这样的东西。Class#allocate分配一个对象并返回该对象。这就是程序

  7. git使用常见问题（提交代码，合并冲突） - 2

    文章目录git常用命令（简介，详细参数往下看）Git提交代码步骤gitpullgitstatusgitaddgitcommitgitpushgit代码冲突合并问题方法一：放弃本地代码方法二：合并代码常用命令以及详细参数gitadd将文件添加到仓库：gitdiff比较文件异同gitlog查看历史记录gitreset代码回滚版本库相关操作远程仓库相关操作分支相关操作创建分支查看分支：gitbranch合并分支：gitmerge删除分支：gitbranch-ddev查看分支合并图：gitlog–graph–pretty=oneline–abbrev-commit撤消某次提交git用户名密码相关配置g

  8. ruby - 有没有办法从 ruby​​ case 语句中访问表达式？ - 2

    我想从then子句中访问c​​ase语句表达式，即food="cheese"casefoodwhen"dip"then"carrotsticks"when"cheese"then"#{expr}crackers"else"mayo"end在这种情况下，expr是食物的当前值(value)。在这种情况下，我知道，我可以简单地访问变量food，但是在某些情况下，该值可能无法再访问(array.shift等)。除了将expr移出到局部变量然后访问它之外，是否有直接访问caseexpr值的方法？罗亚附注我知道这个具体示例很简单，只是一个示例场景。 最佳答案

  9. ruby - 从外部访问类的实例变量 - 2

    我理解(我认为)Ruby中类变量和类的实例变量之间的区别。我想知道如何从该类外部访问该类的实例变量。从内部(即在类方法中而不是实例方法中)，它可以直接访问，但是从外部，有没有办法做MyClass.class.[@$#]variablename?我没有任何具体原因要这样做，只是学习Ruby并想知道是否可行。 最佳答案 classMyClass@my_class_instance_var="foo"class上述yield:>>foo我相信Arkku演示了如何从类外部访问类变量(@@)，而不是类实例变量(@)。我从这篇文章中提取了上述内

  10. ruby-on-rails - 使用 HTTP.get_response 检索 Facebook 访问 token 时出现 Rails EOF 错误 - 2

    我试图在我的网站上实现使用Facebook登录功能，但在尝试从Facebook取回访问token时遇到障碍。这是我的代码:ifparams[:error_reason]=="user_denied"thenflash[:error]="TologinwithFacebook,youmustclick'Allow'toletthesiteaccessyourinformation"redirect_to:loginelsifparams[:code]thentoken_uri=URI.parse("https://graph.facebook.com/oauth/access_token

随机推荐