CU-LAUN0001 · raw_barrier_over_cooperative_groups

Count raw barriers and warp intrinsics that Cooperative Groups states more safely.

This is a deterministic rule for cuda. Read its implementation.

Definition

Report a call to a raw block barrier or a masked warp intrinsic. Cooperative Groups expresses the same synchronization through a typed group object, which makes the scope explicit in the signature instead of implicit in a mask argument. A wrong mask is silent, because the threads that were left out keep running and the result is a race that reproduces only under some occupancy.

Evidence

Each finding records the call range and the intrinsic. The value is the number of raw synchronization calls.

Exceptions

A kernel that must run on a toolkit older than Cooperative Groups keeps the raw form. A performance-critical kernel may also keep a hand-written intrinsic after measurement, which is a decision worth recording rather than a finding to suppress silently.

Examples

Bad

__syncthreads();
value = __shfl_down_sync(0xffffffff, value, offset);

Good

auto block = cooperative_groups::this_thread_block();
block.sync();
auto warp = cooperative_groups::tiled_partition<32>(block);
value = warp.shfl_down(value, offset);

References

• Cites “CUDA C++ Programming Guide”, Cooperative Groups. Open reference
• Cites “The NVIDIA Technical Blog”, Cooperative Groups, flexible CUDA thread programming. Open reference
• Cites “CUDA C++ Programming Guide”, warp shuffle functions. Open reference