GPU: make the Metal atomics operate on plain counters

[ktf]

GPUAtomic() expanded to metal::atomic, but the counters it guards are plain
integers in structures that the host allocates and transfers, and the code reads
them directly outside the atomic operations. That is exactly the CUDA and HIP
situation, where GPUAtomic() is a no-op and atomicAdd() takes a plain pointer,
so Metal now does the same.

MSL's atomic_*_explicit do want metal::atomic, which has the same size and
alignment, so the operand is cast at the one point where the operation is
issued. That cast has to be made in the address space the pointer is really in,
since casting a threadgroup pointer to a device atomic is a hard error, hence
the two overloads: threadgroup stays threadgroup for the *Shared entry points,
anything else resolves to device.

[davidrohr]

The atomic part inside the #idef METAL is of course fine.
You added GPUsharedref in the clusterizer. I assume this was only needed for OpenCL1, which we don't use anymore, so in that sense you are lucky, that we still have the GPUsharedref in most places. But you might need to fix it up in more places, let's see. Only point is: I would separate it into 2 commits, for atomics and sharedref.

[ktf]

Notice that the shared ref is actually a requirement for the atomics to work. It's there to avoid that some threadgroup pointer is silently then compiled against a device one.

With the threadgroup qualifier visible, casting to a device atomic is a hard error ("converts between mismatching address spaces"), but a generic pointer both misses the threadgroup overload and launders that cast past the compiler, so the counter would get a device atomic with no diagnostic.

[davidrohr]

Yes, my point was that it is still 2 fixes: 1. you need to add the sharedref, since it is simply missing, and the other backends do not care since they do not need it (after OpenCL1 was removed).
And then the METAL atomic change is a totally different topic IMHO.

[ktf]

ok, I will split it.

[ktf]

Ok, splitted.

[ktf]

Split done, fullCI works. All changes are no-op when running on linux. Merging.