Fix warnings with Intel 2021.4

[alexxy/gromacs.git] / src / gromacs / gpu_utils / gmxsycl.h

diff --git a/src/gromacs/gpu_utils/gmxsycl.h b/src/gromacs/gpu_utils/gmxsycl.h
index 8eb5cabc9fda44fda285e0bbe0b9c79d9d19bdae..3d790d49450dba7b8147796ba70f1a5052e27553 100644 (file)
--- a/src/gromacs/gpu_utils/gmxsycl.h
+++ b/src/gromacs/gpu_utils/gmxsycl.h
@@ -52,23 +52,31 @@
 #ifndef GMX_GPU_UTILS_GMXSYCL_H
 #define GMX_GPU_UTILS_GMXSYCL_H
 
-/* Some versions of Intel ICPX compiler (at least 2021.1.1 and 2021.1.2) fail to unroll a loop
- * in sycl::accessor::__init, and emit -Wpass-failed=transform-warning. This is a useful
- * warning, but mostly noise right now. Probably related to using shared memory accessors.
- * The unroll directive was introduced in https://github.com/intel/llvm/pull/2449. */
-#if defined(__INTEL_LLVM_COMPILER)
-#    include <CL/sycl/version.hpp>
-#    define DISABLE_UNROLL_WARNINGS \
-        ((__SYCL_COMPILER_VERSION >= 20201113) && (__SYCL_COMPILER_VERSION <= 20201214))
-#else
-#    define DISABLE_UNROLL_WARNINGS 0
-#endif
+#include "config.h"
 
-#if DISABLE_UNROLL_WARNINGS
+// For hipSYCL, we need to activate floating-point atomics
+#if GMX_SYCL_HIPSYCL
+#    define HIPSYCL_EXT_FP_ATOMICS
 #    pragma clang diagnostic push
-#    pragma clang diagnostic ignored "-Wpass-failed"
+#    pragma clang diagnostic ignored "-Wunused-variable"
+#    pragma clang diagnostic ignored "-Wunused-parameter"
+#    pragma clang diagnostic ignored "-Wmissing-noreturn"
+#    pragma clang diagnostic ignored "-Wshadow-field"
+#    pragma clang diagnostic ignored "-Wctad-maybe-unsupported"
+#    pragma clang diagnostic ignored "-Wdeprecated-copy-dtor"
+#    pragma clang diagnostic ignored "-Winconsistent-missing-destructor-override"
+#    pragma clang diagnostic ignored "-Wunused-template"
+#    pragma clang diagnostic ignored "-Wsign-compare"
+#    pragma clang diagnostic ignored "-Wundefined-reinterpret-cast"
+#    pragma clang diagnostic ignored "-Wdeprecated-copy"
+#    pragma clang diagnostic ignored "-Wnewline-eof"
+#    pragma clang diagnostic ignored "-Wextra-semi"
+#    pragma clang diagnostic ignored "-Wsuggest-override"
+#    pragma clang diagnostic ignored "-Wsuggest-destructor-override"
+#    pragma clang diagnostic ignored "-Wgcc-compat"
 #endif
 
+
 #ifdef DIM
 #    if DIM != 3
 #        error "The workaround here assumes we use DIM=3."
@@ -81,15 +89,13 @@
 #    include <CL/sycl.hpp>
 #endif
 
-#if DISABLE_UNROLL_WARNINGS
+#if GMX_SYCL_HIPSYCL
 #    pragma clang diagnostic pop
 #endif
 
-#undef DISABLE_UNROLL_WARNINGS
-
 /* Exposing Intel-specific extensions in a manner compatible with SYCL2020 provisional spec.
- * Despite ICPX (up to 2021.1.2 at the least) having SYCL_LANGUAGE_VERSION=202001,
- * some parts of the spec are still in custom sycl::ONEAPI namespace (sycl::intel in beta versions),
+ * Despite ICPX (up to 2021.3.0 at the least) having SYCL_LANGUAGE_VERSION=202001,
+ * some parts of the spec are still in custom sycl::ONEAPI namespace (sycl::ext::oneapi in beta versions),
  * and some functions have different names. To make things easier to upgrade
  * in the future, this thin layer is added.
  * */
@@ -97,10 +103,16 @@ namespace sycl_2020
 {
 namespace detail
 {
-#if defined(__SYCL_COMPILER_VERSION) // Intel DPCPP compiler
-// Confirmed to work for 2021.1-beta10 (20201005), 2021.1.1 (20201113), 2021.1.2 (20201214).
+#if GMX_SYCL_DPCPP
+// Confirmed to work for 2021.1-beta10 (20201005) to 2021.3.0 (20210619).
+// Deprecated in favor of sycl::ext::oneapi on 20210717 in https://github.com/intel/llvm/commit/d703f578.
+// Removed on 20210927 with https://github.com/intel/llvm/pull/4488
+#    if __clang_major__ >= 14
+namespace origin = sycl::ext::oneapi;
+#    else
 namespace origin = cl::sycl::ONEAPI;
-#elif defined(__HIPSYCL__)
+#    endif
+#elif GMX_SYCL_HIPSYCL
 namespace origin = cl::sycl;
 #else
 #    error "Unsupported version of SYCL compiler"
@@ -112,7 +124,7 @@ using detail::origin::memory_scope;
 using detail::origin::plus;
 using detail::origin::sub_group;
 
-#if defined(__SYCL_COMPILER_VERSION) // Intel DPCPP compiler
+#if GMX_SYCL_DPCPP
 using detail::origin::atomic_ref;
 template<typename... Args>
 bool group_any_of(Args&&... args)
@@ -124,8 +136,8 @@ auto group_reduce(Args&&... args) -> decltype(detail::origin::reduce(std::forwar
 {
     return detail::origin::reduce(std::forward<Args>(args)...);
 }
-#elif defined(__HIPSYCL__)
-// No atomic_ref in hipSYCL yet (2021-01-29)
+#elif GMX_SYCL_HIPSYCL
+using detail::origin::atomic_ref;
 using detail::origin::group_any_of;
 using detail::origin::group_reduce;
 #else