CUDA-2.4ΒΆ

All threads must execute the same __syncwarp() in convergence

Required inputs: IR, StaticSemanticAnalysis

CUDA 2.4 [collective.warp.in_convergence] All threads must execute the same __syncwarp() in convergence

In early systems all threads in the mask must execute the same __syncwarp() in convergence, and the union of all values in mask must be equal to the active mask. Otherwise, the behavior is undefined.

Scope: Device.
Audience: CUDA C++.
Category: Mandatory.
Hardware Applicability: Compute Capability 6.0 and below.
Rationale

In earlier systems the __syncwarp() behavior was more strictly defined. Since incorrect invocation can result in undefined behavior, the programmer most avoid invoking the call incorrectly.

Example 1 (Bad)
__global__ void test() {
    if(thread.idx == 0) {
      __syncwarp(~0u); // Non-compliant: call in divergent code.
    }
    else {
      __syncwarp(~1u); // Non-compliant: call in divergent code.
    }
}
Example 2 (Bad)
__global__ void test() {
    unsigned mask;
    if(thread.idx == 0) {
      mask = ~0u;
    }
    else if(thread.idx == 1) {
return;
    }
    else {
      mask = ~1u;
    }
    __syncwarp(mask); // Non-compliant: Union of mask values is not equal to the active mask.
}
Excerpt from NVIDIA CUDA C++ Guidelines for robust and safety-critical programming, Version 3.0.1, Copyright (C) 2018-2023 NVIDIA Corporation.

Possible Messages

Key

Text

Severity

Disabled

divergent_mask

Participating lanes do not agree on sync mask. Lane(s) {lanes} have mask 0b{lane_mask:032b} but union of all masks is 0b{union_mask:032b} in context(s) {contexts}

None

False

not_participating_lane_in_mask

Lane(s) {lanes} named in sync mask 0b{mask:032b} but do not participate in context(s) {contexts}

None

False

potentially_divergent_mask

Potentially participating lanes do not agree on sync mask. Lane(s) {lanes} have mask 0b{lane_mask:032b} but union of all masks is 0b{union_mask:032b} in context(s) {contexts}

None

True

potentially_not_participating_lane_in_mask

Lane(s) {lanes} named in sync mask 0b{mask:032b} but potentially do not participate in context(s) {contexts}

None

True

Options