src/gromacs/gpu_utils/gmxsycl.h

   1 /*
   2  * This file is part of the GROMACS molecular simulation package.
   3  *
   4  * Copyright (c) 2020,2021, by the GROMACS development team, led by
   5  * Mark Abraham, David van der Spoel, Berk Hess, and Erik Lindahl,
   6  * and including many others, as listed in the AUTHORS file in the
   7  * top-level source directory and at http://www.gromacs.org.
   8  *
   9  * GROMACS is free software; you can redistribute it and/or
  10  * modify it under the terms of the GNU Lesser General Public License
  11  * as published by the Free Software Foundation; either version 2.1
  12  * of the License, or (at your option) any later version.
  13  *
  14  * GROMACS is distributed in the hope that it will be useful,
  15  * but WITHOUT ANY WARRANTY; without even the implied warranty of
  16  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
  17  * Lesser General Public License for more details.
  18  *
  19  * You should have received a copy of the GNU Lesser General Public
  20  * License along with GROMACS; if not, see
  21  * http://www.gnu.org/licenses, or write to the Free Software Foundation,
  22  * Inc., 51 Franklin Street, Fifth Floor, Boston, MA  02110-1301  USA.
  23  *
  24  * If you want to redistribute modifications to GROMACS, please
  25  * consider that scientific software is very special. Version
  26  * control is crucial - bugs must be traceable. We will be happy to
  27  * consider code for inclusion in the official distribution, but
  28  * derived work must not be called official GROMACS. Details are found
  29  * in the README & COPYING files - if they are missing, get the
  30  * official version at http://www.gromacs.org.
  31  *
  32  * To help us fund GROMACS development, we humbly ask that you cite
  33  * the research papers on the package. Check out http://www.gromacs.org.
  34  */
  35 /*! \libinternal \file
  36  * \brief
  37  * Wraps the complexity of including SYCL in GROMACS.
  38  *
  39  * The __SYCL_COMPILER_VERSION macro is used to identify Intel DPCPP compiler.
  40  * See https://github.com/intel/llvm/pull/2998 for better proposal.
  41  *
  42  * Intel SYCL headers use symbol DIM as a template parameter, which gets broken by macro DIM defined
  43  * in gromacs/math/vectypes.h. Here, we include the SYCL header while temporary undefining this macro.
  44  * See https://github.com/intel/llvm/issues/2981.
  45  *
  46  * Different compilers, at the time of writing, have different names for some of the proposed features
  47  * of the SYCL2020 standard. For uniformity, they are all aliased in our custom sycl_2020 namespace.
  48  *
  49  * \inlibraryapi
  50  */
  51
  52 #ifndef GMX_GPU_UTILS_GMXSYCL_H
  53 #define GMX_GPU_UTILS_GMXSYCL_H
  54
  55 #include "config.h"
  56
  57 /* Some versions of Intel ICPX compiler (at least 2021.1.1 and 2021.1.2) fail to unroll a loop
  58  * in sycl::accessor::__init, and emit -Wpass-failed=transform-warning. This is a useful
  59  * warning, but mostly noise right now. Probably related to using shared memory accessors.
  60  * The unroll directive was introduced in https://github.com/intel/llvm/pull/2449. */
  61 #if GMX_SYCL_DPCPP
  62 #    include <CL/sycl/version.hpp>
  63 #    define DISABLE_UNROLL_WARNINGS \
  64         ((__SYCL_COMPILER_VERSION >= 20201113) && (__SYCL_COMPILER_VERSION <= 20201214))
  65 #else
  66 #    define DISABLE_UNROLL_WARNINGS 0
  67 #endif
  68
  69 #if DISABLE_UNROLL_WARNINGS
  70 #    pragma clang diagnostic push
  71 #    pragma clang diagnostic ignored "-Wpass-failed"
  72 #endif
  73
  74 // For hipSYCL, we need to activate floating-point atomics
  75 #if GMX_SYCL_HIPSYCL
  76 #    define HIPSYCL_EXT_FP_ATOMICS
  77 #    pragma clang diagnostic push
  78 #    pragma clang diagnostic ignored "-Wunused-variable"
  79 #    pragma clang diagnostic ignored "-Wunused-parameter"
  80 #    pragma clang diagnostic ignored "-Wmissing-noreturn"
  81 #    pragma clang diagnostic ignored "-Wshadow-field"
  82 #    pragma clang diagnostic ignored "-Wctad-maybe-unsupported"
  83 #    pragma clang diagnostic ignored "-Wdeprecated-copy-dtor"
  84 #    pragma clang diagnostic ignored "-Winconsistent-missing-destructor-override"
  85 #    pragma clang diagnostic ignored "-Wunused-template"
  86 #    pragma clang diagnostic ignored "-Wsign-compare"
  87 #    pragma clang diagnostic ignored "-Wundefined-reinterpret-cast"
  88 #    pragma clang diagnostic ignored "-Wdeprecated-copy"
  89 #    pragma clang diagnostic ignored "-Wnewline-eof"
  90 #    pragma clang diagnostic ignored "-Wextra-semi"
  91 #    pragma clang diagnostic ignored "-Wsuggest-override"
  92 #    pragma clang diagnostic ignored "-Wsuggest-destructor-override"
  93 #    pragma clang diagnostic ignored "-Wgcc-compat"
  94 #endif
  95
  96
  97 #ifdef DIM
  98 #    if DIM != 3
  99 #        error "The workaround here assumes we use DIM=3."
 100 #    else
 101 #        undef DIM
 102 #        include <CL/sycl.hpp>
 103 #        define DIM 3
 104 #    endif
 105 #else
 106 #    include <CL/sycl.hpp>
 107 #endif
 108
 109 #if DISABLE_UNROLL_WARNINGS
 110 #    pragma clang diagnostic pop
 111 #endif
 112
 113 #if GMX_SYCL_HIPSYCL
 114 #    pragma clang diagnostic pop
 115 #endif
 116
 117 #undef DISABLE_UNROLL_WARNINGS
 118
 119 /* Exposing Intel-specific extensions in a manner compatible with SYCL2020 provisional spec.
 120  * Despite ICPX (up to 2021.1.2 at the least) having SYCL_LANGUAGE_VERSION=202001,
 121  * some parts of the spec are still in custom sycl::ONEAPI namespace (sycl::intel in beta versions),
 122  * and some functions have different names. To make things easier to upgrade
 123  * in the future, this thin layer is added.
 124  * */
 125 namespace sycl_2020
 126 {
 127 namespace detail
 128 {
 129 #if GMX_SYCL_DPCPP
 130 // Confirmed to work for 2021.1-beta10 (20201005), 2021.1.1 (20201113), 2021.1.2 (20201214).
 131 namespace origin = cl::sycl::ONEAPI;
 132 #elif GMX_SYCL_HIPSYCL
 133 namespace origin = cl::sycl;
 134 #else
 135 #    error "Unsupported version of SYCL compiler"
 136 #endif
 137 } // namespace detail
 138
 139 using detail::origin::memory_order;
 140 using detail::origin::memory_scope;
 141 using detail::origin::plus;
 142 using detail::origin::sub_group;
 143
 144 #if GMX_SYCL_DPCPP
 145 using detail::origin::atomic_ref;
 146 template<typename... Args>
 147 bool group_any_of(Args&&... args)
 148 {
 149     return detail::origin::any_of(std::forward<Args>(args)...);
 150 }
 151 template<typename... Args>
 152 auto group_reduce(Args&&... args) -> decltype(detail::origin::reduce(std::forward<Args>(args)...))
 153 {
 154     return detail::origin::reduce(std::forward<Args>(args)...);
 155 }
 156 #elif GMX_SYCL_HIPSYCL
 157 // No atomic_ref in hipSYCL yet (2021-02-22)
 158 using detail::origin::group_any_of;
 159 using detail::origin::group_reduce;
 160 #else
 161 #    error "Unsupported SYCL compiler"
 162 #endif
 163
 164 } // namespace sycl_2020
 165
 166 #endif