AdaptiveCpp / AdaptiveCpp/AdaptiveCpp

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

Abierto
#1,466 3 comentarios 0 reacciones 0 asignados Ver en GitHub
bug
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

Abrir la guía de contribución

Evaluación

Este issue todavía no se ha evaluado.

Recibe los nuevos issues en tu correo

Un resumen breve de issues de GitHub para principiantes.