AdaptiveCpp / AdaptiveCpp/AdaptiveCpp
Using `sycl::atomic_ref<...>{}.store()` defaults to `__ATOMIC_SEQ_CST` leading to a compiler error.
- Lenguaje dominante
- C++
- Estrellas
- 1.9k
- Forks
- 228
- Merge medio
- 4 d 1 h
- PR fusionados (30 d)
- 25
Descripción
**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.
Guía de contribución
Evaluación
Este issue todavía no se ha evaluado.