Take the 2-minute tour ×

I have a device function that checks a byte array using threads, each thread checking a different byte in the array for a certain value and returns bool true or false.

How can I efficiently decide if all the checks have returned true or otherwise?

share|improve this question
CUDA has warp voting functions that can be used to construct a fairly efficient "any"/"all"/"none" type binary reduction at the block level. You probably can't inspect the results all checks across the entire grid within a running kernel, because it requires synchronisation across the whole grid. A second kernel launch or a small host side reduction would be necessary to get the state across the whole grid. –  talonmies Jul 1 '12 at 15:44
@talonmies: That's an excellent answer. Why a comment? –  Roger Dahl Jul 1 '12 at 15:53
Thanks, I'll go look up on the voting function. Anyway, I'm not trying to check across grids, just within a block. –  gamerx Jul 1 '12 at 15:58
It seems the warp vote as the name suggests only allow you to check within a warp, however I have more than the warp size number of threads running. Can it still be used? –  gamerx Jul 1 '12 at 16:04
An easy implementation would be to have a single thread in each warp "or" the warp's result onto a value in shared memory and then have a single thread in the block check that value in the end. –  Roger Dahl Jul 1 '12 at 16:13

1 Answer 1

up vote 2 down vote accepted
// returns true if predicate is true for all threads in a block
__device__ bool unanimous(bool predicate) { ... }

__device__ bool all_the_same(unsigned char* bytes, unsigned char value, int n) {
    return unanimous(bytes[threadIdx.x] == value);

The implementation of unanimous() depends on the compute capability of your hardware. For compute capability 2.0 or higher devices, it is trivial:

__device__ bool unanimous(bool predicate) { return __syncthreads_and(predicate); }

For compute capability 1.0 and 1.1 devices, you will need to implement an AND reduction (exercise for the reader, since it's well documented). For the special case of compute capability 1.3, you can optimize the AND reduction using warp vote instructions, using the __all() intrinsic function provided in the CUDA headers.


OK, since gamerx is asking in the comments. On sm_13 hardware, you can do this.

// note: supports maximum of 1024 threads in block as written
__device__ bool unanimous(bool predicate) {
    __shared__ bool warp_votes[32];
    if (threadIdx.x < warpSize) warp_votes[threadIdx.x] = true;
    warp_votes[threadIdx.x / warpSize] = __all(pred);
    if (threadIdx.x < warpSize) warp_votes[0] = __all(warp_votes[threadIdx.x];
    return warp_votes[0];
share|improve this answer
I've since figured it out, but thanks anyway. –  gamerx Jul 4 '12 at 3:04

Your Answer


