From 21b6ec1e5d9c998e8670e7286759f2a10cfd0cd7 Mon Sep 17 00:00:00 2001 From: Rafal Bielski Date: Tue, 22 Apr 2025 13:15:18 +0100 Subject: [PATCH 1/6] [DFT] Support AdaptiveCpp in cuFFT backend --- src/dft/backends/cufft/backward.cpp | 6 +++--- src/dft/backends/cufft/commit.cpp | 10 ++++++---- src/dft/backends/cufft/execute_helper.hpp | 8 +++++++- src/dft/backends/cufft/forward.cpp | 6 +++--- src/dft/execute_helper_generic.hpp | 4 +++- tests/unit_tests/dft/source/descriptor_tests.cpp | 5 ++++- 6 files changed, 26 insertions(+), 13 deletions(-) diff --git a/src/dft/backends/cufft/backward.cpp b/src/dft/backends/cufft/backward.cpp index 475f1ea49..1ee2d2ac4 100644 --- a/src/dft/backends/cufft/backward.cpp +++ b/src/dft/backends/cufft/backward.cpp @@ -76,7 +76,7 @@ ONEMATH_EXPORT void compute_backward(descriptor_type& desc, auto stream = detail::setup_stream(func_name, ih, plan); auto inout_native = reinterpret_cast*>( - ih.get_native_mem(inout_acc)); + ih.get_native_mem(inout_acc)); detail::cufft_execute>( func_name, stream, plan, reinterpret_cast(inout_native + offsets[0]), reinterpret_cast(inout_native + offsets[1])); @@ -123,11 +123,11 @@ ONEMATH_EXPORT void compute_backward(descriptor_type& desc, auto in_native = reinterpret_cast( reinterpret_cast*>( - ih.get_native_mem(in_acc)) + + ih.get_native_mem(in_acc)) + offsets[0]); auto out_native = reinterpret_cast( reinterpret_cast*>( - ih.get_native_mem(out_acc)) + + ih.get_native_mem(out_acc)) + offsets[1]); detail::cufft_execute>( func_name, stream, plan, in_native, out_native); diff --git a/src/dft/backends/cufft/commit.cpp b/src/dft/backends/cufft/commit.cpp index b6d2164ff..91cd17971 100644 --- a/src/dft/backends/cufft/commit.cpp +++ b/src/dft/backends/cufft/commit.cpp @@ -34,6 +34,8 @@ #include "oneapi/math/dft/detail/cufft/onemath_dft_cufft.hpp" #include "oneapi/math/dft/types.hpp" +#include "execute_helper.hpp" +#include "../../execute_helper_generic.hpp" #include "../stride_helper.hpp" #include @@ -84,7 +86,7 @@ class cufft_commit final : public dft::detail::commit_impl { if (fix_context) { // cufftDestroy changes the context so change it back. CUdevice interopDevice = - sycl::get_native(this->get_queue().get_device()); + sycl::get_native(this->get_queue().get_device()); CUcontext interopContext; if (cuDevicePrimaryCtxRetain(&interopContext, interopDevice) != CUDA_SUCCESS) { throw math::exception("dft/backends/cufft", __FUNCTION__, @@ -353,8 +355,8 @@ class cufft_commit final : public dft::detail::commit_impl { .submit([&](sycl::handler& cgh) { auto workspace_acc = buffer_workspace.template get_access(cgh); - cgh.host_task([=](sycl::interop_handle ih) { - auto stream = ih.get_native_queue(); + dft::detail::fft_enqueue_task(cgh, [=](sycl::interop_handle ih) { + auto stream = ih.get_native_queue(); auto result = cufftSetStream(plan, stream); if (result != CUFFT_SUCCESS) { throw oneapi::math::exception( @@ -362,7 +364,7 @@ class cufft_commit final : public dft::detail::commit_impl { "cufftSetStream returned " + std::to_string(result)); } auto workspace_native = reinterpret_cast( - ih.get_native_mem(workspace_acc)); + ih.get_native_mem(workspace_acc)); cufftSetWorkArea(plan, workspace_native); }); }) diff --git a/src/dft/backends/cufft/execute_helper.hpp b/src/dft/backends/cufft/execute_helper.hpp index da485fea2..1912ad128 100644 --- a/src/dft/backends/cufft/execute_helper.hpp +++ b/src/dft/backends/cufft/execute_helper.hpp @@ -36,6 +36,12 @@ namespace oneapi::math::dft::cufft::detail { +#ifdef __ADAPTIVECPP__ +constexpr auto sycl_cuda_backend{sycl::backend::cuda}; +#else // DPC++ +constexpr auto sycl_cuda_backend{sycl::backend::ext_oneapi_cuda}; +#endif + template inline dft::detail::commit_impl* checked_get_commit( dft::detail::descriptor& desc) { @@ -142,7 +148,7 @@ void cufft_execute(const std::string& func, CUstream stream, cufftHandle plan, v } inline CUstream setup_stream(const std::string& func, sycl::interop_handle ih, cufftHandle plan) { - auto stream = ih.get_native_queue(); + auto stream = ih.get_native_queue(); auto result = cufftSetStream(plan, stream); if (result != CUFFT_SUCCESS) { throw oneapi::math::exception("dft/backends/cufft", func, diff --git a/src/dft/backends/cufft/forward.cpp b/src/dft/backends/cufft/forward.cpp index 6b2867b5f..5350c0bbb 100644 --- a/src/dft/backends/cufft/forward.cpp +++ b/src/dft/backends/cufft/forward.cpp @@ -79,7 +79,7 @@ ONEMATH_EXPORT void compute_forward(descriptor_type& desc, auto stream = detail::setup_stream(func_name, ih, plan); auto inout_native = reinterpret_cast*>( - ih.get_native_mem(inout_acc)); + ih.get_native_mem(inout_acc)); detail::cufft_execute>( func_name, stream, plan, reinterpret_cast(inout_native + offsets[0]), reinterpret_cast(inout_native + offsets[1])); @@ -126,11 +126,11 @@ ONEMATH_EXPORT void compute_forward(descriptor_type& desc, auto in_native = reinterpret_cast( reinterpret_cast*>( - ih.get_native_mem(in_acc)) + + ih.get_native_mem(in_acc)) + offsets[0]); auto out_native = reinterpret_cast( reinterpret_cast*>( - ih.get_native_mem(out_acc)) + + ih.get_native_mem(out_acc)) + offsets[1]); detail::cufft_execute>( func_name, stream, plan, in_native, out_native); diff --git a/src/dft/execute_helper_generic.hpp b/src/dft/execute_helper_generic.hpp index 4f5ae6727..76c91f83d 100644 --- a/src/dft/execute_helper_generic.hpp +++ b/src/dft/execute_helper_generic.hpp @@ -39,7 +39,9 @@ namespace oneapi::math::dft::detail { */ template static inline void fft_enqueue_task(HandlerT&& cgh, FnT&& f) { -#ifdef SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND +#if defined(__ADAPTIVECPP__) + cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) { +#elif defined(SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND) cgh.ext_codeplay_enqueue_native_command([=](sycl::interop_handle ih) { #else cgh.host_task([=](sycl::interop_handle ih) { diff --git a/tests/unit_tests/dft/source/descriptor_tests.cpp b/tests/unit_tests/dft/source/descriptor_tests.cpp index a4290e553..0ccc659cd 100644 --- a/tests/unit_tests/dft/source/descriptor_tests.cpp +++ b/tests/unit_tests/dft/source/descriptor_tests.cpp @@ -571,7 +571,9 @@ inline void recommit_values(sycl::queue& sycl_queue) { } template -inline void change_queue_causes_wait(sycl::queue& busy_queue) { +inline void change_queue_causes_wait([[maybe_unused]] sycl::queue& busy_queue) { + // Skip this test in AdaptiveCpp, which doesn't support host_task + #ifndef __ADAPTIVECPP__ // create a queue with work on it, and then show that work is waited on when the descriptor // is committed to a new queue. // its possible to have a false positive result, but a false negative should not be possible. @@ -616,6 +618,7 @@ inline void change_queue_causes_wait(sycl::queue& busy_queue) { // busy queue task has now completed. auto after_status = e.template get_info(); ASSERT_EQ(after_status, sycl::info::event_command_status::complete); + #endif } template From 692b0ed966e20954d423c2759693b38739536dd4 Mon Sep 17 00:00:00 2001 From: Rafal Bielski Date: Wed, 23 Apr 2025 16:16:36 +0100 Subject: [PATCH 2/6] [DFT] Support AdaptiveCpp in rocFFT backend --- src/dft/backends/rocfft/commit.cpp | 6 ++++-- src/dft/backends/rocfft/execute_helper.hpp | 10 ++++++++-- 2 files changed, 12 insertions(+), 4 deletions(-) diff --git a/src/dft/backends/rocfft/commit.cpp b/src/dft/backends/rocfft/commit.cpp index 4c5d51d2f..47f12f336 100644 --- a/src/dft/backends/rocfft/commit.cpp +++ b/src/dft/backends/rocfft/commit.cpp @@ -34,6 +34,8 @@ #include "oneapi/math/dft/detail/rocfft/onemath_dft_rocfft.hpp" #include "oneapi/math/dft/types.hpp" +#include "execute_helper.hpp" +#include "../../execute_helper_generic.hpp" #include "../stride_helper.hpp" #include "rocfft_handle.hpp" @@ -557,9 +559,9 @@ class rocfft_commit final : public dft::detail::commit_impl { this->get_queue().submit([&](sycl::handler& cgh) { auto workspace_acc = buffer_workspace.template get_access(cgh); - cgh.host_task([=](sycl::interop_handle ih) { + dft::detail::fft_enqueue_task(cgh, [=](sycl::interop_handle ih) { auto workspace_native = reinterpret_cast( - ih.get_native_mem(workspace_acc)); + ih.get_native_mem(workspace_acc)); set_workspace_impl(handle, workspace_native, workspace_bytes, "set_workspace"); }); }); diff --git a/src/dft/backends/rocfft/execute_helper.hpp b/src/dft/backends/rocfft/execute_helper.hpp index c1ee6302b..c928d9e2d 100644 --- a/src/dft/backends/rocfft/execute_helper.hpp +++ b/src/dft/backends/rocfft/execute_helper.hpp @@ -36,6 +36,12 @@ namespace oneapi::math::dft::rocfft::detail { +#ifdef __ADAPTIVECPP__ +constexpr auto sycl_hip_backend{sycl::backend::hip}; +#else // DPC++ +constexpr auto sycl_hip_backend{sycl::backend::ext_oneapi_hip}; +#endif + template inline dft::detail::commit_impl* checked_get_commit( dft::detail::descriptor& desc) { @@ -60,12 +66,12 @@ inline auto expect_config(DescT& desc, const char* message) { template inline void* native_mem(sycl::interop_handle& ih, Acc& buf) { - return ih.get_native_mem(buf); + return ih.get_native_mem(buf); } inline hipStream_t setup_stream(const std::string& func, sycl::interop_handle& ih, rocfft_execution_info info) { - auto stream = ih.get_native_queue(); + auto stream = ih.get_native_queue(); auto result = rocfft_execution_info_set_stream(info, stream); if (result != rocfft_status_success) { throw oneapi::math::exception( From 49396ea324b10ef3f987e1508b802e3e30cc3e7c Mon Sep 17 00:00:00 2001 From: Rafal Bielski Date: Wed, 23 Apr 2025 17:33:09 +0100 Subject: [PATCH 3/6] [DFT] Update docs with cuFFT and rocFFT support for AdaptiveCpp --- README.md | 4 ++-- docs/building_the_project_with_adaptivecpp.rst | 18 ++++++++++++------ 2 files changed, 14 insertions(+), 8 deletions(-) diff --git a/README.md b/README.md index 7c2766c1e..eef846925 100644 --- a/README.md +++ b/README.md @@ -338,7 +338,7 @@ Supported compilers include: NVIDIA GPU NVIDIA cuFFT - Open DPC++ + Open DPC++
AdaptiveCpp Dynamic, Static @@ -349,7 +349,7 @@ Supported compilers include: AMD GPU AMD rocFFT - Open DPC++ + Open DPC++
AdaptiveCpp Dynamic, Static diff --git a/docs/building_the_project_with_adaptivecpp.rst b/docs/building_the_project_with_adaptivecpp.rst index 41e5b03f7..c4e1ccfca 100644 --- a/docs/building_the_project_with_adaptivecpp.rst +++ b/docs/building_the_project_with_adaptivecpp.rst @@ -57,7 +57,7 @@ additional guidance. The target architectures must be specified with ``HIP_TARGETS``. See the `AdaptiveCpp documentation `_. -If a backend library supports multiple domains (i.e. BLAS, RNG), it may be +If a backend library supports multiple domains (i.e. BLAS, DFT, RNG), it may be desirable to only enable selected domains. For this, the ``TARGET_DOMAINS`` variable should be set. For further details, see :ref:`_build_target_domains`. @@ -81,6 +81,9 @@ The most important supported build options are: * - ENABLE_CUBLAS_BACKEND - True, False - False + * - ENABLE_CUFFT_BACKEND + - True, False + - False * - ENABLE_CURAND_BACKEND - True, False - False @@ -93,6 +96,9 @@ The most important supported build options are: * - ENABLE_ROCBLAS_BACKEND - True, False - False + * - ENABLE_ROCFFT_BACKEND + - True, False + - False * - ENABLE_ROCRAND_BACKEND - True, False - False @@ -106,7 +112,7 @@ The most important supported build options are: - True, False - True * - TARGET_DOMAINS (list) - - blas, rng + - blas, dft, rng - All supported domains Some additional build options are given in @@ -120,8 +126,8 @@ Backends Building for CUDA ~~~~~~~~~~~~~~~~~ -The CUDA backends can be enabled with ``ENABLE_CUBLAS_BACKEND`` and -``ENABLE_CURAND_BACKEND``. +The CUDA backends can be enabled with ``ENABLE_CUBLAS_BACKEND``, +``ENABLE_CUFFT_BACKEND`` and ``ENABLE_CURAND_BACKEND``. The target architecture must be set using the ``HIPSYCL_TARGETS`` parameter. For example, to target a Nvidia A100 (Ampere architecture), set @@ -140,8 +146,8 @@ the CUDA libraries should be found automatically by CMake. Building for ROCm ~~~~~~~~~~~~~~~~~ -The ROCm backends can be enabled with ``ENABLE_ROCBLAS_BACKEND`` and -``ENABLE_ROCRAND_BACKEND``. +The ROCm backends can be enabled with ``ENABLE_ROCBLAS_BACKEND``, +``ENABLE_ROCFFT_BACKEND`` and ``ENABLE_ROCRAND_BACKEND``. The target architecture must be set using the ``HIPSYCL_TARGETS`` parameter. See the `AdaptiveCpp documentation From a902e6b3750cfed0beb78a682d797121c53ff257 Mon Sep 17 00:00:00 2001 From: Rafal Bielski Date: Thu, 24 Apr 2025 15:58:06 +0100 Subject: [PATCH 4/6] [DFT] Apply clang-format --- src/dft/backends/cufft/backward.cpp | 16 ++++++++-------- src/dft/backends/cufft/execute_helper.hpp | 4 ++-- src/dft/backends/cufft/forward.cpp | 16 ++++++++-------- src/dft/backends/rocfft/execute_helper.hpp | 4 ++-- tests/unit_tests/dft/source/descriptor_tests.cpp | 6 +++--- 5 files changed, 23 insertions(+), 23 deletions(-) diff --git a/src/dft/backends/cufft/backward.cpp b/src/dft/backends/cufft/backward.cpp index 1ee2d2ac4..41483062e 100644 --- a/src/dft/backends/cufft/backward.cpp +++ b/src/dft/backends/cufft/backward.cpp @@ -121,14 +121,14 @@ ONEMATH_EXPORT void compute_backward(descriptor_type& desc, dft::detail::fft_enqueue_task(cgh, [=](sycl::interop_handle ih) { auto stream = detail::setup_stream(func_name, ih, plan); - auto in_native = reinterpret_cast( - reinterpret_cast*>( - ih.get_native_mem(in_acc)) + - offsets[0]); - auto out_native = reinterpret_cast( - reinterpret_cast*>( - ih.get_native_mem(out_acc)) + - offsets[1]); + auto in_native = + reinterpret_cast(reinterpret_cast*>( + ih.get_native_mem(in_acc)) + + offsets[0]); + auto out_native = + reinterpret_cast(reinterpret_cast*>( + ih.get_native_mem(out_acc)) + + offsets[1]); detail::cufft_execute>( func_name, stream, plan, in_native, out_native); }); diff --git a/src/dft/backends/cufft/execute_helper.hpp b/src/dft/backends/cufft/execute_helper.hpp index 1912ad128..e5c47d4de 100644 --- a/src/dft/backends/cufft/execute_helper.hpp +++ b/src/dft/backends/cufft/execute_helper.hpp @@ -37,9 +37,9 @@ namespace oneapi::math::dft::cufft::detail { #ifdef __ADAPTIVECPP__ -constexpr auto sycl_cuda_backend{sycl::backend::cuda}; +constexpr auto sycl_cuda_backend{ sycl::backend::cuda }; #else // DPC++ -constexpr auto sycl_cuda_backend{sycl::backend::ext_oneapi_cuda}; +constexpr auto sycl_cuda_backend{ sycl::backend::ext_oneapi_cuda }; #endif template diff --git a/src/dft/backends/cufft/forward.cpp b/src/dft/backends/cufft/forward.cpp index 5350c0bbb..324c83142 100644 --- a/src/dft/backends/cufft/forward.cpp +++ b/src/dft/backends/cufft/forward.cpp @@ -124,14 +124,14 @@ ONEMATH_EXPORT void compute_forward(descriptor_type& desc, dft::detail::fft_enqueue_task(cgh, [=](sycl::interop_handle ih) { auto stream = detail::setup_stream(func_name, ih, plan); - auto in_native = reinterpret_cast( - reinterpret_cast*>( - ih.get_native_mem(in_acc)) + - offsets[0]); - auto out_native = reinterpret_cast( - reinterpret_cast*>( - ih.get_native_mem(out_acc)) + - offsets[1]); + auto in_native = + reinterpret_cast(reinterpret_cast*>( + ih.get_native_mem(in_acc)) + + offsets[0]); + auto out_native = + reinterpret_cast(reinterpret_cast*>( + ih.get_native_mem(out_acc)) + + offsets[1]); detail::cufft_execute>( func_name, stream, plan, in_native, out_native); }); diff --git a/src/dft/backends/rocfft/execute_helper.hpp b/src/dft/backends/rocfft/execute_helper.hpp index c928d9e2d..ebc1f072b 100644 --- a/src/dft/backends/rocfft/execute_helper.hpp +++ b/src/dft/backends/rocfft/execute_helper.hpp @@ -37,9 +37,9 @@ namespace oneapi::math::dft::rocfft::detail { #ifdef __ADAPTIVECPP__ -constexpr auto sycl_hip_backend{sycl::backend::hip}; +constexpr auto sycl_hip_backend{ sycl::backend::hip }; #else // DPC++ -constexpr auto sycl_hip_backend{sycl::backend::ext_oneapi_hip}; +constexpr auto sycl_hip_backend{ sycl::backend::ext_oneapi_hip }; #endif template diff --git a/tests/unit_tests/dft/source/descriptor_tests.cpp b/tests/unit_tests/dft/source/descriptor_tests.cpp index 0ccc659cd..5ca9c790a 100644 --- a/tests/unit_tests/dft/source/descriptor_tests.cpp +++ b/tests/unit_tests/dft/source/descriptor_tests.cpp @@ -572,8 +572,8 @@ inline void recommit_values(sycl::queue& sycl_queue) { template inline void change_queue_causes_wait([[maybe_unused]] sycl::queue& busy_queue) { - // Skip this test in AdaptiveCpp, which doesn't support host_task - #ifndef __ADAPTIVECPP__ +// Skip this test in AdaptiveCpp, which doesn't support host_task +#ifndef __ADAPTIVECPP__ // create a queue with work on it, and then show that work is waited on when the descriptor // is committed to a new queue. // its possible to have a false positive result, but a false negative should not be possible. @@ -618,7 +618,7 @@ inline void change_queue_causes_wait([[maybe_unused]] sycl::queue& busy_queue) { // busy queue task has now completed. auto after_status = e.template get_info(); ASSERT_EQ(after_status, sycl::info::event_command_status::complete); - #endif +#endif } template From e100870354fc9eb14cce53afc34841c6ace41cbd Mon Sep 17 00:00:00 2001 From: Rafal Bielski Date: Mon, 28 Apr 2025 19:44:14 +0100 Subject: [PATCH 5/6] [DFT] Support both hipSYCL and AdaptiveCpp macros Still support older releases with different macro naming --- src/dft/backends/cufft/execute_helper.hpp | 2 +- src/dft/backends/rocfft/execute_helper.hpp | 2 +- src/dft/execute_helper_generic.hpp | 2 ++ tests/unit_tests/dft/source/descriptor_tests.cpp | 2 +- 4 files changed, 5 insertions(+), 3 deletions(-) diff --git a/src/dft/backends/cufft/execute_helper.hpp b/src/dft/backends/cufft/execute_helper.hpp index e5c47d4de..a588c35d0 100644 --- a/src/dft/backends/cufft/execute_helper.hpp +++ b/src/dft/backends/cufft/execute_helper.hpp @@ -36,7 +36,7 @@ namespace oneapi::math::dft::cufft::detail { -#ifdef __ADAPTIVECPP__ +#if defined(__ADAPTIVECPP__) || defined (__HIPSYCL__) constexpr auto sycl_cuda_backend{ sycl::backend::cuda }; #else // DPC++ constexpr auto sycl_cuda_backend{ sycl::backend::ext_oneapi_cuda }; diff --git a/src/dft/backends/rocfft/execute_helper.hpp b/src/dft/backends/rocfft/execute_helper.hpp index ebc1f072b..8cf089ce9 100644 --- a/src/dft/backends/rocfft/execute_helper.hpp +++ b/src/dft/backends/rocfft/execute_helper.hpp @@ -36,7 +36,7 @@ namespace oneapi::math::dft::rocfft::detail { -#ifdef __ADAPTIVECPP__ +#if defined(__ADAPTIVECPP__) || defined (__HIPSYCL__) constexpr auto sycl_hip_backend{ sycl::backend::hip }; #else // DPC++ constexpr auto sycl_hip_backend{ sycl::backend::ext_oneapi_hip }; diff --git a/src/dft/execute_helper_generic.hpp b/src/dft/execute_helper_generic.hpp index 76c91f83d..a40a3f465 100644 --- a/src/dft/execute_helper_generic.hpp +++ b/src/dft/execute_helper_generic.hpp @@ -41,6 +41,8 @@ template static inline void fft_enqueue_task(HandlerT&& cgh, FnT&& f) { #if defined(__ADAPTIVECPP__) cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) { +#elif defined(__HIPSYCL__) + cgh.hipSYCL_enqueue_custom_operation([=](sycl::interop_handle ih) { #elif defined(SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND) cgh.ext_codeplay_enqueue_native_command([=](sycl::interop_handle ih) { #else diff --git a/tests/unit_tests/dft/source/descriptor_tests.cpp b/tests/unit_tests/dft/source/descriptor_tests.cpp index 5ca9c790a..3e728be43 100644 --- a/tests/unit_tests/dft/source/descriptor_tests.cpp +++ b/tests/unit_tests/dft/source/descriptor_tests.cpp @@ -573,7 +573,7 @@ inline void recommit_values(sycl::queue& sycl_queue) { template inline void change_queue_causes_wait([[maybe_unused]] sycl::queue& busy_queue) { // Skip this test in AdaptiveCpp, which doesn't support host_task -#ifndef __ADAPTIVECPP__ +#if !defined(__ADAPTIVECPP__) && !defined (__HIPSYCL__) // create a queue with work on it, and then show that work is waited on when the descriptor // is committed to a new queue. // its possible to have a false positive result, but a false negative should not be possible. From 45ab6ea5cdbe8c2bcc5516d97716526fcbb07696 Mon Sep 17 00:00:00 2001 From: Rafal Bielski Date: Mon, 28 Apr 2025 20:07:16 +0100 Subject: [PATCH 6/6] [DFT] Apply clang-format --- src/dft/backends/cufft/execute_helper.hpp | 2 +- src/dft/backends/rocfft/execute_helper.hpp | 2 +- tests/unit_tests/dft/source/descriptor_tests.cpp | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) diff --git a/src/dft/backends/cufft/execute_helper.hpp b/src/dft/backends/cufft/execute_helper.hpp index a588c35d0..3d5f6e791 100644 --- a/src/dft/backends/cufft/execute_helper.hpp +++ b/src/dft/backends/cufft/execute_helper.hpp @@ -36,7 +36,7 @@ namespace oneapi::math::dft::cufft::detail { -#if defined(__ADAPTIVECPP__) || defined (__HIPSYCL__) +#if defined(__ADAPTIVECPP__) || defined(__HIPSYCL__) constexpr auto sycl_cuda_backend{ sycl::backend::cuda }; #else // DPC++ constexpr auto sycl_cuda_backend{ sycl::backend::ext_oneapi_cuda }; diff --git a/src/dft/backends/rocfft/execute_helper.hpp b/src/dft/backends/rocfft/execute_helper.hpp index 8cf089ce9..27218d0ad 100644 --- a/src/dft/backends/rocfft/execute_helper.hpp +++ b/src/dft/backends/rocfft/execute_helper.hpp @@ -36,7 +36,7 @@ namespace oneapi::math::dft::rocfft::detail { -#if defined(__ADAPTIVECPP__) || defined (__HIPSYCL__) +#if defined(__ADAPTIVECPP__) || defined(__HIPSYCL__) constexpr auto sycl_hip_backend{ sycl::backend::hip }; #else // DPC++ constexpr auto sycl_hip_backend{ sycl::backend::ext_oneapi_hip }; diff --git a/tests/unit_tests/dft/source/descriptor_tests.cpp b/tests/unit_tests/dft/source/descriptor_tests.cpp index 3e728be43..a825a1539 100644 --- a/tests/unit_tests/dft/source/descriptor_tests.cpp +++ b/tests/unit_tests/dft/source/descriptor_tests.cpp @@ -573,7 +573,7 @@ inline void recommit_values(sycl::queue& sycl_queue) { template inline void change_queue_causes_wait([[maybe_unused]] sycl::queue& busy_queue) { // Skip this test in AdaptiveCpp, which doesn't support host_task -#if !defined(__ADAPTIVECPP__) && !defined (__HIPSYCL__) +#if !defined(__ADAPTIVECPP__) && !defined(__HIPSYCL__) // create a queue with work on it, and then show that work is waited on when the descriptor // is committed to a new queue. // its possible to have a false positive result, but a false negative should not be possible.