AdaptiveCpp / AdaptiveCpp/AdaptiveCpp

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

Đang mở
#1,466 3 bình luận 0 reaction 0 người được giao Xem trên GitHub
bug
Ngôn ngữ chính
C++
Star
1.9k
Fork
228
Merge trung bình
4 ngày 1 giờ
Pull request đã merge (30 ngày)
25

Mô tả

**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.

Hướng dẫn đóng góp

Mở hướng dẫn đóng góp

Đánh giá

Issue này chưa được đánh giá.

Nhận issue mới trong hộp thư của bạn

Bản tóm tắt ngắn những issue GitHub phù hợp với người mới.