GPU-programmers use popcount-based programming all the time these days, but the abstractions are built on top and are hardware accelerated.
CUDA's __activemask(); returns the 32-bit value of your current 32-wide EXEC mask. That is to say, if your current warp is:
int foo = 0;
if(threadIdx.x %= 2){
foo = __activemask();
}
foo will be "0b01010101...." or 0x55555555. This __activemask() has a number of useful properties should you use __popc with it.
popcount(__activemask()); returns the number of threads executing.
lanemask_lt() returns "0b0000000000000001" for the 0th lane. 0b0000000000000011 for the 1st lane. 0b0000000000000111... for the 2nd lane... and 111111111...111 for the last 31st lane.
popcount(__activemask() & lanemask_lt()); returns the "active lane count". All together now, we can make a parallel SIMD-stack that can push/pop together in parallel.
int head = 0;
char buffer[0x1000];
while(fooBar()){ // Dynamic! We don't know who is, or is not active anymore
int localPrefix = __popc(__activemask() & __lanemask_lt());
int totalWarpActive = __popc(__activemask());
buffer[head + localPrefix] = generateValueThisThread();
if(localPrefix == 0){
head += totalWarpActive; // Move the head forward, much like a "push" operation in single-thread land
// Only one thread should move the head
}
__syncthreads(); // Thread barrier, make sure everyone is waiting on activeThread#0 before continuing.
}
------------
As such, you can dynamically load-balance between GPU threads (!!!) from a shared stack with minimal overheads.
If you want to extend this larger than one 32-wide CUDA-warp, you'll need to use __shared __ memory to share the prefix with the rest of the block.
It is a bad idea (too much overhead) to extend this much larger than a block, as there's no quick way to communicate outside of your block. Still though, having chunks of up to 1024 threads synchronized through a shared data-structure that only has nanoseconds of overhead is a nifty trick.
-----------
EDIT: Oh right, and this concept is now replicated very, very quickly in the dedicated __ballot_sync(...) function (which compiles down to just a few assembly instructions).
Playing with the "Exec-mask" is a hugely efficient way to synchronously, and dynamically gather information across your warp. So lots of little tricks have been built around this.