#include "gromacs/gpu_utils/cudautils.cuh"
#include "gromacs/gpu_utils/devicebuffer.h"
#include "gromacs/gpu_utils/gpueventsynchronizer.cuh"
+#include "gromacs/gpu_utils/typecasts.cuh"
#include "gromacs/gpu_utils/vectype_ops.cuh"
#include "gromacs/math/vectypes.h"
#include "gromacs/pbcutil/ishift.h"
#include <cassert>
#include "gromacs/gpu_utils/cuda_kernel_utils.cuh"
-#include "gromacs/gpu_utils/cudautils.cuh"
+#include "gromacs/gpu_utils/typecasts.cuh"
#include "pme.cuh"
#include "pme_calculate_splines.cuh"
#include <cassert>
#include "gromacs/gpu_utils/cuda_kernel_utils.cuh"
-#include "gromacs/gpu_utils/cudautils.cuh"
+#include "gromacs/gpu_utils/typecasts.cuh"
#include "pme.cuh"
#include "pme_calculate_splines.cuh"
* Optional second stage of the spline_and_spread_kernel.
*
* \tparam[in] order PME interpolation order.
- * \tparam[in] wrapX A boolean which tells if the grid overlap in dimension X should
- * be wrapped. \tparam[in] wrapY A boolean which tells if the grid overlap in
- * dimension Y should be wrapped. \tparam[in] useOrderThreads A boolean which Tells if we
- * should use order threads per atom (order*order used if false) \param[in] kernelParams Input PME
- * CUDA data in constant memory. \param[in] atomIndexOffset Starting atom index for the
- * execution block w.r.t. global memory. \param[in] atomCharge Atom charge/coefficient of
- * atom processed by thread. \param[in] sm_gridlineIndices Atom gridline indices in the shared
- * memory. \param[in] sm_theta Atom spline values in the shared memory.
+ * \tparam[in] wrapX Whether the grid overlap in dimension X should be wrapped.
+ * \tparam[in] wrapY Whether the grid overlap in dimension Y should be wrapped.
+ * \tparam[in] useOrderThreads Whether we should use order threads per atom (order*order used if false).
+ *
+ * \param[in] kernelParams Input PME CUDA data in constant memory.
+ * \param[in] atomIndexOffset Starting atom index for the execution block w.r.t. global memory.
+ * \param[in] atomCharge Atom charge/coefficient of atom processed by thread.
+ * \param[in] sm_gridlineIndices Atom gridline indices in the shared memory.
+ * \param[in] sm_theta Atom spline values in the shared memory.
*/
template<const int order, const bool wrapX, const bool wrapY, const bool useOrderThreads>
__device__ __forceinline__ void spread_charges(const PmeGpuCudaKernelParams kernelParams,
rvec tmp = { b.x, b.y, b.z };
rvec_inc(a, tmp);
}
-/*! \brief Cast RVec buffer to float3 buffer.
- *
- * \param[in] in The RVec buffer to cast.
- *
- * \returns Buffer, casted to float3*.
- */
-static inline __host__ __device__ float3* asFloat3(gmx::RVec* in)
-{
- static_assert(sizeof(in[0]) == sizeof(float3),
- "Size of the host-side data-type is different from the size of the device-side "
- "counterpart.");
- return reinterpret_cast<float3*>(in);
-}
/*! \brief Wait for all taks in stream \p s to complete.
*
# CUDA-only test
list(APPEND SOURCES_FROM_CXX
pinnedmemorychecker.cpp
+ typecasts.cpp
+ typecasts_runner.cpp
)
+ gmx_add_libgromacs_sources(typecasts_runner.cu)
# 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.
*/
#include "gmxpre.h"
-#include "config.h"
-
#include <vector>
-#ifndef __CUDA_ARCH__
-/*! \brief Dummy definition to avoid compiler error
- *
- * \todo Find a better solution. Probably, move asFloat3(...) function to different header.
- */
-# define __CUDA_ARCH__ -1
-# include <cuda_runtime.h>
-# undef __CUDA_ARCH__
-#else
-# include <cuda_runtime.h>
-#endif
#include <gtest/gtest.h>
-#include "gromacs/gpu_utils/cudautils.cuh"
-#include "gromacs/math/vectypes.h"
-#include "gromacs/utility/real.h"
+#include "gromacs/utility/exceptions.h"
-#if GMX_GPU == GMX_GPU_CUDA
+#include "testutils/testasserts.h"
+#include "testutils/testmatchers.h"
+
+#include "typecasts_runner.h"
namespace gmx
{
namespace test
{
-TEST(GpuDataTypesCompatibilityTest, RVecAndFloat3)
+//! Test data in RVec format
+static const std::vector<RVec> rVecInput = { { 1.0, 2.0, 3.0 }, { 4.0, 5.0, 6.0 } };
+
+TEST(GpuDataTypesCompatibilityTest, RVecAndFloat3OnHost)
{
- std::vector<RVec> dataRVec;
- dataRVec.emplace_back(1.0, 2.0, 3.0);
- dataRVec.emplace_back(4.0, 5.0, 6.0);
- float3* dataFloat3 = asFloat3(dataRVec.data());
- EXPECT_EQ(dataFloat3[0].x, dataRVec[0][XX]);
- EXPECT_EQ(dataFloat3[0].y, dataRVec[0][YY]);
- EXPECT_EQ(dataFloat3[0].z, dataRVec[0][ZZ]);
- EXPECT_EQ(dataFloat3[1].x, dataRVec[1][XX]);
- EXPECT_EQ(dataFloat3[1].y, dataRVec[1][YY]);
- EXPECT_EQ(dataFloat3[1].z, dataRVec[1][ZZ]);
+ std::vector<RVec> rVecOutput(rVecInput.size());
+ convertRVecToFloat3OnHost(rVecOutput, rVecInput);
+ EXPECT_THAT(rVecInput, testing::Pointwise(RVecEq(ulpTolerance(0)), rVecOutput));
}
-} // namespace test
-} // namespace gmx
+TEST(GpuDataTypesCompatibilityTest, RVecAndFloat3OnDevice)
+{
+ std::vector<RVec> rVecOutput(rVecInput.size());
+ convertRVecToFloat3OnDevice(rVecOutput, rVecInput);
+ EXPECT_THAT(rVecInput, testing::Pointwise(RVecEq(ulpTolerance(0)), rVecOutput));
+}
-#endif // GMX_GPU == GMX_GPU_CUDA
\ No newline at end of file
+} // namespace test
+} // namespace gmx
\ No newline at end of file
--- /dev/null
+/*
+ * This file is part of the GROMACS molecular simulation package.
+ *
+ * Copyright (c) 2020, 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
+ * Stub for runners for tests types compatibility.
+ *
+ * \author Artem Zhmurov <zhmurov@gmail.com>
+ */
+#include "gmxpre.h"
+
+#include "typecasts_runner.h"
+
+#include "config.h"
+
+#include <vector>
+
+#include "testutils/testasserts.h"
+
+#if GMX_GPU != GMX_GPU_CUDA
+
+namespace gmx
+{
+
+namespace test
+{
+
+void convertRVecToFloat3OnHost(std::vector<gmx::RVec>& /* rVecOutput */,
+ const std::vector<gmx::RVec>& /* rVecInput */)
+{
+ FAIL() << "Can't test float3 and RVec compatibility without CUDA.");
+}
+
+void convertRVecToFloat3OnDevice(std::vector<gmx::RVec>& /* rVecOutput */,
+ const std::vector<gmx::RVec>& /* rVecInput*/)
+{
+ FAIL() << "Can't test float3 and RVec compatibility without CUDA.");
+}
+
+} // namespace test
+} // namespace gmx
+
+#endif // GMX_GPU != GMX_GPU_CUDA
\ No newline at end of file
--- /dev/null
+/*
+ * This file is part of the GROMACS molecular simulation package.
+ *
+ * Copyright (c) 2020, 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
+ * Runners for tests of CUDA types compatibility.
+ *
+ * \author Artem Zhmurov <zhmurov@gmail.com>
+ */
+#include "gmxpre.h"
+
+#include "typecasts_runner.h"
+
+#include "config.h"
+
+#include <vector>
+
+#include "gromacs/gpu_utils/cudautils.cuh"
+#include "gromacs/gpu_utils/devicebuffer.h"
+#include "gromacs/gpu_utils/gpu_testutils.h"
+#include "gromacs/gpu_utils/typecasts.cuh"
+#include "gromacs/utility/exceptions.h"
+#include "gromacs/utility/stringutil.h"
+
+#if GMX_GPU == GMX_GPU_CUDA
+
+namespace gmx
+{
+
+namespace test
+{
+
+/* \brief Perform a component-wise conversion of the float3 vector back to RVec format.
+ *
+ * This is needed to pass the data back to the CPU testing code for comparison with the initial input.
+ *
+ * \param[out] rVecOutput Output data in RVec format for the output.
+ * \param[in] float3Output Output data in float3 format.
+ * \param[in] numElements Size of the data buffers.
+ */
+void inline saveFloat3InRVecFormat(std::vector<gmx::RVec>& rVecOutput, const float3* float3Output, int numElements)
+{
+ for (int i = 0; i < numElements; i++)
+ {
+ rVecOutput[i][XX] = float3Output[i].x;
+ rVecOutput[i][YY] = float3Output[i].y;
+ rVecOutput[i][ZZ] = float3Output[i].z;
+ }
+}
+
+void convertRVecToFloat3OnHost(std::vector<gmx::RVec>& rVecOutput, const std::vector<gmx::RVec>& rVecInput)
+{
+ const int numElements = rVecInput.size();
+
+ float3* dataFloat3 = asFloat3(const_cast<RVec*>(rVecInput.data()));
+
+ saveFloat3InRVecFormat(rVecOutput, dataFloat3, numElements);
+}
+
+//! Number of CUDA threads in a block.
+constexpr static int c_threadsPerBlock = 256;
+
+/*! \brief GPU kernel to perform type conversion on the device.
+ *
+ * \param[out] gm_float3Output Buffer to write the output into.
+ * \param[in] gm_rVecInput Input data in RVec format.
+ * \param[in] size Size of the data buffers.
+ *
+ */
+static __global__ void convertRVecToFloat3OnDevice_kernel(DeviceBuffer<float3> gm_float3Output,
+ DeviceBuffer<RVec> gm_rVecInput,
+ const int size)
+{
+ int threadIndex = blockIdx.x * blockDim.x + threadIdx.x;
+ if (threadIndex < size)
+ {
+ gm_float3Output[threadIndex] = asFloat3(gm_rVecInput)[threadIndex];
+ }
+}
+
+void convertRVecToFloat3OnDevice(std::vector<gmx::RVec>& h_rVecOutput, const std::vector<gmx::RVec>& h_rVecInput)
+{
+ if (canComputeOnGpu())
+ {
+ const int numElements = h_rVecInput.size();
+
+ DeviceBuffer<RVec> d_rVecInput;
+ allocateDeviceBuffer(&d_rVecInput, numElements, nullptr);
+ copyToDeviceBuffer(&d_rVecInput, h_rVecInput.data(), 0, numElements, nullptr,
+ GpuApiCallBehavior::Sync, nullptr);
+
+ DeviceBuffer<float3> d_float3Output;
+ allocateDeviceBuffer(&d_float3Output, numElements * DIM, nullptr);
+
+ std::vector<float3> h_float3Output(numElements);
+
+ KernelLaunchConfig kernelLaunchConfig;
+ kernelLaunchConfig.gridSize[0] = (numElements + c_threadsPerBlock - 1) / c_threadsPerBlock;
+ kernelLaunchConfig.blockSize[0] = c_threadsPerBlock;
+ kernelLaunchConfig.blockSize[1] = 1;
+ kernelLaunchConfig.blockSize[2] = 1;
+ kernelLaunchConfig.sharedMemorySize = 0;
+ kernelLaunchConfig.stream = nullptr;
+
+ auto kernelPtr = convertRVecToFloat3OnDevice_kernel;
+ const auto kernelArgs = prepareGpuKernelArguments(
+ kernelPtr, kernelLaunchConfig, &d_float3Output, &d_rVecInput, &numElements);
+ launchGpuKernel(kernelPtr, kernelLaunchConfig, nullptr,
+ "convertRVecToFloat3OnDevice_kernel", kernelArgs);
+
+ copyFromDeviceBuffer(h_float3Output.data(), &d_float3Output, 0, numElements, nullptr,
+ GpuApiCallBehavior::Sync, nullptr);
+
+ saveFloat3InRVecFormat(h_rVecOutput, h_float3Output.data(), numElements);
+
+ freeDeviceBuffer(&d_rVecInput);
+ freeDeviceBuffer(&d_float3Output);
+ }
+}
+
+} // namespace test
+} // namespace gmx
+
+#endif // GMX_GPU == GMX_GPU_CUDA
\ No newline at end of file
--- /dev/null
+/*
+ * This file is part of the GROMACS molecular simulation package.
+ *
+ * Copyright (c) 2020, 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
+ * Header for runner for CUDA float3 type layout tests.
+ *
+ * \author Artem Zhmurov <zhmurov@gmail.com>
+ */
+#ifndef GMX_GPU_UTILS_TESTS_TYPECASTS_RUNNER_H
+#define GMX_GPU_UTILS_TESTS_TYPECASTS_RUNNER_H
+
+#include "gmxpre.h"
+
+#include <vector>
+
+#include "gromacs/math/vectypes.h"
+
+namespace gmx
+{
+
+namespace test
+{
+
+/*! \brief Tests the compatibility of RVec and float3 using the conversion on host.
+ *
+ * \param[out] rVecOutput Data in RVec format for the output.
+ * \param[in] rVecInput Data in RVec format with the input.
+ */
+void convertRVecToFloat3OnHost(std::vector<gmx::RVec>& rVecOutput, const std::vector<gmx::RVec>& rVecInput);
+
+/*! \brief Tests the compatibility of RVec and float3 using the conversion on device.
+ *
+ * \param[out] rVecOutput Data in RVec format for the output.
+ * \param[in] rVecInput Data in RVec format with the input.
+ */
+void convertRVecToFloat3OnDevice(std::vector<gmx::RVec>& rVecOutput, const std::vector<gmx::RVec>& rVecInput);
+
+} // namespace test
+} // namespace gmx
+
+#endif // GMX_GPU_UTILS_TESTS_TYPECASTS_RUNNER_H
\ No newline at end of file
--- /dev/null
+/*
+ * This file is part of the GROMACS molecular simulation package.
+ *
+ * Copyright (c) 2020, 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 Declare functions to be used to cast CPU types to compatible GPU types.
+ *
+ * \author Artem Zhmurov <zhmurov@gmail.com>
+ *
+ * \inlibraryapi
+ */
+#ifndef GMX_GPU_UTILS_TYPECASTS_CUH
+#define GMX_GPU_UTILS_TYPECASTS_CUH
+
+#include "gmxpre.h"
+
+#include "gromacs/math/vectypes.h"
+
+/*! \brief Cast RVec buffer to float3 buffer.
+ *
+ * \param[in] in The RVec buffer to cast.
+ *
+ * \returns Buffer, casted to float3*.
+ */
+static inline __host__ __device__ float3* asFloat3(gmx::RVec* in)
+{
+ static_assert(sizeof(in[0]) == sizeof(float3),
+ "Size of the host-side data-type is different from the size of the device-side "
+ "counterpart.");
+ return reinterpret_cast<float3*>(in);
+}
+
+#endif // GMX_GPU_UTILS_TYPECASTS_CUH
#include "gromacs/gpu_utils/cudautils.cuh"
#include "gromacs/gpu_utils/gpueventsynchronizer.cuh"
+#include "gromacs/gpu_utils/typecasts.cuh"
#include "gromacs/gpu_utils/vectype_ops.cuh"
#include "gromacs/mdtypes/simulation_workload.h"
#include "gromacs/nbnxm/atomdata.h"
biod.pnpi.spb.ru / alexxy / gromacs.git / commitdiff