草庐IT

c++ - CUB (CUDA UnBound) 相当于 thrust::gather

coder 2024-02-21 原文

由于 Thrust 库存在一些性能问题(有关详细信息，请参阅 this page)，我计划重构一个 CUDA 应用程序以使用 CUB 而不是 Thrust。具体来说，就是替换 thrust::sort_by_key 和 thrust::inclusive_scan 调用)。在我的应用程序的特定点上，我需要按键对 3 个数组进行排序。这就是我用推力做到这一点的方式:

thrust::sort_by_key(key_iter, key_iter + numKeys, indices);
thrust::gather_wrapper(indices, indices + numKeys, 
      thrust::make_zip_iterator(thrust::make_tuple(values1Ptr, values2Ptr, values3Ptr)),
      thrust::make_zip_iterator(thrust::make_tuple(valuesOut1Ptr, valuesOut2Ptr, valuesOut3Ptr))
);

在哪里

  • key iter 是一个 thrust::device_ptr，它指向我要排序的键
  • indices 指向设备内存中的一个序列(从 0 到 numKeys-1)
  • values{1,2,3}Ptr 是我要排序的值的 device_ptr
  • values{1,2,3}OutPtr 是排序值的 device_ptr

随着CUB SortPairs功能我可以对单个值缓冲区进行排序，但不能一次性对所有 3 个进行排序。问题是我没有看到任何 CUB“类似收集”的实用程序。有什么建议吗？

编辑:

我想我可以实现自己的 gather 内核，但是除了以下方法还有其他更好的方法吗:

template <typename Index, typename Value> 
__global__ void  gather_kernel(const unsigned int N, const Index * map, 
const Value * src, Value * dst) 
{ 
    unsigned int i = blockDim.x * blockIdx.x + threadIdx.x; 
    if (i < N) 
    { 
        dst[i] = src[map[i]]; 
    } 
}

非合并的加载和存储让我感到畏缩，但如果 map 上没有已知结构，这可能是不可避免的。

最佳答案

看来你想要实现的目标取决于thrust::zip_iterator .你也可以

  1. 只替换thrust::sort_by_key通过 cub::DeviceRadixSort::SortPairs并保持 thrust::gather , 或
  2. zip values{1,2,3}在使用 cub::DeviceRadixSort::SortPairs 之前进入结构数组

更新

看完thrust::gather的执行后,

$CUDA_HOME/include/thrust/system/detail/generic/gather.inl

你可以看到它只是一个朴素的内核

__global__ gather(int* index, float* in, float* out, int len) {
  int i=...;
  if (i<len) { out[i] = in[index[i]]; }
}

然后我认为你上面的代码可以用一个内核替换而不需要太多努力。

在此内核中，您可以首先使用 CUB block 大小原语 cub::BlockRadixSort<...>::SortBlockedToStriped 获取存储在寄存器中的排序索引，然后执行天真的重新排序拷贝作为 thrust::gather填写values{1,2,3}Out .

使用 SortBlockedToStriped而不是 Sort 复制 values 时可以合并写入(虽然不是为了阅读) .

关于c++ - CUB (CUDA UnBound) 相当于 thrust::gather，我们在Stack Overflow上找到一个类似的问题： https://stackoverflow.com/questions/19210652/

amp相当codethrustgatherc++cudaparallel-processinggpu

有关c++ - CUB (CUDA UnBound) 相当于 thrust::gather的更多相关文章

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

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

  2. Python 相当于 Perl/Ruby ||= - 2

    这个问题在这里已经有了答案:关闭10年前。PossibleDuplicate:Pythonconditionalassignmentoperator对于这样一个简单的问题表示歉意，但是谷歌搜索||=并不是很有帮助；)Python中是否有与Ruby和Perl中的||=语句等效的语句？例如:foo="hey"foo||="what"#assignfooifit'sundefined#fooisstill"hey"bar||="yeah"#baris"yeah"另外，类似这样的东西的通用术语是什么？条件分配是我的第一个猜测，但Wikipediapage跟我想的不太一样。

  3. java - 什么相当于 ruby​​ 的 rack 或 python 的 Java wsgi？ - 2

    什么是ruby​​的rack或python的Java的wsgi？还有一个路由库。 最佳答案 来自Python标准PEP333:Bycontrast,althoughJavahasjustasmanywebapplicationframeworksavailable,Java's"servlet"APImakesitpossibleforapplicationswrittenwithanyJavawebapplicationframeworktoruninanywebserverthatsupportstheservletAPI.ht

  4. ruby - 使用 `+=` 和 `send` 方法 - 2

    如何将send与+=一起使用？a=20;a.send"+=",10undefinedmethod`+='for20:Fixnuma=20;a+=10=>30 最佳答案 恐怕你不能。+=不是方法，而是语法糖。参见http://www.ruby-doc.org/docs/ProgrammingRuby/html/tut_expressions.html它说Incommonwithmanyotherlanguages,Rubyhasasyntacticshortcut:a=a+2maybewrittenasa+=2.你能做的最好的事情是:

  5. java - Ruby 相当于 Java 的 Collections.unmodifiableList 和 Collections.unmodifiableMap - 2

    Java的Collections.unmodifiableList和Collections.unmodifiableMap在Ruby标准API中是否有等价物？ 最佳答案 使用freeze应用程序接口(interface):Preventsfurthermodificationstoobj.ARuntimeErrorwillberaisedifmodificationisattempted.Thereisnowaytounfreezeafrozenobject.SeealsoObject#frozen?.Thismethodretur

  6. ruby - 如何计算 Liquid 中的变量 +1 - 2

    我对如何计算通过{%assignvar=0%}赋值的变量加一完全感到困惑。这应该是最简单的任务。到目前为止，这是我尝试过的:{%assignamount=0%}{%forvariantinproduct.variants%}{%assignamount=amount+1%}{%endfor%}Amount:{{amount}}结果总是0。也许我忽略了一些明显的东西。也许有更好的方法。我想要存档的只是获取运行的迭代次数。 最佳答案 因为{{incrementamount}}将输出您的变量值并且不会影响{%assign%}定义的变量，我

  7. python - Ruby 相当于 Python str[3 :] - 2

    是否有Ruby等效于Python的方法来获取在字符串末尾结束的子字符串，如str[3:]？必须输入字符串的长度并不方便。 最佳答案 传递最后一个元素=-1的范围str[3..-1] 关于python-Ruby相当于Pythonstr[3:]，我们在StackOverflow上找到一个类似的问题： https://stackoverflow.com/questions/12978768/

  8. arrays - Ruby 数组 += vs 推送 - 2

    我有一个数组数组，想将元素附加到子数组。+=做我想做的，但我想了解为什么push不做。我期望的行为(并与+=一起工作):b=Array.new(3,[])b[0]+=["apple"]b[1]+=["orange"]b[2]+=["frog"]b=>[["苹果"],["橙子"],["Frog"]]通过推送，我将推送的元素附加到每个子数组(为什么？):a=Array.new(3,[])a[0].push("apple")a[1].push("orange")a[2].push("frog")a=>[[“苹果”、“橙子”、“Frog”]、[“苹果”、“橙子”、“Frog”]、[“苹果”、“

  9. += 的 Ruby 方法 - 2

    有没有办法让Ruby能够做这样的事情？classPlane@moved=0@x=0defx+=(v)#thisiserror@x+=v@moved+=1enddefto_s"moved#{@moved}times,currentxis#{@x}"endendplane=Plane.newplane.x+=5plane.x+=10putsplane.to_s#moved2times,currentxis15 最佳答案 您不能在Ruby中覆盖复合赋值运算符。任务在内部处理。您应该覆盖+，而不是+=。plane.a+=b与plane.a=

  10. ruby - Java 8 相当于 ruby​​ each_with_index - 2

    我想知道，是否有一些流操作可以像ruby​​中的each_with_index那样做。其中each_with_index遍历值以及值的索引。 最佳答案 没有专门用于该目的的流操作。但您可以通过多种方式模仿该功能。索引变量:以下方法适用于顺序流。int[]index={0};stream.forEach(item->System.out.printf("%s%d\n",item,index[0]++));外部迭代:以下方法适用于并行流，只要原始集合支持随机访问。Listtokens=...;IntStream.range(0,toke

随机推荐