64bit atomics in metal

Graphics & Games Metal
alexfs123 OP
Created
Replies 1
Boosts 0
Views 911
Participants 2

According to the official Metal Feature Set Tables, the Apple9 family introduces full 64-bit atomics such as atomic_compare_exchange_weak_explicit() and others. However when trying to use it the compiler throws an error "no matching function call...". Example:

kernel void test_64bit_cas(device atomic_ulong* atomic_var [[buffer(0)]]) {
    ulong expected = 0;
    ulong desired = 1;
    // This fails to compile despite Apple9 hardware supporting 64-bit buffer atomics
    atomic_compare_exchange_weak_explicit(atomic_var, &expected, desired, memory_order_relaxed, memory_order_relaxed);
}

Am i doing something wrong ?

Answered by DTS Engineer in 908459022

Thanks for the minimal kernel.

Your reading of the Metal Feature Set Tables matches mine: footnote 7 says "The full set of 64-bit atomic operations is supported on all platforms starting with Apple9." The Metal Shading Language Specification (version 4.1) is narrower:

  • Section 6.16.4.4 says the compare-and-exchange functions "support atomic types A of atomic_int, atomic_uint, atomic_bool, and atomic_float."
  • Section 6.16.4.6, "Atomic Modify Functions (64 Bits)," lists atomic_max_explicit and atomic_min_explicit for atomic_ulong in device memory, with memory_order_relaxed as the only supported order. Both return void.

So for atomic_ulong, the compare-and-exchange call has no matching function, and the compiler matches the specification. With Xcode 27.0, your kernel fails with the same error under -std=metal3.1 through -std=metal4.1. Calls to atomic_max_explicit and atomic_min_explicit on the same device atomic_ulong* compile.

I suggest filing two separate Feedback Assistant reports:

  • a report about the difference between the Metal Feature Set Tables and the Metal Shading Language Specification for 64-bit atomics
  • an enhancement request for 64-bit compare-and-exchange on atomic_ulong, describing your use case

Please post the Feedback numbers here so we can make sure they get routed to the right team.

Please note that filing an enhancement request does not guarantee Apple will add any requested functionality to Apple products in the future. However, with that said, filing an enhancement request is a good way to get your ideas in front of the folks who make decisions about that sort of thing.

If you're not familiar with filing reports, take a look at Bug Reporting: How and Why?

Replies  1
Boosts 0
Views  911
Participants  2
DTS Engineer OP
Apple
Recommended

Thanks for the minimal kernel.

Your reading of the Metal Feature Set Tables matches mine: footnote 7 says "The full set of 64-bit atomic operations is supported on all platforms starting with Apple9." The Metal Shading Language Specification (version 4.1) is narrower:

  • Section 6.16.4.4 says the compare-and-exchange functions "support atomic types A of atomic_int, atomic_uint, atomic_bool, and atomic_float."
  • Section 6.16.4.6, "Atomic Modify Functions (64 Bits)," lists atomic_max_explicit and atomic_min_explicit for atomic_ulong in device memory, with memory_order_relaxed as the only supported order. Both return void.

So for atomic_ulong, the compare-and-exchange call has no matching function, and the compiler matches the specification. With Xcode 27.0, your kernel fails with the same error under -std=metal3.1 through -std=metal4.1. Calls to atomic_max_explicit and atomic_min_explicit on the same device atomic_ulong* compile.

I suggest filing two separate Feedback Assistant reports:

  • a report about the difference between the Metal Feature Set Tables and the Metal Shading Language Specification for 64-bit atomics
  • an enhancement request for 64-bit compare-and-exchange on atomic_ulong, describing your use case

Please post the Feedback numbers here so we can make sure they get routed to the right team.

Please note that filing an enhancement request does not guarantee Apple will add any requested functionality to Apple products in the future. However, with that said, filing an enhancement request is a good way to get your ideas in front of the folks who make decisions about that sort of thing.

If you're not familiar with filing reports, take a look at Bug Reporting: How and Why?