CUDA上多线程使用设备变量

Question

我正在玩cuda。

目前我有一个问题。我正在测试一个大型数组的特定响应，当我收到响应时，我必须将数据复制到另一个数组上。

例如，我的 5 个元素的测试数组如下所示：
[ ][ ][v1][ ][ ][v2]

结果必须如下所示：
[v1][v2]

问题是如何计算第二个数组存储结果的地址？并行检查第一个数组的所有元素。

我正在考虑声明一个设备变量 int addr = 0。每次我找到响应时，我都会增加 addr。但我对此不确定，因为这意味着 addr 可以同时被多个线程访问。这会引起问题吗？或者线程会等到另一个线程完成使用该变量吗？

Answer 1

并不像看起来那么微不足道。我刚刚完成一个实施，我可以告诉你需要什么
阅读扫描 Gpu Gems 3 文章，特别是章节 39.3.1流压缩。

要实现您自己的从 SDK 中的 LargeArrayScan 示例开始，这将只为您提供预扫描。假设设备内存中有选择数组（1 和 0 的数组，表示 1-选择 0-丢弃），dev_selection_array 一个 dev_elements_array 元素将被选择一个 dev_prescan_array 和一个 dev_result_array 所有大小 N 然后你就可以

prescan(dev_prescan_array,dev_selection_array, N);
scatter(dev_result_array, dev_prescan_array,
         dev_selection_array, dev_elements_array, N);

在分散的地方

 __global__ void scatter_kernel( T*dev_result_array, 
                   const T* dev_prescan_array, 
                   const T* dev_selection_array,
                   const T* dev_elements_array, std::size_t size){

unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx >= size) return;
if (dev_selection_array[idx] == 1){
    dev_result_array[dev_prescan_array[idx]] = dev_elements_array[idx];
}
}

进行预扫描的其他很好的应用，请参阅论文 Ble93

玩得开心！

Answer 2

您正在谈论经典的流压缩。一般来说，我建议查看 Thrust 或CUDPP（这些链接转到压缩文档）。这两个都是开源的，如果您想自己开发，我还建议您查看“扫描”SDK 示例。