Cuda - 选择性内存存储答案

【问题标题】：Cuda - selective memory storeCuda - 选择性内存存储
【发布时间】：2012-08-18 09:54:18
【问题描述】：

在我的内核中，如果满足条件，我会更新输出缓冲区的一项

if (condition(input[i])) //?
    output[i] = 1;

否则输出可能保持不变，值为 0。

更新的密度非常不可预测，具体取决于输入。此外，哪个输出位置将被更新也是未知的。（在某些情况下，我可能会强迫他们）

我的问题是，是写所有项目更好，实现合并，还是选择性地写？

output[i] = condition(input[i]); //?

您介意讨论一下您的陈述吗？

【问题讨论】：

标签： memory if-statement cuda store coalescing

【解决方案1】：

即使 warp 中的某些线程不参与加载或存储，只要所有参与的线程都满足合并的要求，就可以实现合并。所以条件写入应该对内存吞吐量没有影响。

但是，由于涉及分支，执行条件写入可能涉及额外的指令（例如，这可能解释了 Eugene 在他的回答中测量的性能差异）。

【讨论】：

如果写入输出缓冲区的速率很低怎么办。
您的意思是，“如果每个经纱执行存储的线程比例很低怎么办”？因为内存事务是每个 warp 执行的，假设每个 warp 至少有一个线程进行存储，那么我的答案是一样的。如果每个 warp 至少有一个线程执行存储，则成本（假设完美合并）与所有线程都完成存储相同。但是，如果很大一部分 warp 没有存储，那么if() 方法绝对是优越的。
谢谢@Harrism。那么，将结果存储在共享内存中并一次性将它们存储回全局内存没有任何好处吗？不过，它可以将写入类型从 float 更改为 float4 以获取 4 个值。在全局内存写入之前仅进行一次同步就足够了吗？您能否为全局内存存储、限制写入或不限制写入的方法指出一个好方法？（我也可以将其转换为问题：）
我认为你应该做一些实验。您的问题的答案实际上取决于您的计算细节。
是的，做一些实验会更好，我可能会分享结果以供进一步讨论。亲切的问候。

【解决方案2】：

在执行条件设置（选项 1）的设置内核上运行 1.727 us 和选项 2 1.399 us。这是我的代码（setConditional 更快）：

__global__ void conditionalSet(unsigned int* array) {
    if ((threadIdx.x & 3) == 0) {
        array[threadIdx.x] = 1;
    }
}

__global__ void setConditional(unsigned int* array) {
    array[threadIdx.x] = (threadIdx.x & 3) == 0 ? 1 : 0;
}

【讨论】：

"ms" 代表“毫秒”，“us”代表“微秒”。那么什么是“mus”呢？
您的时间安排不正确——内核启动开销远远超过几微秒。你的时间安排如何？
我从视觉分析器中获得了这些数字（当单击时间轴中的内核时 - 它们不包括开销）。但是，正如您所指出的，它们完全是虚假的，因为运行时间的差异很可能归因于指令数的差异，而不是内存访问模式的差异。一旦我修复了本地环境，我将进行更多实验。