在 CUDA 上通过多个线程使用设备变量答案

【问题标题】：Using device variable by multiple threads on CUDA在 CUDA 上通过多个线程使用设备变量
【发布时间】：2010-03-16 04:43:22
【问题描述】：

我在玩 cuda。

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

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

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

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

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

【问题讨论】：

标签： cuda race-condition

【解决方案1】：

并不像看起来那么简单。我刚刚完成了一个，我可以告诉你需要什么阅读scan Gpu Gems 3 Article 的特别章节39.3.1 Stream Compaction。

从 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

玩得开心！

【讨论】：

【解决方案2】：

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

【讨论】：