Mitigation for MSL atomic bounds check.

[glalonde]

Will CI trigger automatically?

[glalonde]

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.

[glalonde]

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.

[kvark]

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

[jimblandy]

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

[glalonde]

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...

[jimblandy]

Yeah, I think as long as we do the same things that an ordinary load would do, we should be fine. I'll finish the review in a few minutes.

…

On Mon, Jan 31, 2022, 6:48 PM glalonde ***@***.***> wrote: 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... — Reply to this email directly, view it on GitHub <#1703 (comment)>, or unsubscribe <https://github.com/notifications/unsubscribe-auth/AAFXNKCHOWISLCSNZ4IXVR3UY5CZBANCNFSM5M42YR2A> . Triage notifications on the go with GitHub Mobile for iOS <https://apps.apple.com/app/apple-store/id1477376905?ct=notification-email&mt=8&pt=524675> or Android <https://play.google.com/store/apps/details?id=com.github.android&referrer=utm_campaign%3Dnotification-email%26utm_medium%3Demail%26utm_source%3Dgithub>. You are receiving this because you were mentioned.Message ID: ***@***.***>

[jimblandy]

This looks like just the right approach. I have a few small comments, and a suggestion on how to fix the tests.

[glalonde]

Thanks for the review, ptal

[jimblandy]

This looks great. I wrote the bounds-checking code for Metal, so I humbly thank you for catching my mistake. I have one last change to request, and then I'll merge it.

[jimblandy]

Might want to add:

// for atomic types. These are separate from `bounds-check-zero.wgsl because
// SPIR-V does not yet support `ReadZeroSkipWrite` for atomics. Once it does,
// the test files could be combined.

[jimblandy]

Yes, I noticed the duplicated code as well. Naga's style is to not use closures in this way, to keep control flow simple. This is certainly subjective, but it's what we've settled on.

Instead, please declare a bool, and then have two if statements on either side of the common code.

[jimblandy]

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

Thank you very much!

[glalonde]

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