Clang scoped atomics should accept vector types
- Dominant language
- LLVM
- Stars
- 40.5k
- Forks
- 18.7k
- PR merge metrics
- PR metrics pending
Description
Now that atomicrmw properly supports vector operations, the clang atomic builtins should accept them:
```c
// Compile: clang -x hip --offload-arch=gfx942 -nogpuinc -nogpulib \
// --cuda-device-only -O1 -S -emit-llvm scoped_atomic_vector_reject.hip
#define __device__ __attribute__((device))
typedef _Float16 __attribute__((ext_vector_type(2))) half2;
typedef __bf16 __attribute__((ext_vector_type(2))) bfloat2;
// half2 via the scoped-atomic builtin:
__device__ half2 scoped_half2(half2 *p, half2 v) {
[[clang::atomic(remote_memory, no_fine_grained_memory)]] {
return __scoped_atomic_fetch_add(p, v, __ATOMIC_RELAXED, __MEMORY_SCOPE_DEVICE);
}
}
// bfloat2 via the scoped-atomic builtin:
__device__ bfloat2 scoped_bfloat2(bfloat2 *p, bfloat2 v) {
[[clang::atomic(remote_memory, no_fine_grained_memory)]] {
return __scoped_atomic_fetch_add(p, v, __ATOMIC_RELAXED, __MEMORY_SCOPE_DEVICE);
}
}
// half2 via the GNU __atomic builtin:
__device__ half2 gnu_half2(half2 *p, half2 v) {
return __atomic_fetch_add(p, v, __ATOMIC_RELAXED);
}
// bfloat2 via the C11 form with an _Atomic vector:
__device__ bfloat2 c11_bfloat2(_Atomic(bfloat2) *p, bfloat2 v) {
return __c11_atomic_fetch_add(p, v, __ATOMIC_RELAXED);
}
```
These all error today:
```
error: address argument to atomic operation must be a pointer to integer,
pointer or supported floating point type ('half2 *' (aka
'_Float16 __attribute__((ext_vector_type(2))) *') invalid)
return __scoped_atomic_fetch_add(p, v, __ATOMIC_RELAXED, __MEMORY_SCOPE_DEVICE);
^ ~
error: address argument to atomic operation must be a pointer to integer,
pointer or supported floating point type ('bfloat2 *' (aka
'__bf16 __attribute__((ext_vector_type(2))) *') invalid)
return __scoped_atomic_fetch_add(p, v, __ATOMIC_RELAXED, __MEMORY_SCOPE_DEVICE);
^ ~
error: address argument to atomic operation must be a pointer to integer,
pointer or supported floating point type ('half2 *' invalid)
return __atomic_fetch_add(p, v, __ATOMIC_RELAXED);
^ ~
error: address argument to atomic operation must be a pointer to atomic integer,
pointer or supported floating point type ('_Atomic(bfloat2) *' invalid)
return __c11_atomic_fetch_add(p, v, __ATOMIC_RELAXED);
```
Contributor guide
Research direction
Compile scoped_atomic_vector_reject.hip with the command in the issue to reproduce the diagnostics, then inspect Clang's atomic builtin validation alongside the mentioned atomicrmw vector support. Add regression coverage for the half2 and bfloat2 scoped, GNU, and C11 cases, and verify those examples compile successfully.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- c
- Domain
- compilers
- Issue type
- Feature
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Quiet
- Clarity
- Mostly clear
- Newbie friendliness
- 48/100