naga: Comparison-expressions on atomic types are not treated as an error

[armansito]

Description
According to the WGSL Specification, Section 8.8, comparison expressions are defined to operate on AbstractInt, AbstractFloat, bool, i32, u32, f32, or f16 (and apply component-wise on their vector equivalents). It looks like naga (0.19.2) permits an instance of atomic<u32> and a u32 to be passed as the operands of a comparison operator and injects an atomic load before the comparison while translating the output. Tint treats this as an error and rejects the program.

I believe Tint's behavior matches the specification while Naga's behavior does not. Please see the details below.

Details

The offending WGSL is from a compute program that applies comparison between two bindings (// ... denotes sections that are omitted from the rest of the program):

struct Config {
    // ...
    lines_size: u32,
    // ...
}

struct BumpAllocators {
    // ...
    lines: atomic<u32>,
    // ...
}

@group(0) @binding(0) var<uniform> config: Config;
@group(0) @binding(1) var<storage, read_write> bump: BumpAllocators;

var<workgroup> sh_previous_failed: u32;

@compute @workgroup_size(256)
fn main(@builtin(local_invocation_id) local_id: vec3<u32>) {
    // ...
    if local_id.x == 0u {
        let failed = bump.lines > config.lines_size;    // offending line
        sh_previous_failed = u32(failed);
    }
    // ...
}

Tint Behavior:

Given the above, tint outputs the following error, which I believe matches the spec:

Errors:
line 243:22 no matching overload for 'operator > (atomic<u32>, u32)'

2 candidate operators:
 • 'operator > (T  ✗ , T  ✗ ) -> bool' where:
      ✗  'T' is 'abstract-float', 'abstract-int', 'f32', 'i32', 'u32' or 'f16'
 • 'operator > (vecN<T>  ✗ , vecN<T>  ✗ ) -> vecN<bool>' where:
      ✗  'T' is 'abstract-float', 'abstract-int', 'f32', 'i32', 'u32' or 'f16'

Naga Behavior:

Naga accepts the program. Looking at the generated MSL, it looks like Naga inserts an atomic load and performs the comparison on the result of that. I haven't attempted to reproduce this on other backends:

    if (local_id.x == 0u) {
        uint _e22 = metal::atomic_load_explicit(&bump.lines, metal::memory_order_relaxed);
        uint _e25 = config.lines_size;
        bool failed = _e22 > _e25;
        sh_previous_failed = static_cast<uint>(failed);
    }

[armansito]

It looks like this may apply to numerical operators in general. For instance, I see that the following compound-assignment to an atomic (where the type of bump.failed is atomic<u32> and STAGE_FLATTEN is a u32) is treated similarly:

WGSL:

bump.failed |= STAGE_FLATTEN;

This is rejected by Tint with the following message:

line 250:25 no matching overload for 'operator |= (atomic<u32>, u32)'

4 candidate operators:
 • 'operator |= (T  ✗ , T  ✗ ) -> T' where:
      ✗  'T' is 'abstract-int', 'i32' or 'u32'
 • 'operator |= (bool  ✗ , bool  ✗ ) -> bool'
 • 'operator |= (vecN<bool>  ✗ , vecN<bool>  ✗ ) -> vecN<bool>'
 • 'operator |= (vecN<T>  ✗ , vecN<T>  ✗ ) -> vecN<T>' where:
      ✗  'T' is 'abstract-int', 'i32' or 'u32'

but Naga accepts it and translates it to the following MSL:

uint _e39 = metal::atomic_load_explicit(&bump.failed, metal::memory_order_relaxed);
metal::atomic_store_explicit(&bump.failed, _e39 | STAGE_FLATTEN, metal::memory_order_relaxed);

[ErichDonGubler]

Going to add the kind: diagnostics label here, since this is an interesting example of argument type mismatch diagnostics. 🙂