Some codes in dpct/util.hpp use DPC++ features which don't support AMD GPU.
mazhaojia123 opened this issue · 4 comments
mazhaojia123 commented
Describe the bug
Some codes in dpct/util.hpp use DPC++ features which don't support AMD GPU.
The call stack is as follow.
In file included from main.dp.cpp:1:
In file included from /home/mazhaojia/pkg/dpct_workspace/c2s_install_20240321/include/dpct/dpct.hpp:12:
In file included from /home/mazhaojia/pkg/sycl_workspace/build-20240321/bin/../include/sycl/sycl.hpp:15:
/home/mazhaojia/pkg/sycl_workspace/build-20240321/bin/../include/sycl/atomic_ref.hpp:131:17: error: static assertion failed due to requirement '(sycl::memory_order)5 != sycl::memory_order::seq_cst': seq_cst memory order is not supported on AMDGPU
131 | static_assert(DefaultOrder != sycl::memory_order::seq_cst,
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/mazhaojia/pkg/sycl_workspace/build-20240321/bin/../include/sycl/atomic_ref.hpp:278:14: note: in instantiation of template class 'sycl::detail::atomic_ref_base<unsigned int, sycl::memory_order::seq_cst, sycl::memory_scope::device, sycl::access::address_space::global_space>' requested here
278 | : public atomic_ref_base<T, DefaultOrder, DefaultScope, AddressSpace> {
| ^
/home/mazhaojia/pkg/sycl_workspace/build-20240321/bin/../include/sycl/atomic_ref.hpp:714:14: note: in instantiation of template class 'sycl::detail::atomic_ref_impl<unsigned int, false, sycl::memory_order::seq_cst, sycl::memory_scope::device, sycl::access::address_space::global_space>' requested here
714 | : public detail::atomic_ref_impl<T, sizeof(T) == 8, DefaultOrder,
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_20240321/include/dpct/util.hpp:652:25: note: in instantiation of template class 'sycl::atomic_ref<unsigned int, sycl::memory_order::seq_cst, sycl::memory_scope::device, sycl::access::address_space::global_space>' requested here
652 | old_arrive = counter.fetch_add(inc);
| ^
/home/mazhaojia/pkg/sycl_workspace/build-20240321/bin/../include/sycl/atomic_ref.hpp:131:30: note: expression evaluates to '5 != 5'
131 | static_assert(DefaultOrder != sycl::memory_order::seq_cst,
| ~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
In file included from main.dp.cpp:1:
In file included from /home/mazhaojia/pkg/dpct_workspace/c2s_install_20240321/include/dpct/dpct.hpp:22:
In file included from /home/mazhaojia/pkg/dpct_workspace/c2s_install_20240321/include/dpct/image.hpp:15:
/home/mazhaojia/pkg/dpct_workspace/c2s_install_20240321/include/dpct/util.hpp:654:35: error: no member named 'load' in 'sycl::atomic_ref<unsigned int, sycl::memory_order::seq_cst, sycl::memory_scope::device, sycl::access::address_space::global_space>'
654 | while (((old_arrive ^ counter.load()) & 0x80000000) == 0)
| ~~~~~~~ ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_20240321/include/dpct/util.hpp:687:26: error: no member named 'fetch_add' in 'sycl::atomic_ref<unsigned int, sycl::memory_order::seq_cst, sycl::memory_scope::device, sycl::access::address_space::global_space>'
687 | old_arrive = counter.fetch_add(inc);
| ~~~~~~~ ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_20240321/include/dpct/util.hpp:689:35: error: no member named 'load' in 'sycl::atomic_ref<unsigned int, sycl::memory_order::seq_cst, sycl::memory_scope::device, sycl::access::address_space::global_space>'
689 | while (((old_arrive ^ counter.load()) & 0x80000000) == 0)
| ~~~~~~~ ^
4 errors generated.
make: *** [Makefile:10: main] Error 1
To reproduce
main.dp.cpp
#include <dpct/dpct.hpp>
int main() {}Compile command I use.
clang++ -fsycl -fsycl-targets=amdgcn-amd-amdhsa -Xsycl-target-backend --offload-arch=gfx908 main.dp.cpp -o main
Environment
- OS: Ubuntu 20.04
- DPCT: SYCLomatic 20240321 build.
- DPC++: Intel/llvm 20240321 build
- sycl-ls --verbose
# sycl-ls --verbose
[opencl:cpu][opencl:0] Intel(R) OpenCL, Intel(R) Xeon(R) Silver 4110 CPU @ 2.10GHz OpenCL 3.0 (Build 0) [2023.16.10.0.17_160000]
[opencl:cpu][opencl:1] Intel(R) OpenCL, Intel(R) Xeon(R) Silver 4110 CPU @ 2.10GHz OpenCL 3.0 (Build 0) [2023.16.6.0.22_223734]
[opencl:fpga][opencl:2] Intel(R) FPGA Emulation Platform for OpenCL(TM), Intel(R) FPGA Emulation Device OpenCL 1.2 [2023.16.6.0.22_223734]
[hip:gpu][hip:0] AMD HIP BACKEND, AMD Instinct MI100 gfx908:sramecc+:xnack- [HIP 50422.80]
Platforms: 4
Platform [#1]:
Version : OpenCL 3.0 LINUX
Name : Intel(R) OpenCL
Vendor : Intel(R) Corporation
Devices : 1
Device [#0]:
Type : cpu
Version : OpenCL 3.0 (Build 0)
Name : Intel(R) Xeon(R) Silver 4110 CPU @ 2.10GHz
Vendor : Intel(R) Corporation
Driver : 2023.16.10.0.17_160000
Aspects : cpu fp64 online_compiler online_linker queue_profiling usm_device_allocations usm_host_allocations usm_shared_allocations usm_system_allocations t_allocations usm_atomic_shared_allocations atomic64 ext_oneapi_srgb ext_intel_legacy_image ext_oneapi_ballot_group ext_oneapi_fixed_size_group ext_oneapi_opportunistneapi_tangle_group
info::device::sub_group_sizes: 4 8 16 32 64
Platform [#2]:
Version : OpenCL 3.0 LINUX
Name : Intel(R) OpenCL
Vendor : Intel(R) Corporation
Devices : 1
Device [#1]:
Type : cpu
Version : OpenCL 3.0 (Build 0)
Name : Intel(R) Xeon(R) Silver 4110 CPU @ 2.10GHz
Vendor : Intel(R) Corporation
Driver : 2023.16.6.0.22_223734
Aspects : cpu fp64 online_compiler online_linker queue_profiling usm_device_allocations usm_host_allocations usm_shared_allocations usm_system_allocations t_allocations usm_atomic_shared_allocations atomic64 ext_oneapi_srgb ext_intel_legacy_image ext_oneapi_ballot_group ext_oneapi_fixed_size_group ext_oneapi_opportunistneapi_tangle_group
info::device::sub_group_sizes: 4 8 16 32 64
Platform [#3]:
Version : OpenCL 1.2 Intel(R) FPGA SDK for OpenCL(TM), Version 20.3
Name : Intel(R) FPGA Emulation Platform for OpenCL(TM)
Vendor : Intel(R) Corporation
Devices : 1
Device [#2]:
Type : fpga
Version : OpenCL 1.2
Name : Intel(R) FPGA Emulation Device
Vendor : Intel(R) Corporation
Driver : 2023.16.6.0.22_223734
Aspects : accelerator fp64 online_compiler online_linker queue_profiling usm_device_allocations usm_host_allocations usm_shared_allocations usm_atomic_hostsm_atomic_shared_allocations ext_oneapi_srgb ext_oneapi_ballot_group ext_oneapi_fixed_size_group ext_oneapi_opportunistic_group ext_oneapi_tangle_group ext_intel_fpga
info::device::sub_group_sizes: 4 8 16 32 64
Platform [#4]:
Version : HIP 50422.80
Name : AMD HIP BACKEND
Vendor : AMD Corporation
Devices : 1
Device [#0]:
Type : gpu
Version : gfx908:sramecc+:xnack-
Name : AMD Instinct MI100
Vendor : AMD Corporation
Driver : HIP 50422.80
Aspects : gpu fp16 fp64 online_compiler online_linker queue_profiling usm_device_allocations usm_host_allocations usm_shared_allocations ext_intel_pci_addr_host_allocations usm_atomic_shared_allocations atomic64 ext_intel_device_info_uuid ext_oneapi_native_assert ext_intel_free_memory ext_intel_device_id ext_intel_memorxt_intel_memory_bus_width ext_intel_legacy_image ext_oneapi_graph
info::device::sub_group_sizes: 64
default_selector() : gpu, AMD HIP BACKEND, AMD Instinct MI100 gfx908:sramecc+:xnack- [HIP 50422.80]
accelerator_selector() : fpga, Intel(R) FPGA Emulation Platform for OpenCL(TM), Intel(R) FPGA Emulation Device OpenCL 1.2 [2023.16.6.0.22_223734]
cpu_selector() : cpu, Intel(R) OpenCL, Intel(R) Xeon(R) Silver 4110 CPU @ 2.10GHz OpenCL 3.0 (Build 0) [2023.16.10.0.17_160000]
gpu_selector() : gpu, AMD HIP BACKEND, AMD Instinct MI100 gfx908:sramecc+:xnack- [HIP 50422.80]
custom_selector(gpu) : gpu, AMD HIP BACKEND, AMD Instinct MI100 gfx908:sramecc+:xnack- [HIP 50422.80]
custom_selector(cpu) : cpu, Intel(R) OpenCL, Intel(R) Xeon(R) Silver 4110 CPU @ 2.10GHz OpenCL 3.0 (Build 0) [2023.16.10.0.17_160000]
custom_selector(acc) : fpga, Intel(R) FPGA Emulation Platform for OpenCL(TM), Intel(R) FPGA Emulation Device OpenCL 1.2 [2023.16.6.0.22_223734]
Additional context
No response
mazhaojia123 commented
Anyway, I commented these codes. It's seems that there is no functional impact.
mazhaojia123 commented
The solution from this PR seems to cause some other errors.
dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:61:8: error: duplicate case value 'acq_rel'
61 | case sycl::memory_order::acq_rel:
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:58:8: note: previous case defined here
58 | case sycl::memory_order::acq_rel:
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:124:8: error: duplicate case value 'acq_rel'
124 | case sycl::memory_order::acq_rel:
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:121:8: note: previous case defined here
121 | case sycl::memory_order::acq_rel:
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:187:8: error: duplicate case value 'acq_rel'
187 | case sycl::memory_order::acq_rel:
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:184:8: note: previous case defined here
184 | case sycl::memory_order::acq_rel:
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:250:8: error: duplicate case value 'acq_rel'
250 | case sycl::memory_order::acq_rel:
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:247:8: note: previous case defined here
247 | case sycl::memory_order::acq_rel:
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:313:8: error: duplicate case value 'acq_rel'
313 | case sycl::memory_order::acq_rel:
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:310:8: note: previous case defined here
310 | case sycl::memory_order::acq_rel:
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:376:8: error: duplicate case value 'acq_rel'
376 | case sycl::memory_order::acq_rel:
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:373:8: note: previous case defined here
373 | case sycl::memory_order::acq_rel:
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:439:8: error: duplicate case value 'acq_rel'
439 | case sycl::memory_order::acq_rel:
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:436:8: note: previous case defined here
436 | case sycl::memory_order::acq_rel:
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:531:8: error: duplicate case value 'acq_rel'
531 | case sycl::memory_order::acq_rel:
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:527:8: note: previous case defined here
527 | case sycl::memory_order::acq_rel:
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:585:8: error: duplicate case value 'acq_rel'
585 | case sycl::memory_order::acq_rel:
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:582:8: note: previous case defined here
582 | case sycl::memory_order::acq_rel:
| ^
9 errors generated.
make: *** [Makefile:10: main] Error 1
mazhaojia123 commented
There are still some bugs from the commit of PR1818.
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:699:45: error: no type named 'acq_rel' in 'sycl::memory_order'
699 | unsigned int, sycl::memory_order::acq_rel,
| ~~~~~~~~~~~~~~~~~~~~^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:699:23: error: template parameter missing a default argument
699 | unsigned int, sycl::memory_order::acq_rel,
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:697:45: note: previous default template argument defined here
697 | sycl::memory_scope DefaultScope = sycl::memory_scope::system,
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:699:52: error: template parameter missing a default argument
699 | unsigned int, sycl::memory_order::acq_rel,
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:697:45: note: previous default template argument defined here
697 | sycl::memory_scope DefaultScope = sycl::memory_scope::system,
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:716:27: error: use of undeclared identifier 'DefaultOrder'
716 | sycl::atomic_ref<T, DefaultOrder, DefaultScope, Space>::default_read_order;
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:715:39: error: constexpr variable 'default_read_order' must be initialized by a constant expression
715 | static constexpr sycl::memory_order default_read_order =
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:718:27: error: use of undeclared identifier 'DefaultOrder'
718 | sycl::atomic_ref<T, DefaultOrder, DefaultScope, Space>::default_write_order;
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:717:39: error: constexpr variable 'default_write_order' must be initialized by a constant expression
717 | static constexpr sycl::memory_order default_write_order =
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:721:7: error: use of undeclared identifier 'DefaultOrder'
721 | DefaultOrder;
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:735:25: error: use of undeclared identifier 'DefaultOrder'
735 | sycl::atomic_ref<T, DefaultOrder, DefaultScope, Space> atm(__d);
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:745:25: error: use of undeclared identifier 'DefaultOrder'
745 | sycl::atomic_ref<T, DefaultOrder, DefaultScope, Space> atm(
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:759:25: error: use of undeclared identifier 'DefaultOrder'
759 | sycl::atomic_ref<T, DefaultOrder, DefaultScope, Space> atm(__d);
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:775:25: error: use of undeclared identifier 'DefaultOrder'
775 | sycl::atomic_ref<T, DefaultOrder, DefaultScope, Space> atm(__d);
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:786:25: error: use of undeclared identifier 'DefaultOrder'
786 | sycl::atomic_ref<T, DefaultOrder, DefaultScope, Space> atm(__d);
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:803:25: error: use of undeclared identifier 'DefaultOrder'
803 | sycl::atomic_ref<T, DefaultOrder, DefaultScope, Space> atm(__d);
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:814:25: error: use of undeclared identifier 'DefaultOrder'
814 | sycl::atomic_ref<T, DefaultOrder, DefaultScope, Space> atm(__d);
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:827:25: error: use of undeclared identifier 'DefaultOrder'
827 | sycl::atomic_ref<T, DefaultOrder, DefaultScope, Space> atm(__d);
| ^
/home/mazhaojia/pkg/dpct_workspace/c2s_install_tom/include/dpct/atomic.hpp:840:25: error: use of undeclared identifier 'DefaultOrder'
840 | sycl::atomic_ref<T, DefaultOrder, DefaultScope, Space> atm(__d);
| ^
17 errors generated.
make: *** [Makefile:10: main] Error 1