title: In CUDA, how can I get this warp's thread mask in conditionally executed code (in order to execute e.g., __shfl_sync or <cg>.shfl?
body: <p>I'm trying to update some older CUDA code (pre CUDA 9.0), and I'm having some difficulty updating usage of warp shuffles (e.g., <code>__shfl</code>).</p>
  <p>Basically the relevant part of the kernel might be something like this:</p>
  <pre><code>int f = d[threadIdx.x];
  int warpLeader = &lt;something in [0,32)&gt;;
  
  // Point being, some threads in the warp get removed by i &lt; stop
  for(int i = k; i &lt; stop; i+=skip)
  {
     // Point being, potentially more threads don't see the shuffle below.
     if(mem[threadIdx.x + i/2] == foo)
     {
        // Pre CUDA 9.0.
        f = __shfl(f, warpLeader); 
    }
  }
  </code></pre>
  <p>Maybe that's not the best example (real code too complex to post), but the two things accomplished easily with the old intrinsics were:</p>
  <ul>
  <li>Shuffle/broadcast to whatever threads happen to be here at this time.</li>
  <li>Still get to use the warp-relative thread index.</li>
  </ul>
  <p>I'm not sure how to do the above post CUDA 9.0.</p>
  <p>This question is almost/partially answered here: <a href="https://stackoverflow.com/questions/57709125/how-can-i-synchronize-threads-within-warp-in-conditional-while-statement-in-cuda">How can I synchronize threads within warp in conditional while statement in CUDA?</a>, but I think that post has a few unresolved questions.</p>
  <p>I don't believe <code>__shfl_sync(__activemask(), ...)</code> will work. This was noted in the linked question and many other places online.</p>
  <p>The linked question says to use <code>coalesced_group</code>, but my understanding is that this type of <code>cooperative_group</code> re-ranks the threads, so if you had a <code>warpLeader</code> (on [0, 32)) in mind as above, I'm not sure there's a way to &quot;figure out&quot; its new rank in the <code>coalesced_group</code>.</p>
  <p>(Also, based on the truncated comment conversation in the linked question, it seems unclear if <code>coalesced_group</code> is just a nice wrapper for <code>__activemask()</code> or not anyway ...)</p>
  <p>It is possible to iteratively build up a mask using <code>__ballot_sync</code> as described in the linked question, but for code similar to the above, that can become pretty tedious. Is this our only way forward for CUDA &gt; 9.0?</p>