From 77ac5aae23cb1d6a133e64f2e8fd80ad9d9294e5 Mon Sep 17 00:00:00 2001 From: Mark Abraham Date: Fri, 30 Aug 2019 18:27:05 +0200 Subject: [PATCH] Extend GPU traits class Now GPU traits provide a non-GPU header, so that generic code can use CommandStream, CommandEvent and DeviceContext types. The header also diverges to a platform-specific version when needed upon compilation. This change allows for passing the variables of the above types in the general (non-GPU) parts of the code and can be included where the code is shared between different platforms. Renamed a Context variable to DeviceContext for greater clarity. Change-Id: If21b9dacac66ff7203948eb03de96f9473b7359a --- src/gromacs/ewald/pme_gpu_program_impl.h | 12 +--- src/gromacs/ewald/pme_gpu_types_host_impl.h | 2 +- src/gromacs/ewald/pme_only.cpp | 4 +- src/gromacs/gpu_utils/devicebuffer.cuh | 4 +- src/gromacs/gpu_utils/devicebuffer.h | 8 +-- src/gromacs/gpu_utils/devicebuffer_ocl.h | 6 +- src/gromacs/gpu_utils/gputraits.cuh | 5 +- src/gromacs/gpu_utils/gputraits.h | 67 +++++++++++++++++++ src/gromacs/gpu_utils/gputraits_ocl.h | 5 +- src/gromacs/mdrun/runner.cpp | 4 +- .../mdtypes/state_propagator_data_gpu.h | 14 ++-- .../state_propagator_data_gpu_impl.cpp | 2 +- .../mdtypes/state_propagator_data_gpu_impl.h | 16 ++--- .../state_propagator_data_gpu_impl_gpu.cpp | 22 +++--- .../nbnxm/cuda/nbnxm_cuda_data_mgmt.cu | 2 +- .../nbnxm/opencl/nbnxm_ocl_data_mgmt.cpp | 2 +- 16 files changed, 118 insertions(+), 57 deletions(-) create mode 100644 src/gromacs/gpu_utils/gputraits.h diff --git a/src/gromacs/ewald/pme_gpu_program_impl.h b/src/gromacs/ewald/pme_gpu_program_impl.h index f1c114115f..bfe4815af7 100644 --- a/src/gromacs/ewald/pme_gpu_program_impl.h +++ b/src/gromacs/ewald/pme_gpu_program_impl.h @@ -44,17 +44,9 @@ #include "config.h" +#include "gromacs/gpu_utils/gputraits.h" #include "gromacs/utility/classhelpers.h" -#if GMX_GPU == GMX_GPU_CUDA -#include "gromacs/gpu_utils/gputraits.cuh" -#elif GMX_GPU == GMX_GPU_OPENCL -#include "gromacs/gpu_utils/gputraits_ocl.h" -#elif GMX_GPU == GMX_GPU_NONE -// TODO place in gputraits_stub.h -using Context = void *; -#endif - struct gmx_device_info_t; /*! \internal @@ -85,7 +77,7 @@ struct PmeGpuProgramImpl * TODO: Later we want to be able to own the context at a higher level and not here, * but this class would still need the non-owning context handle to build the kernels. */ - Context context; + DeviceContext context; //! Conveniently all the PME kernels use the same single argument type #if GMX_GPU == GMX_GPU_CUDA diff --git a/src/gromacs/ewald/pme_gpu_types_host_impl.h b/src/gromacs/ewald/pme_gpu_types_host_impl.h index be865c8778..b113fcba4e 100644 --- a/src/gromacs/ewald/pme_gpu_types_host_impl.h +++ b/src/gromacs/ewald/pme_gpu_types_host_impl.h @@ -76,7 +76,7 @@ struct PmeGpuSpecific * but should be a constructor parameter to PmeGpu, as well as PmeGpuProgram, * managed by high-level code. */ - Context context; + DeviceContext context; /* Synchronization events */ /*! \brief Triggered after the PME Force Calculations have been completed */ diff --git a/src/gromacs/ewald/pme_only.cpp b/src/gromacs/ewald/pme_only.cpp index dc2a2ec7e8..01518afbb7 100644 --- a/src/gromacs/ewald/pme_only.cpp +++ b/src/gromacs/ewald/pme_only.cpp @@ -548,7 +548,7 @@ int gmx_pmeonly(struct gmx_pme_t *pme, //TODO the variable below should be queried from the task assignment info const bool useGpuForPme = (runMode == PmeRunMode::GPU) || (runMode == PmeRunMode::Mixed); const void *commandStream = useGpuForPme ? pme_gpu_get_device_stream(pme) : nullptr; - const void *gpuContext = useGpuForPme ? pme_gpu_get_device_context(pme) : nullptr; + const void *deviceContext = useGpuForPme ? pme_gpu_get_device_context(pme) : nullptr; const int paddingSize = pme_gpu_get_padding_size(pme); if (useGpuForPme) { @@ -557,7 +557,7 @@ int gmx_pmeonly(struct gmx_pme_t *pme, } // Unconditionally initialize the StatePropagatorDataGpu object to get more verbose message if it is used from CPU builds - auto stateGpu = std::make_unique(commandStream, gpuContext, GpuApiCallBehavior::Sync, paddingSize); + auto stateGpu = std::make_unique(commandStream, deviceContext, GpuApiCallBehavior::Sync, paddingSize); clear_nrnb(mynrnb); diff --git a/src/gromacs/gpu_utils/devicebuffer.cuh b/src/gromacs/gpu_utils/devicebuffer.cuh index ea95f021d6..7e7cfe91a7 100644 --- a/src/gromacs/gpu_utils/devicebuffer.cuh +++ b/src/gromacs/gpu_utils/devicebuffer.cuh @@ -57,12 +57,12 @@ * \tparam ValueType Raw value type of the \p buffer. * \param[in,out] buffer Pointer to the device-side buffer. * \param[in] numValues Number of values to accomodate. - * \param[in] context The buffer's dummy context - not managed explicitly in CUDA RT. + * \param[in] deviceContext The buffer's dummy device context - not managed explicitly in CUDA RT. */ template void allocateDeviceBuffer(DeviceBuffer *buffer, size_t numValues, - Context /* context */) + DeviceContext /* deviceContext */) { GMX_ASSERT(buffer, "needs a buffer pointer"); cudaError_t stat = cudaMalloc((void **)buffer, numValues * sizeof(ValueType)); diff --git a/src/gromacs/gpu_utils/devicebuffer.h b/src/gromacs/gpu_utils/devicebuffer.h index 7e0ff99e77..5c469cead9 100644 --- a/src/gromacs/gpu_utils/devicebuffer.h +++ b/src/gromacs/gpu_utils/devicebuffer.h @@ -65,7 +65,7 @@ * Allocation is buffered and therefore freeing is only needed * if the previously allocated space is not enough. * \p currentNumValues and \p currentMaxNumValues are updated. - * TODO: \p currentNumValues, \p currentMaxNumValues, \p context + * TODO: \p currentNumValues, \p currentMaxNumValues, \p deviceContext * should all be encapsulated in a host-side class together with the buffer. * * \tparam ValueType Raw value type of the \p buffer. @@ -73,14 +73,14 @@ * \param[in] numValues Number of values to accommodate. * \param[in,out] currentNumValues The pointer to the buffer's number of values. * \param[in,out] currentMaxNumValues The pointer to the buffer's capacity. - * \param[in] context The buffer's context. + * \param[in] deviceContext The buffer's device context. */ template void reallocateDeviceBuffer(DeviceBuffer *buffer, size_t numValues, int *currentNumValues, int *currentMaxNumValues, - Context context) + DeviceContext deviceContext) { GMX_ASSERT(buffer, "needs a buffer pointer"); GMX_ASSERT(currentNumValues, "needs a size pointer"); @@ -95,7 +95,7 @@ void reallocateDeviceBuffer(DeviceBuffer *buffer, } *currentMaxNumValues = over_alloc_large(numValues); - allocateDeviceBuffer(buffer, *currentMaxNumValues, context); + allocateDeviceBuffer(buffer, *currentMaxNumValues, deviceContext); } /* size could have changed without actual reallocation */ *currentNumValues = numValues; diff --git a/src/gromacs/gpu_utils/devicebuffer_ocl.h b/src/gromacs/gpu_utils/devicebuffer_ocl.h index 1c92a7aa00..723b391b13 100644 --- a/src/gromacs/gpu_utils/devicebuffer_ocl.h +++ b/src/gromacs/gpu_utils/devicebuffer_ocl.h @@ -57,17 +57,17 @@ * \tparam ValueType Raw value type of the \p buffer. * \param[in,out] buffer Pointer to the device-side buffer. * \param[in] numValues Number of values to accomodate. - * \param[in] context The buffer's context-to-be. + * \param[in] deviceContext The buffer's device context-to-be. */ template void allocateDeviceBuffer(DeviceBuffer *buffer, size_t numValues, - Context context) + DeviceContext deviceContext) { GMX_ASSERT(buffer, "needs a buffer pointer"); void *hostPtr = nullptr; cl_int clError; - *buffer = clCreateBuffer(context, CL_MEM_READ_WRITE, numValues * sizeof(ValueType), hostPtr, &clError); + *buffer = clCreateBuffer(deviceContext, CL_MEM_READ_WRITE, numValues * sizeof(ValueType), hostPtr, &clError); GMX_RELEASE_ASSERT(clError == CL_SUCCESS, "clCreateBuffer failure"); } diff --git a/src/gromacs/gpu_utils/gputraits.cuh b/src/gromacs/gpu_utils/gputraits.cuh index f3bb0437e7..559cad265f 100644 --- a/src/gromacs/gpu_utils/gputraits.cuh +++ b/src/gromacs/gpu_utils/gputraits.cuh @@ -1,7 +1,7 @@ /* * This file is part of the GROMACS molecular simulation package. * - * Copyright (c) 2018, by the GROMACS development team, led by + * Copyright (c) 2018,2019, 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,6 +40,7 @@ * \author Aleksei Iupinov * * \inlibraryapi + * \ingroup module_gpu_utils */ //! \brief GPU command stream @@ -47,7 +48,7 @@ using CommandStream = cudaStream_t; //! \brief Single GPU call timing event - meaningless in CUDA using CommandEvent = void; //! \brief Context used explicitly in OpenCL, does nothing in CUDA -using Context = void *; +using DeviceContext = void *; /*! \internal \brief * GPU kernels scheduling description. This is same in OpenCL/CUDA. diff --git a/src/gromacs/gpu_utils/gputraits.h b/src/gromacs/gpu_utils/gputraits.h new file mode 100644 index 0000000000..0a8b36517d --- /dev/null +++ b/src/gromacs/gpu_utils/gputraits.h @@ -0,0 +1,67 @@ +/* + * This file is part of the GROMACS molecular simulation package. + * + * Copyright (c) 2018,2019, 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. + */ +#ifndef GMX_GPU_UTILS_GPUTRAITS_H +#define GMX_GPU_UTILS_GPUTRAITS_H + +/*! \libinternal \file + * \brief Declares the GPU type traits for non-GPU builds + * \author Mark Abraham + * + * \inlibraryapi + * \ingroup module_gpu_utils + */ + +#include "config.h" + +#if GMX_GPU == GMX_GPU_CUDA + +#include "gromacs/gpu_utils/gputraits.cuh" + +#elif GMX_GPU == GMX_GPU_OPENCL + +#include "gromacs/gpu_utils/gputraits_ocl.h" + +#else + +//! \brief GPU command stream +using CommandStream = void *; +//! \brief Single GPU call timing event +using CommandEvent = void *; +//! \brief GPU context +using DeviceContext = void *; + +#endif // GMX_GPU + +#endif // GMX_GPU_UTILS_GPUTRAITS_H diff --git a/src/gromacs/gpu_utils/gputraits_ocl.h b/src/gromacs/gpu_utils/gputraits_ocl.h index 5287d32727..5e93698318 100644 --- a/src/gromacs/gpu_utils/gputraits_ocl.h +++ b/src/gromacs/gpu_utils/gputraits_ocl.h @@ -1,7 +1,7 @@ /* * This file is part of the GROMACS molecular simulation package. * - * Copyright (c) 2018, by the GROMACS development team, led by + * Copyright (c) 2018,2019, 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,6 +40,7 @@ * \author Aleksei Iupinov * * \inlibraryapi + * \ingroup module_gpu_utils */ #include "gromacs/gpu_utils/gmxopencl.h" @@ -49,7 +50,7 @@ using CommandStream = cl_command_queue; //! \brief Single GPU call timing event using CommandEvent = cl_event; //! \brief Context used explicitly in OpenCL -using Context = cl_context; +using DeviceContext = cl_context; /*! \internal \brief * GPU kernels scheduling description. This is same in OpenCL/CUDA. diff --git a/src/gromacs/mdrun/runner.cpp b/src/gromacs/mdrun/runner.cpp index 56b399151d..f7f24405b5 100644 --- a/src/gromacs/mdrun/runner.cpp +++ b/src/gromacs/mdrun/runner.cpp @@ -1502,7 +1502,7 @@ int Mdrunner::mdrunner() fcd->disres.nsystems != 0); const void *commandStream = ((GMX_GPU == GMX_GPU_OPENCL) && thisRankHasPmeGpuTask) ? pme_gpu_get_device_stream(fr->pmedata) : nullptr; - const void *gpuContext = ((GMX_GPU == GMX_GPU_OPENCL) && thisRankHasPmeGpuTask) ? pme_gpu_get_device_context(fr->pmedata) : nullptr; + const void *deviceContext = ((GMX_GPU == GMX_GPU_OPENCL) && thisRankHasPmeGpuTask) ? pme_gpu_get_device_context(fr->pmedata) : nullptr; const int paddingSize = pme_gpu_get_padding_size(fr->pmedata); const bool inputIsCompatibleWithModularSimulator = ModularSimulator::isInputCompatible( @@ -1517,7 +1517,7 @@ int Mdrunner::mdrunner() // We initialize GPU state even for the CPU runs so we will have a more verbose // error if someone will try accessing it from the CPU codepath gmx::StatePropagatorDataGpu stateGpu(commandStream, - gpuContext, + deviceContext, transferKind, paddingSize); fr->stateGpu = &stateGpu; diff --git a/src/gromacs/mdtypes/state_propagator_data_gpu.h b/src/gromacs/mdtypes/state_propagator_data_gpu.h index a7cda77663..212fde53dd 100644 --- a/src/gromacs/mdtypes/state_propagator_data_gpu.h +++ b/src/gromacs/mdtypes/state_propagator_data_gpu.h @@ -87,22 +87,22 @@ class StatePropagatorDataGpu * PME work on the GPU, and if that rank also does PP work that is the only * rank. So all coordinates are always transferred. * - * \note \p commandStream and \p gpuContext are allowed to be nullptr if + * \note \p commandStream and \p deviceContext are allowed to be nullptr if * StatePropagatorDataGpu is not used in the OpenCL run (e.g. if PME * does not run on the GPU). * - * \todo Make \p CommandStream visible in the CPU parts of the code so we - * will not have to pass a void*. - * \todo Make \p Context visible in CPU parts of the code so we will not - * have to pass a void*. + * \todo A CommandStream is now visible in the CPU parts of the code so we + * can stop passing a void*. + * \todo A DeviceContext object is visible in CPU parts of the code so we + * can stop passing a void*. * * \param[in] commandStream GPU stream, nullptr allowed. - * \param[in] gpuContext GPU context, nullptr allowed. + * \param[in] deviceContext GPU context, nullptr allowed. * \param[in] transferKind H2D/D2H transfer call behavior (synchronous or not). * \param[in] paddingSize Padding size for coordinates buffer. */ StatePropagatorDataGpu(const void *commandStream, - const void *gpuContext, + const void *deviceContext, GpuApiCallBehavior transferKind, int paddingSize); //! Move constructor diff --git a/src/gromacs/mdtypes/state_propagator_data_gpu_impl.cpp b/src/gromacs/mdtypes/state_propagator_data_gpu_impl.cpp index 9a11956649..d66cfc552f 100644 --- a/src/gromacs/mdtypes/state_propagator_data_gpu_impl.cpp +++ b/src/gromacs/mdtypes/state_propagator_data_gpu_impl.cpp @@ -55,7 +55,7 @@ class StatePropagatorDataGpu::Impl }; StatePropagatorDataGpu::StatePropagatorDataGpu(const void * /* commandStream */, - const void * /* gpuContext */, + const void * /* deviceContext */, GpuApiCallBehavior /* transferKind */, int /* paddingSize */) : impl_(nullptr) diff --git a/src/gromacs/mdtypes/state_propagator_data_gpu_impl.h b/src/gromacs/mdtypes/state_propagator_data_gpu_impl.h index 6ba40d987e..f32d6df9de 100644 --- a/src/gromacs/mdtypes/state_propagator_data_gpu_impl.h +++ b/src/gromacs/mdtypes/state_propagator_data_gpu_impl.h @@ -71,22 +71,22 @@ class StatePropagatorDataGpu::Impl * PME work on the GPU, and if that rank also does PP work that is the only * rank. So all coordinates are always transferred. * - * \note \p commandStream and \p gpuContext are allowed to be nullptr if + * \note \p commandStream and \p deviceContext are allowed to be nullptr if * StatePropagatorDataGpu is not used in the OpenCL run (e.g. if PME * does not run on the GPU). * - * \todo Make CommandStream visible in the CPU parts of the code so we - * will not have to pass a void*. - * \todo Make a Context object visible in CPU parts of the code so we - * will not have to pass a void*. + * \todo A CommandStream is now visible in the CPU parts of the code so we + * can stop passing a void*. + * \todo A DeviceContext object is visible in CPU parts of the code so we + * can stop passing a void*. * * \param[in] commandStream GPU stream, nullptr allowed. - * \param[in] gpuContext GPU context, nullptr allowed. + * \param[in] deviceContext GPU context, nullptr allowed. * \param[in] transferKind H2D/D2H transfer call behavior (synchronous or not). * \param[in] paddingSize Padding size for coordinates buffer. */ Impl(const void *commandStream, - const void *gpuContext, + const void *deviceContext, GpuApiCallBehavior transferKind, int paddingSize); @@ -209,7 +209,7 @@ class StatePropagatorDataGpu::Impl /*! \brief GPU context (for OpenCL builds) * \todo Make a Context class usable in CPU code */ - Context gpuContext_ = nullptr; + DeviceContext deviceContext_ = nullptr; //! Default GPU calls behavior GpuApiCallBehavior transferKind_ = GpuApiCallBehavior::Async; //! Padding size for the coordinates buffer diff --git a/src/gromacs/mdtypes/state_propagator_data_gpu_impl_gpu.cpp b/src/gromacs/mdtypes/state_propagator_data_gpu_impl_gpu.cpp index 66cdfda9e9..89446f2477 100644 --- a/src/gromacs/mdtypes/state_propagator_data_gpu_impl_gpu.cpp +++ b/src/gromacs/mdtypes/state_propagator_data_gpu_impl_gpu.cpp @@ -63,7 +63,7 @@ namespace gmx { StatePropagatorDataGpu::Impl::Impl(gmx_unused const void *commandStream, - gmx_unused const void *gpuContext, + gmx_unused const void *deviceContext, GpuApiCallBehavior transferKind, int paddingSize) : transferKind_(transferKind), @@ -79,9 +79,9 @@ StatePropagatorDataGpu::Impl::Impl(gmx_unused const void *commandStream, { commandStream_ = *static_cast(commandStream); } - if (gpuContext != nullptr) + if (deviceContext != nullptr) { - gpuContext_ = *static_cast(gpuContext); + deviceContext_ = *static_cast(deviceContext); } #endif @@ -94,7 +94,7 @@ StatePropagatorDataGpu::Impl::~Impl() void StatePropagatorDataGpu::Impl::reinit(int numAtomsLocal, int numAtomsAll) { #if GMX_GPU == GMX_GPU_OPENCL - GMX_ASSERT(gpuContext_ != nullptr, "GPU context should be set in OpenCL builds."); + GMX_ASSERT(deviceContext_ != nullptr, "GPU context should be set in OpenCL builds."); #endif numAtomsLocal_ = numAtomsLocal; numAtomsAll_ = numAtomsAll; @@ -109,7 +109,7 @@ void StatePropagatorDataGpu::Impl::reinit(int numAtomsLocal, int numAtomsAll) numAtomsPadded = numAtomsAll_; } - reallocateDeviceBuffer(&d_x_, DIM*numAtomsPadded, &d_xSize_, &d_xCapacity_, gpuContext_); + reallocateDeviceBuffer(&d_x_, DIM*numAtomsPadded, &d_xSize_, &d_xCapacity_, deviceContext_); const size_t paddingAllocationSize = numAtomsPadded - numAtomsAll_; if (paddingAllocationSize > 0) @@ -117,8 +117,8 @@ void StatePropagatorDataGpu::Impl::reinit(int numAtomsLocal, int numAtomsAll) clearDeviceBufferAsync(&d_x_, DIM*numAtomsAll_, DIM*paddingAllocationSize, commandStream_); } - reallocateDeviceBuffer(&d_v_, DIM*numAtomsAll_, &d_vSize_, &d_vCapacity_, gpuContext_); - reallocateDeviceBuffer(&d_f_, DIM*numAtomsAll_, &d_fSize_, &d_fCapacity_, gpuContext_); + reallocateDeviceBuffer(&d_v_, DIM*numAtomsAll_, &d_vSize_, &d_vCapacity_, deviceContext_); + reallocateDeviceBuffer(&d_f_, DIM*numAtomsAll_, &d_fSize_, &d_fCapacity_, deviceContext_); } @@ -155,7 +155,7 @@ void StatePropagatorDataGpu::Impl::copyToDevice(DeviceBuffer { #if GMX_GPU == GMX_GPU_OPENCL - GMX_ASSERT(gpuContext_ != nullptr, "GPU context should be set in OpenCL builds."); + GMX_ASSERT(deviceContext_ != nullptr, "GPU context should be set in OpenCL builds."); #endif GMX_UNUSED_VALUE(dataSize); @@ -187,7 +187,7 @@ void StatePropagatorDataGpu::Impl::copyFromDevice(gmx::ArrayRef h_da { #if GMX_GPU == GMX_GPU_OPENCL - GMX_ASSERT(gpuContext_ != nullptr, "GPU context should be set in OpenCL builds."); + GMX_ASSERT(deviceContext_ != nullptr, "GPU context should be set in OpenCL builds."); #endif GMX_UNUSED_VALUE(dataSize); @@ -284,11 +284,11 @@ int StatePropagatorDataGpu::Impl::numAtomsAll() StatePropagatorDataGpu::StatePropagatorDataGpu(const void *commandStream, - const void *gpuContext, + const void *deviceContext, GpuApiCallBehavior transferKind, int paddingSize) : impl_(new Impl(commandStream, - gpuContext, + deviceContext, transferKind, paddingSize)) { diff --git a/src/gromacs/nbnxm/cuda/nbnxm_cuda_data_mgmt.cu b/src/gromacs/nbnxm/cuda/nbnxm_cuda_data_mgmt.cu index 3641d5eb5f..71e25d23b1 100644 --- a/src/gromacs/nbnxm/cuda/nbnxm_cuda_data_mgmt.cu +++ b/src/gromacs/nbnxm/cuda/nbnxm_cuda_data_mgmt.cu @@ -550,7 +550,7 @@ void gpu_init_pairlist(gmx_nbnxn_cuda_t *nb, iTimers.didPairlistH2D = true; } - Context context = nullptr; + DeviceContext context = nullptr; reallocateDeviceBuffer(&d_plist->sci, h_plist->sci.size(), &d_plist->nsci, &d_plist->sci_nalloc, context); diff --git a/src/gromacs/nbnxm/opencl/nbnxm_ocl_data_mgmt.cpp b/src/gromacs/nbnxm/opencl/nbnxm_ocl_data_mgmt.cpp index a26ab41c99..4123a65e39 100644 --- a/src/gromacs/nbnxm/opencl/nbnxm_ocl_data_mgmt.cpp +++ b/src/gromacs/nbnxm/opencl/nbnxm_ocl_data_mgmt.cpp @@ -823,7 +823,7 @@ void gpu_init_pairlist(gmx_nbnxn_ocl_t *nb, } // TODO most of this function is same in CUDA and OpenCL, move into the header - Context context = nb->dev_rundata->context; + DeviceContext context = nb->dev_rundata->context; reallocateDeviceBuffer(&d_plist->sci, h_plist->sci.size(), &d_plist->nsci, &d_plist->sci_nalloc, context); -- 2.22.0