Move pmalloc(..)/pfree(...) to separate source files in CUDA/OpenCL/SYCL

Having pmalloc(..) and pfree(..) functions in unrelated files in CUDA,
OpenCL and SYCL makes it hard to use these functions in platform-agnostic
code. This creates common header and moves these functions to separate
source files.

diff --git a/src/gromacs/ewald/pme_gpu_internal.cpp b/src/gromacs/ewald/pme_gpu_internal.cpp
index 5f4a4fd28986214931ceb68c33acfe16989e0459..8eea806ac1934177f4f10364b09017354983efe6 100644 (file)
--- a/src/gromacs/ewald/pme_gpu_internal.cpp
+++ b/src/gromacs/ewald/pme_gpu_internal.cpp
@@ -60,6 +60,10 @@
 #include "gromacs/gpu_utils/device_context.h"
 #include "gromacs/gpu_utils/device_stream.h"
 #include "gromacs/gpu_utils/gpu_utils.h"
+#include "gromacs/gpu_utils/pmalloc.h"
+#if GMX_GPU_SYCL
+#    include "gromacs/gpu_utils/syclutils.h"
+#endif
 #include "gromacs/hardware/device_information.h"
 #include "gromacs/math/invertmatrix.h"
 #include "gromacs/math/units.h"
@@ -69,19 +73,12 @@
 #include "gromacs/utility/gmxassert.h"
 #include "gromacs/utility/logger.h"
 #include "gromacs/utility/stringutil.h"
+#include "gromacs/ewald/pme.h"
 
 #if GMX_GPU_CUDA
-#    include "gromacs/gpu_utils/pmalloc_cuda.h"
-
 #    include "pme.cuh"
-#elif GMX_GPU_OPENCL
-#    include "gromacs/gpu_utils/gmxopencl.h"
-#elif GMX_GPU_SYCL
-#    include "gromacs/gpu_utils/syclutils.h"
 #endif
 
-#include "gromacs/ewald/pme.h"
-
 #include "pme_gpu_3dfft.h"
 #include "pme_gpu_calculate_splines.h"
 #include "pme_gpu_constants.h"

diff --git a/src/gromacs/gpu_utils/CMakeLists.txt b/src/gromacs/gpu_utils/CMakeLists.txt
index 3b59b9279e95467d73fe03b732b7dc1339473ec1..7791042ae4c64f09132657423a280ee1b2137d4b 100644 (file)
--- a/src/gromacs/gpu_utils/CMakeLists.txt
+++ b/src/gromacs/gpu_utils/CMakeLists.txt
@@ -49,6 +49,7 @@ if(GMX_GPU_OPENCL)
     gmx_add_libgromacs_sources(
         device_context_ocl.cpp
         device_stream_ocl.cpp
+        pmalloc_ocl.cpp
         ocl_compiler.cpp
         ocl_caching.cpp
         oclutils.cpp
@@ -59,7 +60,7 @@ elseif(GMX_GPU_CUDA)
         device_stream.cu
         gpu_utils.cu
         pinning.cu
-        pmalloc_cuda.cu
+        pmalloc.cu
         )
     _gmx_add_files_to_property(CUDA_SOURCES
         device_stream_manager.cpp
@@ -69,7 +70,7 @@ elseif(GMX_GPU_SYCL)
         devicebuffer_sycl.cpp
         device_context_sycl.cpp
         device_stream_sycl.cpp
-        syclutils.cpp
+        pmalloc_sycl.cpp
         )
     _gmx_add_files_to_property(SYCL_SOURCES
         devicebuffer_sycl.cpp
@@ -77,7 +78,7 @@ elseif(GMX_GPU_SYCL)
         device_context_sycl.cpp
         device_stream_manager.cpp
         device_stream_sycl.cpp
-        syclutils.cpp
+        pmalloc_sycl.cpp
         )
 else()
     gmx_add_libgromacs_sources(

diff --git a/src/gromacs/gpu_utils/gpu_utils.cu b/src/gromacs/gpu_utils/gpu_utils.cu
index 96611b6f91b9cc5567fe15bc899f279a2ff083ac..b35fcabd4ab2d357378c33a65eb79a97e2692cca 100644 (file)
--- a/src/gromacs/gpu_utils/gpu_utils.cu
+++ b/src/gromacs/gpu_utils/gpu_utils.cu
@@ -52,7 +52,6 @@
 #include "gromacs/gpu_utils/cudautils.cuh"
 #include "gromacs/gpu_utils/device_context.h"
 #include "gromacs/gpu_utils/device_stream.h"
-#include "gromacs/gpu_utils/pmalloc_cuda.h"
 #include "gromacs/hardware/device_information.h"
 #include "gromacs/hardware/device_management.h"
 #include "gromacs/utility/basedefinitions.h"

diff --git a/src/gromacs/gpu_utils/oclutils.cpp b/src/gromacs/gpu_utils/oclutils.cpp
index 726e4f2cff47fdfd65afc75b5b80dacc20823700..d5ee96fd55b5d16e3d2b774ba71e010bd79e491e 100644 (file)
--- a/src/gromacs/gpu_utils/oclutils.cpp
+++ b/src/gromacs/gpu_utils/oclutils.cpp
@@ -2,7 +2,7 @@
  * This file is part of the GROMACS molecular simulation package.
  *
  * Copyright (c) 2014,2015,2016,2017,2018 by the GROMACS development team.
- * Copyright (c) 2019,2020, by the GROMACS development team, led by
+ * Copyright (c) 2019,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.
@@ -50,45 +50,7 @@
 #include <string>
 
 #include "gromacs/gpu_utils/gpu_utils.h"
-#include "gromacs/utility/fatalerror.h"
-#include "gromacs/utility/smalloc.h"
 
-/*! \brief \brief Allocates nbytes of host memory. Use ocl_free to free memory allocated with this function.
- *
- *  \todo
- *  This function should allocate page-locked memory to help reduce D2H and H2D
- *  transfer times, similar with pmalloc from pmalloc_cuda.cu.
- *
- * \param[in,out]    h_ptr   Pointer where to store the address of the newly allocated buffer.
- * \param[in]        nbytes  Size in bytes of the buffer to be allocated.
- */
-void pmalloc(void** h_ptr, size_t nbytes)
-{
-    /* Need a temporary type whose size is 1 byte, so that the
-     * implementation of snew_aligned can cope without issuing
-     * warnings. */
-    char** temporary = reinterpret_cast<char**>(h_ptr);
-
-    /* 16-byte alignment is required by the neighbour-searching code,
-     * because it uses four-wide SIMD for bounding-box calculation.
-     * However, when we organize using page-locked memory for
-     * device-host transfers, it will probably need to be aligned to a
-     * 4kb page, like CUDA does. */
-    snew_aligned(*temporary, nbytes, 16);
-}
-
-/*! \brief Frees memory allocated with pmalloc.
- *
- * \param[in]    h_ptr   Buffer allocated with pmalloc that needs to be freed.
- */
-void pfree(void* h_ptr)
-{
-
-    if (h_ptr)
-    {
-        sfree_aligned(h_ptr);
-    }
-}
 
 /*! \brief Convert error code to diagnostic string */
 std::string ocl_get_error_string(cl_int error)

diff --git a/src/gromacs/gpu_utils/oclutils.h b/src/gromacs/gpu_utils/oclutils.h
index c40816716139d9fabd51255e1d13b22f7ccf6bdd..5f881c33c4fc7921693e11da06c7049c043767ef 100644 (file)
--- a/src/gromacs/gpu_utils/oclutils.h
+++ b/src/gromacs/gpu_utils/oclutils.h
@@ -2,7 +2,7 @@
  * This file is part of the GROMACS molecular simulation package.
  *
  * Copyright (c) 2014,2015,2016,2017,2018 by the GROMACS development team.
- * Copyright (c) 2019,2020, by the GROMACS development team, led by
+ * Copyright (c) 2019,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.
@@ -71,12 +71,6 @@ struct gmx_device_runtime_data_t
     cl_program program;
 };
 
-/*! \brief Allocate host memory in malloc style */
-void pmalloc(void** h_ptr, size_t nbytes);
-
-/*! \brief Free host memory in malloc style */
-void pfree(void* h_ptr);
-
 /*! \brief Convert error code to diagnostic string */
 std::string ocl_get_error_string(cl_int error);
 

diff --git a/src/gromacs/gpu_utils/pmalloc_cuda.cu b/src/gromacs/gpu_utils/pmalloc.cu
similarity index 78%
rename from src/gromacs/gpu_utils/pmalloc_cuda.cu
rename to src/gromacs/gpu_utils/pmalloc.cu
index 2d5e1220529a245011656c69e6c66d40b9ab3b5c..3a8f1058fe992c19a1822f90df4c6ae9461688cc 100644 (file)
--- a/src/gromacs/gpu_utils/pmalloc_cuda.cu
+++ b/src/gromacs/gpu_utils/pmalloc.cu
@@ -1,7 +1,7 @@
 /*
  * This file is part of the GROMACS molecular simulation package.
  *
- * Copyright (c) 2012,2014,2015,2018,2019,2020, by the GROMACS development team, led by
+ * Copyright (c) 2012,2014,2015,2018,2019,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.
@@ -40,7 +40,7 @@
 
 #include "gmxpre.h"
 
-#include "pmalloc_cuda.h"
+#include "pmalloc.h"
 
 #include <stdlib.h>
 
@@ -70,29 +70,6 @@ void pmalloc(void** h_ptr, size_t nbytes)
     CU_RET_ERR(stat, strbuf);
 }
 
-/*! Allocates nbytes of page-locked memory with write-combining.
- *  This memory should always be freed using pfree (or with the page-locked
- *  free functions provied by the CUDA library).
- */
-void pmalloc_wc(void** h_ptr, size_t nbytes)
-{
-    cudaError_t stat;
-    char        strbuf[STRLEN];
-    int         flag = cudaHostAllocDefault | cudaHostAllocWriteCombined;
-
-    if (nbytes == 0)
-    {
-        *h_ptr = nullptr;
-        return;
-    }
-
-    gmx::ensureNoPendingDeviceError("Could not allocate page-locked memory with write-combining.");
-
-    stat = cudaMallocHost(h_ptr, nbytes, flag);
-    sprintf(strbuf, "cudaMallocHost of size %d bytes failed", (int)nbytes);
-    CU_RET_ERR(stat, strbuf);
-}
-
 /*! Frees page locked memory allocated with pmalloc.
  *  This function can safely be called also with a pointer to a page-locked
  *  memory allocated directly with CUDA API calls.

diff --git a/src/gromacs/gpu_utils/pmalloc_cuda.h b/src/gromacs/gpu_utils/pmalloc.h
similarity index 85%
rename from src/gromacs/gpu_utils/pmalloc_cuda.h
rename to src/gromacs/gpu_utils/pmalloc.h
index d9ab9e088d21526db95dd5064d68a1f1321df884..d6b2707556b21accf5600396f59b3fb06b890382 100644 (file)
--- a/src/gromacs/gpu_utils/pmalloc_cuda.h
+++ b/src/gromacs/gpu_utils/pmalloc.h
@@ -2,7 +2,7 @@
  * This file is part of the GROMACS molecular simulation package.
  *
  * Copyright (c) 2012,2013,2014,2015,2018 by the GROMACS development team.
- * Copyright (c) 2019,2020, by the GROMACS development team, led by
+ * Copyright (c) 2019,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.
@@ -34,13 +34,13 @@
  * the research papers on the package. Check out http://www.gromacs.org.
  */
 /*! \libinternal \file
- *  \brief Declare functions for host-side memory handling when using CUDA devices.
+ *  \brief Declare functions for host-side memory handling.
  *
  *  \author Szilard Pall <pall.szilard@gmail.com>
  *  \inlibraryapi
  */
-#ifndef GMX_GPU_UTILS_PMALLOC_CUDA_H
-#define GMX_GPU_UTILS_PMALLOC_CUDA_H
+#ifndef GMX_GPU_UTILS_PMALLOC_H
+#define GMX_GPU_UTILS_PMALLOC_H
 
 #include <stdlib.h>
 
@@ -51,9 +51,6 @@
 /*! \brief Allocates nbytes of page-locked memory. */
 void pmalloc(void** h_ptr, size_t nbytes);
 
-/*! \brief Allocates nbytes of page-locked memory with write-combining. */
-void pmalloc_wc(void** h_ptr, size_t nbytes);
-
 /*! \brief Frees page locked memory allocated with pmalloc. */
 void pfree(void* h_ptr);
 

diff --git a/src/gromacs/gpu_utils/pmalloc_ocl.cpp b/src/gromacs/gpu_utils/pmalloc_ocl.cpp
new file mode 100644 (file)
index 0000000..9ac9d77
--- /dev/null
+++ b/src/gromacs/gpu_utils/pmalloc_ocl.cpp
@@ -0,0 +1,82 @@
+/*
+ * This file is part of the GROMACS molecular simulation package.
+ *
+ * Copyright (c) 2012,2014,2015,2018,2019,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.
+ *
+ * GROMACS is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public License
+ * as published by the Free Software Foundation; either version 2.1
+ * of the License, or (at your option) any later version.
+ *
+ * GROMACS is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with GROMACS; if not, see
+ * http://www.gnu.org/licenses, or write to the Free Software Foundation,
+ * Inc., 51 Franklin Street, Fifth Floor, Boston, MA  02110-1301  USA.
+ *
+ * If you want to redistribute modifications to GROMACS, please
+ * consider that scientific software is very special. Version
+ * control is crucial - bugs must be traceable. We will be happy to
+ * consider code for inclusion in the official distribution, but
+ * derived work must not be called official GROMACS. Details are found
+ * in the README & COPYING files - if they are missing, get the
+ * official version at http://www.gromacs.org.
+ *
+ * To help us fund GROMACS development, we humbly ask that you cite
+ * the research papers on the package. Check out http://www.gromacs.org.
+ */
+/*! \internal \file
+ *  \brief Define functions for host-side memory handling when using OpenCL devices.
+ *
+ *  \author Anca Hamuraru <anca@streamcomputing.eu>
+ */
+
+#include "gmxpre.h"
+
+#include "pmalloc.h"
+
+#include "gromacs/utility/smalloc.h"
+
+/*! \brief \brief Allocates nbytes of host memory. Use pfree to free memory allocated with this function.
+ *
+ *  \todo
+ *  This function should allocate page-locked memory to help reduce D2H and H2D
+ *  transfer times, similar with pmalloc from pmalloc.cu.
+ *
+ * \param[in,out]    h_ptr   Pointer where to store the address of the newly allocated buffer.
+ * \param[in]        nbytes  Size in bytes of the buffer to be allocated.
+ */
+void pmalloc(void** h_ptr, size_t nbytes)
+{
+    /* Need a temporary type whose size is 1 byte, so that the
+     * implementation of snew_aligned can cope without issuing
+     * warnings. */
+    char** temporary = reinterpret_cast<char**>(h_ptr);
+
+    /* 16-byte alignment is required by the neighbour-searching code,
+     * because it uses four-wide SIMD for bounding-box calculation.
+     * However, when we organize using page-locked memory for
+     * device-host transfers, it will probably need to be aligned to a
+     * 4kb page, like CUDA does. */
+    snew_aligned(*temporary, nbytes, 16);
+}
+
+/*! \brief Frees memory allocated with pmalloc.
+ *
+ * \param[in]    h_ptr   Buffer allocated with pmalloc that needs to be freed.
+ */
+void pfree(void* h_ptr)
+{
+
+    if (h_ptr)
+    {
+        sfree_aligned(h_ptr);
+    }
+}

diff --git a/src/gromacs/gpu_utils/syclutils.cpp b/src/gromacs/gpu_utils/pmalloc_sycl.cpp
similarity index 96%
rename from src/gromacs/gpu_utils/syclutils.cpp
rename to src/gromacs/gpu_utils/pmalloc_sycl.cpp
index 0f660c04f5563ca51c105884abb2f068e589c9e6..839409b2452376287a06312a97bb617ab1b47f3f 100644 (file)
--- a/src/gromacs/gpu_utils/syclutils.cpp
+++ b/src/gromacs/gpu_utils/pmalloc_sycl.cpp
@@ -1,7 +1,7 @@
 /*
  * This file is part of the GROMACS molecular simulation package.
  *
- * Copyright (c) 2020,2021, by the GROMACS development team, led by
+ * Copyright (c) 2012,2014,2015,2018,2019,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.
@@ -39,7 +39,7 @@
  */
 #include "gmxpre.h"
 
-#include "syclutils.h"
+#include "pmalloc.h"
 
 #include "gromacs/utility/smalloc.h"
 

diff --git a/src/gromacs/gpu_utils/syclutils.h b/src/gromacs/gpu_utils/syclutils.h
index 1318a2a9b4d4888ef8cf50ec53e930ca21518192..ffa4be2e838779506b133a4f49e47719496c4bf7 100644 (file)
--- a/src/gromacs/gpu_utils/syclutils.h
+++ b/src/gromacs/gpu_utils/syclutils.h
@@ -67,12 +67,6 @@ struct gmx_device_runtime_data_t
 
 #ifndef DOXYGEN
 
-/*! \brief Allocate host memory in malloc style */
-void pmalloc(void** h_ptr, size_t nbytes);
-
-/*! \brief Free host memory in malloc style */
-void pfree(void* h_ptr);
-
 /* To properly mark function as [[noreturn]], we must do it everywhere it is declared, which
  * will pollute common headers.*/
 #    pragma clang diagnostic push

diff --git a/src/gromacs/gpu_utils/tests/pinnedmemorychecker.cpp b/src/gromacs/gpu_utils/tests/pinnedmemorychecker.cpp
index b4493e961e460bbe5da26aef9190a53abcb20a0a..699dff0613e97c1f9475a0283e4c2863b397324d 100644 (file)
--- a/src/gromacs/gpu_utils/tests/pinnedmemorychecker.cpp
+++ b/src/gromacs/gpu_utils/tests/pinnedmemorychecker.cpp
@@ -1,7 +1,7 @@
 /*
  * This file is part of the GROMACS molecular simulation package.
  *
- * Copyright (c) 2017,2018,2019,2020, by the GROMACS development team, led by
+ * Copyright (c) 2017,2018,2019,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.
@@ -50,7 +50,7 @@
 
 #    include "gromacs/gpu_utils/gpu_utils.h"
 #    include "gromacs/gpu_utils/hostallocator.h"
-#    include "gromacs/gpu_utils/pmalloc_cuda.h"
+#    include "gromacs/gpu_utils/pmalloc.h"
 #    include "gromacs/utility/real.h"
 #    include "gromacs/utility/smalloc.h"
 

diff --git a/src/gromacs/nbnxm/cuda/nbnxm_cuda_data_mgmt.cu b/src/gromacs/nbnxm/cuda/nbnxm_cuda_data_mgmt.cu
index 692f4076a6faa58b06802b44c884812012d6e146..ac77badc413858441700703719b89e807c7ad936 100644 (file)
--- a/src/gromacs/nbnxm/cuda/nbnxm_cuda_data_mgmt.cu
+++ b/src/gromacs/nbnxm/cuda/nbnxm_cuda_data_mgmt.cu
@@ -55,7 +55,7 @@
 #include "gromacs/gpu_utils/device_stream_manager.h"
 #include "gromacs/gpu_utils/gpu_utils.h"
 #include "gromacs/gpu_utils/gpueventsynchronizer.cuh"
-#include "gromacs/gpu_utils/pmalloc_cuda.h"
+#include "gromacs/gpu_utils/pmalloc.h"
 #include "gromacs/hardware/device_information.h"
 #include "gromacs/hardware/device_management.h"
 #include "gromacs/math/vectypes.h"

diff --git a/src/gromacs/nbnxm/opencl/nbnxm_ocl_data_mgmt.cpp b/src/gromacs/nbnxm/opencl/nbnxm_ocl_data_mgmt.cpp
index b823973a6812fe2707598370e1ab03c6e2329d5a..b927bd319645ebc687e961892251c3634fd71ba4 100644 (file)
--- a/src/gromacs/nbnxm/opencl/nbnxm_ocl_data_mgmt.cpp
+++ b/src/gromacs/nbnxm/opencl/nbnxm_ocl_data_mgmt.cpp
@@ -53,7 +53,7 @@
 #include <cmath>
 
 #include "gromacs/gpu_utils/device_stream_manager.h"
-#include "gromacs/gpu_utils/oclutils.h"
+#include "gromacs/gpu_utils/pmalloc.h"
 #include "gromacs/hardware/device_information.h"
 #include "gromacs/hardware/device_management.h"
 #include "gromacs/math/vectypes.h"

diff --git a/src/gromacs/nbnxm/sycl/nbnxm_sycl_data_mgmt.cpp b/src/gromacs/nbnxm/sycl/nbnxm_sycl_data_mgmt.cpp
index b7ad72a6a0f8f622bedac3d74c01d2648d8b804f..e5a1a4e34411f46b9d1b93071bf56e6882168fdf 100644 (file)
--- a/src/gromacs/nbnxm/sycl/nbnxm_sycl_data_mgmt.cpp
+++ b/src/gromacs/nbnxm/sycl/nbnxm_sycl_data_mgmt.cpp
@@ -42,6 +42,7 @@
 #include "gmxpre.h"
 
 #include "gromacs/gpu_utils/device_stream_manager.h"
+#include "gromacs/gpu_utils/pmalloc.h"
 #include "gromacs/hardware/device_information.h"
 #include "gromacs/mdtypes/interaction_const.h"
 #include "gromacs/nbnxm/atomdata.h"