memory - Cuda - 选择性内存存储

Question

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

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

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

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

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

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

你介意讨论你的陈述吗？

Answer 1

在我的设置内核上，条件设置（选项 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;
}

Answer 2

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

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