比较 CUDA 中 2 个线程之间的值

Question

考虑一个块中的所有线程都有一个整数变量 a，其值在线程之间可能不同。如何比较块中的第一个线程是否与块中的最后一个线程具有相同的 a 值？我只需要比较这 2 个线程，因为 a 的值不会随着线程索引一起递减。

Answer 1

这是一个包含多个扭曲的块的解决方案，您必须使用共享内存或全局内存，其中共享内存会更快。可以使用静态共享内存，因为大小已知：

// .. 
thread_block block = this_thread_block();
__shared__ int s[2];
__shared__ bool result;
uint32_t tid = threadIdx.x;

if(tid == 0)
    s[0] = a;
if(tid == blockDim.x - 1)
    s[1] = a;
block.sync(); // synchronize the block instead the whole grid

if(tid == 0)
    result = s[0] == s[1];
block.sync();

如果块只有一个 warp，或者如果您必须比较 warp 的第一个线程和最后一个线程的 a，cooperative groups

有更好的解决方案

#include <cooperative_groups.h>
namespace cg = cooperative_groups;

auto active = cg::coalesced_threads();
int theOtherA = active.shfl(a, active.size() -1); // everybody has a copy a from 
//last thread of the warp
bool result = a == theOtherA; // thread 0 has correct result
result = active.shfl(result, 0); // get a copy from the register of thread 0

比较 CUDA 中 2 个线程之间的值

Compare value between 2 threads in CUDA

comparison

synchronization

cuda