llvm / llvm/llvm-project

Clang scoped atomics should accept vector types

Open
#213,237 2 comments 0 reactions 0 assignees View on GitHub
clang
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

Open the contributing 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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.