Mostly untested attempt for #1643. This moves the bounds check outside of the atomic expression, but I'm unsure if it is correct to just immediately go to self.put_access_chain(expr_handle, policy, context)?; within or if there can be a more complicated expression than just direct member or array access that wouldn't work in this case.

Also not sure about the is_scoped logic. I figured it should be safe to always wrap in () but for some reason that's discouraged.

Looks like the CI failure is because I added the atomic tests which aren't implemented for spv. I guess I can remove those before submission, or maybe make a separate group of msl tests.

@jimblandy please review this
@glalonde sorry about the delay! Please ping us if this gets lost again.

@glalonde The PR title says "Mitigation", but are you aware of any cases where this doesn't address the problem?

I'm not aware of situations it doesn't handle, but I wouldn't be surprised if they exist.

For example, I'm not sure if there are more complicated access patterns than just array access or struct access ( in the tests) which might break.

I just tried this locally as a more convoluted access:

struct AtomicArray {
   a: array<atomic<u32>, 10>;
};

struct Globals {
...
    e: array<AtomicArray, 10>;
};

fn fetch_add_atomic_array2(i: i32, j: i32) -> u32 {
   return atomicAdd(&globals.e[i].a[j], 1u);
}

metal::uint fetch_add_atomic_array2_(
    int i_14,
    int j_2,
    device Globals& globals,
    constant _mslBufferSizes& _buffer_sizes
) {
    metal::uint _e8 = metal::uint(j_2) < 10 && metal::uint(i_14) < 10 ? metal::atomic_fetch_add_explicit(&globals.e.inner[i_14].a.inner[j_2], 1u, metal::memory_order_relaxed) : DefaultConstructible();
    return _e8;
}

And that looks like it would work...

Thanks for the review, ptal

I sorted out the style points - this is a good fix and I want to get it in right away.

Thank you very much!

awesome, thanks! looking forward to the next release : )