AdaptiveCpp / AdaptiveCpp/AdaptiveCpp

Using `sycl::atomic_ref<...>{}.store()` defaults to `__ATOMIC_SEQ_CST` leading to a compiler error.

Open
#1,466 3 comments 0 reactions 0 assignees View on GitHub
bug
Dominant language
C++
Stars
1.9k
Forks
228
Avg merge
4d 1h
Merged PRs (30d)
25

Description

**Bug summary**
Calling `store()` on `sycl::atomic_ref` unexpectedly defaults to `__ATOMIC_SEQ_CST`, causing the following error:

```
fatal error: error in backend: Cannot select: 0x7ee560: ch = AtomicStore<(store seq_cst (s32) into %ir.30)> 0x7eec48:1, 0x7eeaa8, 0x7eec48
0x7eeaa8: i64,ch = CopyFromReg 0x6252198, Register:i64 %0
0x7ee220: i64 = Register %0
0x7eec48: i32,ch = load<(dereferenceable load (s32) from %ir.11)> 0x6252198, FrameIndex:i64<6>, undef:i64
0x8aa1f0: i64 = FrameIndex<6>
0x8a9b70: i64 = undef
In function: _ZNK7hipsycl4sycl10atomic_refIiLNS0_12memory_orderE0ELNS0_12memory_scopeE3ELNS0_6access13address_spaceE4EE5storeEiS2_S3_
```

**To Reproduce**
Compile the following code:
```cpp
sycl::queue q{sycl::gpu_selector{}};
auto g = sycl::malloc_device(1, q);
q.submit([&](sycl::handler& cgh) {
sycl::local_accessor l(1, cgh);
cgh.parallel_for(sycl::nd_range<1>(1, 1), [&](sycl::nd_item<1>) {
sycl::atomic_ref{g[0]}.store(1, sycl::memory_order::relaxed);
});
});
```
With the following CMake configuration:
```cmake
cmake_minimum_required(VERSION 2.25)
cmake_policy(SET CMP0058 NEW)

project(sycl_test)

set(CMAKE_CXX_STANDARD 17)
set(CMAKE_C_STANDARD_REQUIRED TRUE)
set(CMAKE_CXX_STANDARD_REQUIRED TRUE)

set(CMAKE_EXPORT_COMPILE_COMMANDS ON)
set(CMAKE_INSTALL_PREFIX ${CMAKE_CURRENT_SOURCE_DIR}/install)

add_compile_options(-fdiagnostics-color=always)

find_package(adaptivecpp CONFIG REQUIRED)

add_executable(sycl_test sycl_test.cpp)

add_sycl_to_target(TARGET sycl_test SOURCES sycl_test.cpp)
```

And this command:
`cmake -DCMAKE_PREFIX_PATH="/path/to/AdaptiveCpp" -DACPP_TARGETS=cuda:sm_75 -DCMAKE_C_COMPILER=/usr/lib/llvm-14/bin/clang -DCMAKE_CXX_COMPILER=/usr/lib/llvm-14/bin/clang++ -G Ninja ..`

The same error occurs when compiling the following CUDA code:
```cuda
__global__ void kernel() {
__shared__ int s[1];
__atomic_store_n(s, 1, __ATOMIC_SEQ_CST);
}
```

However, compiling the CUDA code with `__ATOMIC_RELAXED` works:

```cuda
__global__ void kernel() {
__shared__ int s[1];
__atomic_store_n(s, 1, __ATOMIC_RELAXED);
}
```

**Describe your setup**
* Current AdaptiveCpp version: branch = develop, [commit](https://github.com/AdaptiveCpp/AdaptiveCpp/commit/31b7c427ef0a4a066de4771ea5b625f98e451c32)
* Clang version:
```
Ubuntu clang version 14.0.0-1ubuntu1.1
Target: x86_64-pc-linux-gnu
Thread model: posix
InstalledDir: /usr/lib/llvm-14/bin
```

**Additional context**
Using LLVM 17 (clang 17.0.6) does not resolve the issue.

Contributor guide

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.