CUDA中共享内存的并行扫描

Question

我有一个问题，每个线程块(1维)必须在共享内存中的数组上运行扫描和其他几个任务。(数组最多有1024个元素。)

有没有好的库支持这种类型的操作？

我检查了推力和CudPP，但它们都只在数据最初在全局内存中时工作，这不是我想要的，因为我不想启动一个新内核来做很少的工作，然后将数据复制回全局内存，因为它有相当大的开销。

如果没有支持这种类型的操作的公共库，那么如果我只想写一次这个操作，然后在几个不同的问题中使用它，你会建议什么呢？

我的第一个想法是编写一个简单的设备函数来运行扫描操作，但是是否可以从一个不是内核函数的设备函数访问共享内存并执行__syncthreads()操作？我的另一个想法是在宏中编写整个函数，然后它就可以工作了，因为预处理器将代码复制到内核代码中，但编写如此复杂的宏并不是使用宏的最佳方式。

附言:我使用的是流式多处理器2.1版本，所以我不能从另一个内核启动新的内核。

Answer 1

有没有好的库支持这种类型的操作？

如果没有支持这种类型的操作的公共库，那么如果我只想写一次这个操作，然后在几个不同的问题中使用它，你会建议什么呢？

有一个名为CUB的库，它代表CUDA UnBound。你会发现它是here。它提供了扫描功能。

如果你想自己写一些类似的东西，CUB可以帮助你看看它是如何完成的。

我的第一个想法是写一个简单的主机函数来运行扫描操作，但是从一个不是内核函数的设备函数访问共享内存和执行__syncthreads()操作是可能的吗？

我不确定你所说的“运行扫描操作的主机函数”是什么意思。

您可以从主机代码中做的唯一一件事就是启动内核和内存副本。除非设备函数是从拥有相关共享内存的块中调用的，否则答案是否定的。

共享内存是按块分配的，并且只能从块内部访问。