From f9885259b51b0bc7e6c3a7cb24c62b53fd1d1395 Mon Sep 17 00:00:00 2001 From: Mark Abraham Date: Mon, 6 Nov 2017 08:45:49 +0100 Subject: [PATCH] Introduce HostAllocationPolicy This permits host-side standard containers and smart pointers to have their contents placed in memory suitable for efficient GPU transfer. The behaviour can be configured at run time during simulation setup, so that if we are not running on a GPU, then none of the buffers that might be affected actually are. The downside is that all such containers now have state. Change-Id: I9367d0f996de04c21312cef2081cc08148f80561 --- src/gromacs/gpu_utils/CMakeLists.txt | 32 ++- src/gromacs/gpu_utils/hostallocator.cpp | 85 ++++++++ src/gromacs/gpu_utils/hostallocator.cu | 101 +++++++++ src/gromacs/gpu_utils/hostallocator.h | 139 ++++++++++++ src/gromacs/gpu_utils/{ => tests}/CMakeLists.txt | 50 ++++- src/gromacs/gpu_utils/tests/devicetransfers.cpp | 56 +++++ src/gromacs/gpu_utils/tests/devicetransfers.cu | 109 ++++++++++ src/gromacs/gpu_utils/tests/devicetransfers.h | 73 +++++++ .../gpu_utils/tests/devicetransfers_ocl.cpp | 108 ++++++++++ src/gromacs/gpu_utils/tests/gputest.cpp | 76 +++++++ src/gromacs/gpu_utils/tests/gputest.h | 72 +++++++ src/gromacs/gpu_utils/tests/hostallocator.cpp | 233 +++++++++++++++++++++ src/gromacs/utility/allocator.h | 28 ++- src/gromacs/utility/tests/alignedallocator.cpp | 7 + 14 files changed, 1151 insertions(+), 18 deletions(-) create mode 100644 src/gromacs/gpu_utils/hostallocator.cpp create mode 100644 src/gromacs/gpu_utils/hostallocator.cu create mode 100644 src/gromacs/gpu_utils/hostallocator.h copy src/gromacs/gpu_utils/{ => tests}/CMakeLists.txt (54%) create mode 100644 src/gromacs/gpu_utils/tests/devicetransfers.cpp create mode 100644 src/gromacs/gpu_utils/tests/devicetransfers.cu create mode 100644 src/gromacs/gpu_utils/tests/devicetransfers.h create mode 100644 src/gromacs/gpu_utils/tests/devicetransfers_ocl.cpp create mode 100644 src/gromacs/gpu_utils/tests/gputest.cpp create mode 100644 src/gromacs/gpu_utils/tests/gputest.h create mode 100644 src/gromacs/gpu_utils/tests/hostallocator.cpp diff --git a/src/gromacs/gpu_utils/CMakeLists.txt b/src/gromacs/gpu_utils/CMakeLists.txt index 00844e8996..27a190cf7a 100644 --- a/src/gromacs/gpu_utils/CMakeLists.txt +++ b/src/gromacs/gpu_utils/CMakeLists.txt @@ -1,7 +1,7 @@ # # This file is part of the GROMACS molecular simulation package. # -# Copyright (c) 2015,2016, by the GROMACS development team, led by +# Copyright (c) 2015,2016,2017, 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. @@ -32,12 +32,28 @@ # To help us fund GROMACS development, we humbly ask that you cite # the research papers on the package. Check out http://www.gromacs.org. -if (GMX_USE_OPENCL) - gmx_add_libgromacs_sources(gpu_utils_ocl.cpp ocl_compiler.cpp ocl_caching.cpp oclutils.cpp) +if(GMX_USE_OPENCL) + gmx_add_libgromacs_sources( + gpu_utils_ocl.cpp + hostallocator.cpp + ocl_compiler.cpp + ocl_caching.cpp + oclutils.cpp + ) +elseif(GMX_USE_CUDA) + gmx_add_libgromacs_sources( + cudautils.cu + hostallocator.cu + gpu_utils.cu + pmalloc_cuda.cu + ) +else() + gmx_add_libgromacs_sources( + gpu_utils.cpp + hostallocator.cpp + ) endif() -if (GMX_USE_CUDA) - gmx_add_libgromacs_sources(cudautils.cu gpu_utils.cu pmalloc_cuda.cu) -endif() -if (NOT GMX_USE_OPENCL AND NOT GMX_USE_CUDA) - gmx_add_libgromacs_sources(gpu_utils.cpp) + +if (BUILD_TESTING) + add_subdirectory(tests) endif() diff --git a/src/gromacs/gpu_utils/hostallocator.cpp b/src/gromacs/gpu_utils/hostallocator.cpp new file mode 100644 index 0000000000..3ceaefeb7b --- /dev/null +++ b/src/gromacs/gpu_utils/hostallocator.cpp @@ -0,0 +1,85 @@ +/* + * This file is part of the GROMACS molecular simulation package. + * + * Copyright (c) 2017, 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 Implements gmx::HostAllocationPolicy for allocating memory + * suitable for GPU transfers on OpenCL, and when no GPU + * implementation is used. + * + * \todo The same implementation can be used because we do not + * currently attempt to optimize the allocation of host-side buffers + * for OpenCL transfers, but this might be good to do. + * + * \author Mark Abraham + */ +#include "gmxpre.h" + +#include "hostallocator.h" + +#include + +#include "gromacs/utility/alignedallocator.h" + +namespace gmx +{ + +HostAllocationPolicy::HostAllocationPolicy(Impl s) : allocateForGpu_(s) {} + +void * +HostAllocationPolicy::malloc(std::size_t bytes) const +{ + GMX_UNUSED_VALUE(allocateForGpu_); + // TODO if/when this is properly supported for OpenCL, we + // should explore whether it is needed, and if so what + // page size is desirable for alignment. + return AlignedAllocationPolicy::malloc(bytes); +} + +void +HostAllocationPolicy::free(void *buffer) const +{ + if (buffer == nullptr) + { + return; + } + GMX_UNUSED_VALUE(allocateForGpu_); + AlignedAllocationPolicy::free(buffer); +} + +HostAllocationPolicy makeHostAllocationPolicyForGpu() +{ + return HostAllocationPolicy(HostAllocationPolicy::Impl::AllocateForGpu); +} + +} // namespace gmx diff --git a/src/gromacs/gpu_utils/hostallocator.cu b/src/gromacs/gpu_utils/hostallocator.cu new file mode 100644 index 0000000000..bde92207da --- /dev/null +++ b/src/gromacs/gpu_utils/hostallocator.cu @@ -0,0 +1,101 @@ +/* + * This file is part of the GROMACS molecular simulation package. + * + * Copyright (c) 2017, 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 Implements gmx::HostAllocationPolicy for allocating memory + * suitable for GPU transfers on CUDA. + * + * \author Mark Abraham + */ +#include "gmxpre.h" + +#include "hostallocator.h" + +#include + +#include "gromacs/utility/alignedallocator.h" + +namespace gmx +{ + +HostAllocationPolicy::HostAllocationPolicy(Impl s) : allocateForGpu_(s) {} + +void * +HostAllocationPolicy::malloc(std::size_t bytes) const +{ + void *buffer = nullptr; + if (allocateForGpu_ == Impl::AllocateForGpu) + { + if (bytes != 0) + { + // Alternatively, this could become a pair of + // e.g. PageAlignedAllocationPolicy and cudaHostRegister + // calls if that is useful for something. + cudaError_t stat = cudaMallocHost(&buffer, bytes, cudaHostAllocDefault); + // TODO Throw an exception upon failure, particularly + // for cudaErrorMemoryAllocation. + if (stat != cudaSuccess) + { + buffer = nullptr; + } + } + } + else + { + buffer = AlignedAllocationPolicy::malloc(bytes); + } + return buffer; +} + +void +HostAllocationPolicy::free(void *buffer) const +{ + if (buffer == nullptr) + { + return; + } + if (allocateForGpu_ == Impl::AllocateForGpu) + { + cudaFreeHost(buffer); + return; + } + AlignedAllocationPolicy::free(buffer); +} + +HostAllocationPolicy makeHostAllocationPolicyForGpu() +{ + return HostAllocationPolicy(HostAllocationPolicy::Impl::AllocateForGpu); +} + +} // namespace gmx diff --git a/src/gromacs/gpu_utils/hostallocator.h b/src/gromacs/gpu_utils/hostallocator.h new file mode 100644 index 0000000000..dae54f2144 --- /dev/null +++ b/src/gromacs/gpu_utils/hostallocator.h @@ -0,0 +1,139 @@ +/* + * This file is part of the GROMACS molecular simulation package. + * + * Copyright (c) 2017, 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. + */ +/*! \libinternal \file + * \brief Declares gmx::HostAllocationPolicy and gmx::HostAllocator, + * which are used to make standard library containers that can + * allocate memory suitable for GPU transfers. + * + * \author Mark Abraham + * \inlibraryapi + */ +#ifndef GMX_GPU_UTILS_HOSTALLOCATOR_H +#define GMX_GPU_UTILS_HOSTALLOCATOR_H + +#include + +#include "gromacs/utility/allocator.h" + +namespace gmx +{ + +/*! \libinternal + * \brief Policy class for configuring gmx::Allocator, to manage + * allocations of memory that is suitable for GPU transfers. + * + * This allocator has state, so is most useful in cases where it is + * not known at compile time whether the allocated memory will be + * transferred to a GPU. It will increase the size of containers that + * use it. Memory allocated will always be aligned by the GPU + * framework, or by AlignedAllocationPolicy. + * + * \todo Consider also having a stateless version of this policy, + * which might be slightly faster or more convenient to use in the + * cases where it is known at compile time that the allocation will be + * used to transfer to a GPU. + */ +class HostAllocationPolicy +{ + public: + //! Helper construction enum + enum class Impl : int + { + AllocateAligned = 0, + AllocateForGpu = 1 + }; + //! Constructor. + explicit HostAllocationPolicy(Impl s = Impl::AllocateAligned); + /*! \brief Allocate GPU memory + * + * \param bytes Amount of memory (bytes) to allocate. It is + * valid to ask for 0 bytes, which will return a + * non-null pointer that is properly aligned in + * page-locked memory (but that you should not + * use). TODO check this. + * + * \return Valid pointer if the allocation worked, otherwise nullptr. + * + * The memory will always be allocated according to the requirements + * of the acceleration platform in use (e.g. CUDA). + * + * \note Memory allocated with this routine must be released + * with gmx::HostAllocationPolicy::free(), and + * absolutely not the system free(). + */ + void * + malloc(std::size_t bytes) const; + /*! \brief Free GPU memory + * + * \param buffer Memory pointer previously returned from gmx::HostAllocationPolicy::malloc() + * + * \note This routine should only be called with pointers + * obtained from gmx:HostAllocationPolicy::malloc(), + * and absolutely not any pointers obtained the system + * malloc(). + */ + void + free(void *buffer) const; + private: + /*! \brief State of the allocator. + * + * This could change through assignment of one policy to + * another, so isn't const. */ + Impl allocateForGpu_; +}; + +/*! \brief Convenience function + * + * The default construction is for non-GPU allocation, and this + * function makes it less verbose to get allocation intended for use + * with a GPU. */ +HostAllocationPolicy makeHostAllocationPolicyForGpu(); + +/*! \brief Memory allocator for host-side memory for GPU transfers. + * + * \tparam T Type of objects to allocate + * + * This convenience partial specialization can be used for the + * optional allocator template parameter in standard library + * containers whose memory will be used for GPU transfers. The memory + * will always be allocated according to the behavior of + * HostAllocationPolicy. + */ +template +using HostAllocator = Allocator; + +} // namespace gmx + +#endif diff --git a/src/gromacs/gpu_utils/CMakeLists.txt b/src/gromacs/gpu_utils/tests/CMakeLists.txt similarity index 54% copy from src/gromacs/gpu_utils/CMakeLists.txt copy to src/gromacs/gpu_utils/tests/CMakeLists.txt index 00844e8996..1d02bcee68 100644 --- a/src/gromacs/gpu_utils/CMakeLists.txt +++ b/src/gromacs/gpu_utils/tests/CMakeLists.txt @@ -1,7 +1,7 @@ # # This file is part of the GROMACS molecular simulation package. # -# Copyright (c) 2015,2016, by the GROMACS development team, led by +# Copyright (c) 2017, 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. @@ -32,12 +32,48 @@ # To help us fund GROMACS development, we humbly ask that you cite # the research papers on the package. Check out http://www.gromacs.org. -if (GMX_USE_OPENCL) - gmx_add_libgromacs_sources(gpu_utils_ocl.cpp ocl_compiler.cpp ocl_caching.cpp oclutils.cpp) +# Arrange to compile files with test infrastructure as normal C++, and +# other files that sometimes have platform-specific symbols using the +# required infrastructure to compile and link, without needing three +# different files and strategies for the different GPU implementation +# flavours. + +# Always compiled as plain C++ +file(GLOB SOURCES_FROM_CXX + hostallocator.cpp + ) + +if(GMX_USE_CUDA) + # TODO Making a separate library is heavy handed, but nothing else + # seems to work. Also don't use a hyphen in its name, because nvcc + # can't cope with that. + # + # Perhaps FindCUDA's support for single compilation units will help? + cuda_add_library(libgpu_utilstest + devicetransfers.cu + ) +elseif(GMX_USE_OPENCL) + # Do normal compilation of OpenCL files + list(APPEND SOURCES_FROM_CXX + devicetransfers_ocl.cpp + ) +else() + # Do normal compilation of files with null implementations + list(APPEND SOURCES_FROM_CXX + devicetransfers.cpp + ) endif() -if (GMX_USE_CUDA) - gmx_add_libgromacs_sources(cudautils.cu gpu_utils.cu pmalloc_cuda.cu) + +gmx_add_unit_test(GpuUtilsUnitTests gpu_utils-test + # Infrastructure + gputest.cpp + # Tests of code + ${SOURCES_FROM_CXX} + ) + +if(GMX_USE_CUDA) + target_link_libraries(gpu_utils-test libgpu_utilstest) endif() -if (NOT GMX_USE_OPENCL AND NOT GMX_USE_CUDA) - gmx_add_libgromacs_sources(gpu_utils.cpp) +if(GMX_USE_OPENCL) + target_link_libraries(gpu_utils-test ${OPENCL_LIBRARIES}) endif() diff --git a/src/gromacs/gpu_utils/tests/devicetransfers.cpp b/src/gromacs/gpu_utils/tests/devicetransfers.cpp new file mode 100644 index 0000000000..039f4a0583 --- /dev/null +++ b/src/gromacs/gpu_utils/tests/devicetransfers.cpp @@ -0,0 +1,56 @@ +/* + * This file is part of the GROMACS molecular simulation package. + * + * Copyright (c) 2017, 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 Defines helper functionality for device transfers for tests + * for GPU host allocator. + * + * \author Mark Abraham + */ +#include "gmxpre.h" + +#include "devicetransfers.h" + +#include "gromacs/utility/arrayref.h" + +namespace gmx +{ + +void doDeviceTransfers(const gmx_gpu_info_t & /*gpuInfo*/, + ArrayRef /*input*/, + ArrayRef /* output */) +{ +} + +} // namespace gmx diff --git a/src/gromacs/gpu_utils/tests/devicetransfers.cu b/src/gromacs/gpu_utils/tests/devicetransfers.cu new file mode 100644 index 0000000000..940712d489 --- /dev/null +++ b/src/gromacs/gpu_utils/tests/devicetransfers.cu @@ -0,0 +1,109 @@ +/* + * This file is part of the GROMACS molecular simulation package. + * + * Copyright (c) 2017, 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 Defines helper functionality for device transfers for tests + * for GPU host allocator. + * + * Undefined symbols in Google Test, GROMACS use of -Wundef, and the + * implementation of FindCUDA.cmake and/or nvcc mean that no + * compilation unit should include a gtest header while being compiled + * by nvcc. None of -isystem, -Wno-undef, nor the pragma GCC + * diagnostic work. + * + * \author Mark Abraham + */ +#include "gmxpre.h" + +#include "devicetransfers.h" + +#include "gromacs/gpu_utils/cudautils.cuh" +#include "gromacs/hardware/gpu_hw_info.h" +#include "gromacs/utility/arrayref.h" +#include "gromacs/utility/exceptions.h" +#include "gromacs/utility/gmxassert.h" +#include "gromacs/utility/stringutil.h" + +namespace gmx +{ +namespace +{ + +/*! \brief Help give useful diagnostics about error \c status while doing \c message. + * + * \throws InternalError If status indicates failure, supplying + * descriptive text from \c message. */ +static void throwUponFailure(cudaError_t status, const char *message) +{ + if (status != cudaSuccess) + { + GMX_THROW(InternalError(formatString("Failure while %s", message)));; + } +} + +} // namespace + +void doDeviceTransfers(const gmx_gpu_info_t &gpuInfo, + ArrayRef input, + ArrayRef output) +{ + GMX_RELEASE_ASSERT(input.size() == output.size(), "Input and output must have matching size"); + cudaError_t status; + GMX_RELEASE_ASSERT(gpuInfo.n_dev > 0, "Must have a GPU device"); + + const auto &device = gpuInfo.gpu_dev[0]; + int oldDeviceId; + + status = cudaGetDevice(&oldDeviceId); + throwUponFailure(status, "getting old device id"); + status = cudaSetDevice(device.id); + throwUponFailure(status, "setting device id to 0"); + + void *devicePointer; + status = cudaMalloc(&devicePointer, input.size()); + throwUponFailure(status, "creating buffer"); + + status = cudaMemcpy(devicePointer, input.data(), input.size(), cudaMemcpyHostToDevice); + throwUponFailure(status, "transferring host to device"); + status = cudaMemcpy(output.data(), devicePointer, output.size(), cudaMemcpyDeviceToHost); + throwUponFailure(status, "transferring device to host"); + + status = cudaFree(devicePointer); + throwUponFailure(status, "releasing buffer"); + + status = cudaSetDevice(oldDeviceId); + throwUponFailure(status, "setting old device id"); +} + +} // namespace gmx diff --git a/src/gromacs/gpu_utils/tests/devicetransfers.h b/src/gromacs/gpu_utils/tests/devicetransfers.h new file mode 100644 index 0000000000..7ea88a1723 --- /dev/null +++ b/src/gromacs/gpu_utils/tests/devicetransfers.h @@ -0,0 +1,73 @@ +/* + * This file is part of the GROMACS molecular simulation package. + * + * Copyright (c) 2017, 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 Declares helper functionality for device transfers for tests + * for GPU host allocator. + * + * Undefined symbols in Google Test, GROMACS use of -Wundef, and the + * implementation of FindCUDA.cmake and/or nvcc mean that no + * compilation unit should include a gtest header while being compiled + * by nvcc. None of -isystem, -Wno-undef, nor the pragma GCC + * diagnostic work. + * + * Thus, this header isolates CUDA-specific functionality to its own + * translation unit. The OpenCL and no-GPU implementations do not + * require this separation, but do so for consistency. + * + * \author Mark Abraham + */ +#ifndef GMX_GPU_UTILS_TESTS_DEVICETRANSFERS_H +#define GMX_GPU_UTILS_TESTS_DEVICETRANSFERS_H + +#include "gromacs/utility/arrayref.h" + +struct gmx_gpu_info_t; + +namespace gmx +{ + +/*! \brief Helper function for GPU test code to be platform agnostic. + * + * Transfers \c input to device 0, which must be present, and + * transfers it back into \c output. Both sizes must match. + * + * \throws InternalError Upon any GPU API error condition. */ +void doDeviceTransfers(const gmx_gpu_info_t &gpuInfo, + ArrayRef input, + ArrayRef output); + +} // namespace gmx + +#endif diff --git a/src/gromacs/gpu_utils/tests/devicetransfers_ocl.cpp b/src/gromacs/gpu_utils/tests/devicetransfers_ocl.cpp new file mode 100644 index 0000000000..4b31ae3f0c --- /dev/null +++ b/src/gromacs/gpu_utils/tests/devicetransfers_ocl.cpp @@ -0,0 +1,108 @@ +/* + * This file is part of the GROMACS molecular simulation package. + * + * Copyright (c) 2017, 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 Defines helper functionality for device transfers for tests + * for GPU host allocator. + * + * \author Mark Abraham + */ +#include "gmxpre.h" + +#include "gromacs/gpu_utils/gmxopencl.h" +#include "gromacs/gpu_utils/oclutils.h" +#include "gromacs/hardware/gpu_hw_info.h" +#include "gromacs/utility/arrayref.h" +#include "gromacs/utility/exceptions.h" +#include "gromacs/utility/gmxassert.h" +#include "gromacs/utility/stringutil.h" + +#include "devicetransfers.h" + +namespace gmx +{ +namespace +{ + +/*! \brief Help give useful diagnostics about error \c status while doing \c message. + * + * \throws InternalError If status indicates failure, supplying + * descriptive text from \c message. */ +static void throwUponFailure(cl_int status, const char *message) +{ + if (status != CL_SUCCESS) + { + GMX_THROW(InternalError(formatString("Failure while %s, error was %s", message, ocl_get_error_string(status).c_str()))); + } +} + +} // namespace + +void doDeviceTransfers(const gmx_gpu_info_t &gpuInfo, + ArrayRef input, + ArrayRef output) +{ + GMX_RELEASE_ASSERT(input.size() == output.size(), "Input and output must have matching size"); + cl_int status; + GMX_RELEASE_ASSERT(gpuInfo.n_dev > 0, "Must have a GPU device"); + + const auto &device = gpuInfo.gpu_dev[0]; + cl_context_properties properties[] = { + CL_CONTEXT_PLATFORM, + (cl_context_properties) device.ocl_gpu_id.ocl_platform_id, + 0 + }; + // Give uncrustify more space + + auto deviceId = device.ocl_gpu_id.ocl_device_id; + auto context = clCreateContext(properties, 1, &deviceId, NULL, NULL, &status); + throwUponFailure(status, "creating context"); + auto commandQueue = clCreateCommandQueue(context, deviceId, 0, &status); + throwUponFailure(status, "creating command queue"); + + auto devicePointer = clCreateBuffer(context, CL_MEM_READ_WRITE, input.size(), nullptr, &status); + throwUponFailure(status, "creating buffer"); + + status = clEnqueueWriteBuffer(commandQueue, devicePointer, CL_TRUE, 0, input.size(), input.data(), 0, nullptr, nullptr); + throwUponFailure(status, "transferring host to device"); + status = clEnqueueReadBuffer(commandQueue, devicePointer, CL_TRUE, 0, output.size(), output.data(), 0, nullptr, nullptr); + throwUponFailure(status, "transferring device to host"); + + status = clReleaseMemObject(devicePointer); + throwUponFailure(status, "releasing buffer"); + status = clReleaseContext(context); + throwUponFailure(status, "releasing context"); +} + +} // namespace gmx diff --git a/src/gromacs/gpu_utils/tests/gputest.cpp b/src/gromacs/gpu_utils/tests/gputest.cpp new file mode 100644 index 0000000000..2cd80e1609 --- /dev/null +++ b/src/gromacs/gpu_utils/tests/gputest.cpp @@ -0,0 +1,76 @@ +/* + * This file is part of the GROMACS molecular simulation package. + * + * Copyright (c) 2017, 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 + * Tests utilities for GPU device allocation and free. + * + * \author Mark Abraham + */ +#include "gmxpre.h" + +#include "gputest.h" + +#include + +#include "gromacs/gpu_utils/gpu_utils.h" +#include "gromacs/hardware/gpu_hw_info.h" +#include "gromacs/utility/cstringutil.h" +#include "gromacs/utility/smalloc.h" + +namespace gmx +{ +namespace test +{ + +GpuTest::GpuTest() +{ + snew(gpuInfo_, 1); + char errorString[STRLEN]; + detect_gpus(gpuInfo_, errorString); +} + +GpuTest::~GpuTest() +{ + free_gpu_info(gpuInfo_); + sfree(gpuInfo_); +} + +bool GpuTest::haveValidGpus() const +{ + return gpuInfo_->n_dev_compatible > 0; +} + +} // namespace +} // namespace diff --git a/src/gromacs/gpu_utils/tests/gputest.h b/src/gromacs/gpu_utils/tests/gputest.h new file mode 100644 index 0000000000..eafae6a8ba --- /dev/null +++ b/src/gromacs/gpu_utils/tests/gputest.h @@ -0,0 +1,72 @@ +/* + * This file is part of the GROMACS molecular simulation package. + * + * Copyright (c) 2017, 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 + * Declares test fixture testing GPU utility components. + * + * \author Mark Abraham + */ +#ifndef GMX_GPU_UTILS_TESTS_GPUTEST_H +#define GMX_GPU_UTILS_TESTS_GPUTEST_H + +#include "gmxpre.h" + +#include + +struct gmx_gpu_info_t; + +namespace gmx +{ +namespace test +{ + +class GpuTest : public ::testing::Test +{ + public: + //! Information about GPUs that are present. + gmx_gpu_info_t *gpuInfo_; + + //! Constructor + GpuTest(); + //! Destructor + ~GpuTest(); + //! Getter for convenience in testing + bool haveValidGpus() const; +}; + +} // namespace +} // namespace + +#endif diff --git a/src/gromacs/gpu_utils/tests/hostallocator.cpp b/src/gromacs/gpu_utils/tests/hostallocator.cpp new file mode 100644 index 0000000000..689987bec7 --- /dev/null +++ b/src/gromacs/gpu_utils/tests/hostallocator.cpp @@ -0,0 +1,233 @@ +/* + * This file is part of the GROMACS molecular simulation package. + * + * Copyright (c) 2017, 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 + * Tests for GPU host allocator. + * + * \author Mark Abraham + */ +#include "gmxpre.h" + +#include "gromacs/gpu_utils/hostallocator.h" + +#include +#include + +#include + +#include "gromacs/math/vectypes.h" +#include "gromacs/utility/real.h" + +#include "devicetransfers.h" +#include "gputest.h" + +namespace gmx +{ + +namespace +{ + +//! The types used in testing. +typedef ::testing::Types TestTypes; + +//! Typed test fixture +template +class HostAllocatorTest : public test::GpuTest +{ + public: + //! Convenience type + using ValueType = T; + //! Convenience type + using AllocatorType = HostAllocator; + //! Convenience type + using VectorType = std::vector; + //! Convenience type + using ViewType = ArrayRef; + //! Convenience type + using ConstViewType = ArrayRef; + //! Prepare contents of a VectorType. + void fillInput(VectorType *input) const; + //! Compares input and output vectors. + void compareVectors(ConstViewType input, + ConstViewType output) const; + //! Do some transfers and test the results. + void runTest(ConstViewType input, ViewType output) const; +}; + +// Already documented +template +void HostAllocatorTest::fillInput(VectorType *input) const +{ + input->push_back(1); + input->push_back(2); + input->push_back(3); +} + +//! Initialization specialization for RVec +template <> +void HostAllocatorTest::fillInput(VectorType *input) const +{ + input->push_back({1, 2, 3}); +} + +// Already documented +template +void HostAllocatorTest::compareVectors(ConstViewType input, + ConstViewType output) const +{ + for (size_t i = 0; i != input.size(); ++i) + { + EXPECT_EQ(input[i], output[i]) << "for index " << i; + } +} + +//! Comparison specialization for RVec +template <> +void HostAllocatorTest::compareVectors(ConstViewType input, + ConstViewType output) const +{ + for (size_t i = 0; i != input.size(); ++i) + { + EXPECT_EQ(input[i][XX], output[i][XX]) << "for index " << i; + EXPECT_EQ(input[i][YY], output[i][YY]) << "for index " << i; + EXPECT_EQ(input[i][ZZ], output[i][ZZ]) << "for index " << i; + } +} + +/*! \brief Convenience function to transform a view into one with base + * type of (non-const) char. + * + * This transformation is useful for using containers with C APIs + * where the function signature is not declared const even where the + * semantics of the usage actually are const. + * + * \param[in] data The data pointer. + * \param[in] size The size of the data pointer (in T). + * \tparam T The base type of the container + * */ +template +ArrayRef charArrayRefFromArray(T *data, size_t size) +{ + // Make a type like T, but without its possible const qualifier. + using NonConstT = typename std::remove_const::type; + return arrayRefFromArray(reinterpret_cast(const_cast(data)), size * sizeof(T)); +} + +template +void HostAllocatorTest::runTest(ConstViewType input, ViewType output) const +{ + // We can't do a test that does a transfer unless we have a + // compatible device. + if (!this->haveValidGpus()) + { + return; + } + + // Convert the views of input and output to flat non-const chars, + // so that there's no templating when we call doDeviceTransfers. + auto inputRef = charArrayRefFromArray(input.data(), input.size()); + auto outputRef = charArrayRefFromArray(output.data(), output.size()); + + doDeviceTransfers(*this->gpuInfo_, inputRef, outputRef); + this->compareVectors(input, output); +} + +TYPED_TEST_CASE(HostAllocatorTest, TestTypes); + +// Note that in GoogleTest typed tests, the use of TestFixture:: and +// this-> is sometimes required to get access to things in the fixture +// class (or its base classes). + +// Note also that aspects of this code can be tested even when a GPU +// device is not available. + +TYPED_TEST(HostAllocatorTest, EmptyMemoryAlwaysWorks) +{ + typename TestFixture::VectorType v; +} + +TYPED_TEST(HostAllocatorTest, TransfersUsingDefaultHostAllocatorWork) +{ + typename TestFixture::VectorType input = {{1, 2, 3}}, output; + output.resize(input.size()); + + this->runTest(input, output); +} + +TYPED_TEST(HostAllocatorTest, TransfersUsingNormalCpuHostAllocatorWork) +{ + // Make an allocator with a 'normal CPU' allocation policy. This + // might be slower than another policy, but still works. + using AllocatorType = typename TestFixture::AllocatorType; + using AllocatorPolicyType = typename AllocatorType::allocation_policy; + AllocatorPolicyType policy(AllocatorPolicyType::Impl::AllocateAligned); + AllocatorType allocator(policy); + + typename TestFixture::VectorType input(allocator); + this->fillInput(&input); + typename TestFixture::VectorType output(allocator); + output.resize(input.size()); + + this->runTest(input, output); +} + +TYPED_TEST(HostAllocatorTest, TransfersUsingGpuHostAllocatorWork) +{ + // Make an allocator with a 'for GPU' allocation policy. This + // should be more efficient, but we can't test that. + using AllocatorType = typename TestFixture::AllocatorType; + using AllocatorPolicyType = typename AllocatorType::allocation_policy; + AllocatorPolicyType policy(AllocatorPolicyType::Impl::AllocateForGpu); + AllocatorType allocator(policy); + + typename TestFixture::VectorType input(allocator); + this->fillInput(&input); + typename TestFixture::VectorType output(allocator); + output.resize(input.size()); + + this->runTest(input, output); +} + +TYPED_TEST(HostAllocatorTest, StatefulAllocatorUsesMemory) +{ + // The HostAllocator has state, so a container using it will be + // larger than a normal vector, whose default allocator is + // stateless. + EXPECT_LT(sizeof(std::vector), + sizeof(typename TestFixture::VectorType)); +} + +} // namespace +} // namespace diff --git a/src/gromacs/utility/allocator.h b/src/gromacs/utility/allocator.h index c81a5371c7..50ee1f0834 100644 --- a/src/gromacs/utility/allocator.h +++ b/src/gromacs/utility/allocator.h @@ -69,7 +69,23 @@ namespace gmx * e.g. with SIMD alignment, GPU host-side page locking, or perhaps * both, in a way that preserves a common programming interface and * duplicates minimal code. - + * + * AllocationPolicy is used as a base class, so that if + * AllocationPolicy is stateless, then the empty base optimization + * will ensure that Allocation is also stateless, and objects made + * with the Allocator will incur no size penalty. (Embedding an + * AllocationPolicy object incurs a size penalty always, even if the + * object is empty.) Normally a stateless allocator will be used. + * + * However, an AllocationPolicy with state might be desirable for + * simplifying writing code that needs to allocate suitably for a + * transfer to a GPU. That code needs to specify an Allocator that can + * do the right job, which can be stateless. However, if we have code + * that will not know until run time whether a GPU transfer will + * occur, then the allocator needs to be aware of the state. That + * will increase the size of a container that uses the stateful + * allocator. + * * \throws std::bad_alloc Instead of a GROMACS exception object, we * throw the standard one on allocation failures to make it as * compatible as possible with the errors expected by code using the @@ -79,7 +95,7 @@ namespace gmx * \ingroup module_utility */ template -class Allocator +class Allocator : public AllocationPolicy { public: // The standard library specification for a custom allocator @@ -124,7 +140,13 @@ class Allocator * No constructor can be auto-generated in the presence of any * user-defined constructor, but we want the default constructor. */ - Allocator() {}; + Allocator() = default; + + /*! \brief Constructor to accept an AllocationPolicy. + * + * This is useful for AllocationPolicies with state. + */ + Allocator(const AllocationPolicy &p) : AllocationPolicy(p) {} /*! \brief Return address of an object * diff --git a/src/gromacs/utility/tests/alignedallocator.cpp b/src/gromacs/utility/tests/alignedallocator.cpp index d07585d28c..09dc610157 100644 --- a/src/gromacs/utility/tests/alignedallocator.cpp +++ b/src/gromacs/utility/tests/alignedallocator.cpp @@ -122,4 +122,11 @@ TYPED_TEST(AllocatorTest, VectorAllocatesAndReservesWithAlignment) } } +TYPED_TEST(AllocatorTest, StatelessAllocatorUsesNoMemory) +{ + using value_type = typename TypeParam::value_type; + EXPECT_EQ(sizeof(std::vector), + sizeof(std::vector)); +} + } -- 2.11.4.GIT