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 6c1b63573292f37f5b94d0990cf464f4ec855950..3d790d49450dba7b8147796ba70f1a5052e27553 100644 (file)
--- a/src/gromacs/gpu_utils/gmxsycl.h
+++ b/src/gromacs/gpu_utils/gmxsycl.h
@@ -1,7 +1,7 @@
 /*
  * This file is part of the GROMACS molecular simulation package.
  *
- * Copyright (c) 2020, by the GROMACS development team, led by
+ * Copyright (c) 2020,2021, by the GROMACS development team, led by
  * Mark Abraham, David van der Spoel, Berk Hess, and Erik Lindahl,
  * and including many others, as listed in the AUTHORS file in the
  * top-level source directory and at http://www.gromacs.org.
@@ -36,8 +36,15 @@
  * \brief
  * Wraps the complexity of including SYCL in GROMACS.
  *
- * SYCL headers use symbol DIM as a template parameter, which gets broken by macro DIM defined
+ * The __SYCL_COMPILER_VERSION macro is used to identify Intel DPCPP compiler.
+ * See https://github.com/intel/llvm/pull/2998 for better proposal.
+ *
+ * Intel SYCL headers use symbol DIM as a template parameter, which gets broken by macro DIM defined
  * in gromacs/math/vectypes.h. Here, we include the SYCL header while temporary undefining this macro.
+ * See https://github.com/intel/llvm/issues/2981.
+ *
+ * Different compilers, at the time of writing, have different names for some of the proposed features
+ * of the SYCL2020 standard. For uniformity, they are all aliased in our custom sycl_2020 namespace.
  *
  * \inlibraryapi
  */
@@ -45,6 +52,30 @@
 #ifndef GMX_GPU_UTILS_GMXSYCL_H
 #define GMX_GPU_UTILS_GMXSYCL_H
 
+#include "config.h"
+
+// 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 "-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
@@ -58,4 +89,61 @@
 #    include <CL/sycl.hpp>
 #endif
 
+#if GMX_SYCL_HIPSYCL
+#    pragma clang diagnostic pop
+#endif
+
+/* Exposing Intel-specific extensions in a manner compatible with SYCL2020 provisional spec.
+ * 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.
+ * */
+namespace sycl_2020
+{
+namespace detail
+{
+#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;
+#    endif
+#elif GMX_SYCL_HIPSYCL
+namespace origin = cl::sycl;
+#else
+#    error "Unsupported version of SYCL compiler"
+#endif
+} // namespace detail
+
+using detail::origin::memory_order;
+using detail::origin::memory_scope;
+using detail::origin::plus;
+using detail::origin::sub_group;
+
+#if GMX_SYCL_DPCPP
+using detail::origin::atomic_ref;
+template<typename... Args>
+bool group_any_of(Args&&... args)
+{
+    return detail::origin::any_of(std::forward<Args>(args)...);
+}
+template<typename... Args>
+auto group_reduce(Args&&... args) -> decltype(detail::origin::reduce(std::forward<Args>(args)...))
+{
+    return detail::origin::reduce(std::forward<Args>(args)...);
+}
+#elif GMX_SYCL_HIPSYCL
+using detail::origin::atomic_ref;
+using detail::origin::group_any_of;
+using detail::origin::group_reduce;
+#else
+#    error "Unsupported SYCL compiler"
+#endif
+
+} // namespace sycl_2020
+
 #endif