Unify CUDA and OpenCL lookup-table creation

[alexxy/gromacs.git] / src / gromacs / gpu_utils / devicebuffer.cuh

diff --git a/src/gromacs/gpu_utils/devicebuffer.cuh b/src/gromacs/gpu_utils/devicebuffer.cuh
index d4bfe8c35b5b47fc6fd7c3f6190489b99fcc6abf..c9bce141fef3d338a9f19238b77db36c52aed7d2 100644 (file)
--- a/src/gromacs/gpu_utils/devicebuffer.cuh
+++ b/src/gromacs/gpu_utils/devicebuffer.cuh
@@ -45,11 +45,13 @@
  *  \inlibraryapi
  */
 
+#include "gromacs/gpu_utils/cuda_arch_utils.cuh"
 #include "gromacs/gpu_utils/device_context.h"
 #include "gromacs/gpu_utils/devicebuffer_datatype.h"
 #include "gromacs/gpu_utils/gpu_utils.h" //only for GpuApiCallBehavior
 #include "gromacs/gpu_utils/gputraits.cuh"
 #include "gromacs/utility/gmxassert.h"
+#include "gromacs/utility/stringutil.h"
 
 /*! \brief
  * Allocates a device-side buffer.
@@ -139,7 +141,6 @@ void copyToDeviceBuffer(DeviceBuffer<ValueType>* buffer,
     }
 }
 
-
 /*! \brief
  * Performs the device-to-host data copy, synchronous or asynchronously on request.
  *
@@ -231,4 +232,90 @@ static bool checkDeviceBuffer(DeviceBuffer<T> buffer, gmx_unused int requiredSiz
     return buffer != nullptr;
 }
 
+//! Device texture wrapper.
+using DeviceTexture = cudaTextureObject_t;
+
+/*! \brief Create a texture object for an array of type ValueType.
+ *
+ * Creates the device buffer, copies data and binds texture object for an array of type ValueType.
+ *
+ * \todo Test if using textures is still relevant on modern hardware.
+ *
+ * \tparam      ValueType      Raw data type.
+ *
+ * \param[out]  deviceBuffer   Device buffer to store data in.
+ * \param[out]  deviceTexture  Device texture object to initialize.
+ * \param[in]   hostBuffer     Host buffer to get date from
+ * \param[in]   numValues      Number of elements in the buffer.
+ * \param[in]   deviceContext  GPU device context.
+ */
+template<typename ValueType>
+void initParamLookupTable(DeviceBuffer<ValueType>* deviceBuffer,
+                          DeviceTexture*           deviceTexture,
+                          const ValueType*         hostBuffer,
+                          int                      numValues,
+                          const DeviceContext&     deviceContext)
+{
+    if (numValues == 0)
+    {
+        return;
+    }
+    GMX_ASSERT(hostBuffer, "Host buffer should be specified.");
+
+    allocateDeviceBuffer(deviceBuffer, numValues, deviceContext);
+
+    const size_t sizeInBytes = numValues * sizeof(ValueType);
+
+    cudaError_t stat =
+            cudaMemcpy(*((ValueType**)deviceBuffer), hostBuffer, sizeInBytes, cudaMemcpyHostToDevice);
+
+    GMX_RELEASE_ASSERT(
+            stat == cudaSuccess,
+            gmx::formatString("Synchronous H2D copy failed (CUDA error: %s).", cudaGetErrorName(stat))
+                    .c_str());
+
+    if (!c_disableCudaTextures)
+    {
+        cudaResourceDesc rd;
+        cudaTextureDesc  td;
+
+        memset(&rd, 0, sizeof(rd));
+        rd.resType                = cudaResourceTypeLinear;
+        rd.res.linear.devPtr      = *deviceBuffer;
+        rd.res.linear.desc        = cudaCreateChannelDesc<ValueType>();
+        rd.res.linear.sizeInBytes = sizeInBytes;
+
+        memset(&td, 0, sizeof(td));
+        td.readMode = cudaReadModeElementType;
+        stat        = cudaCreateTextureObject(deviceTexture, &rd, &td, nullptr);
+        GMX_RELEASE_ASSERT(stat == cudaSuccess,
+                           gmx::formatString("cudaCreateTextureObject failed (CUDA error: %s).",
+                                             cudaGetErrorName(stat))
+                                   .c_str());
+    }
+}
+
+/*! \brief Unbind the texture and release the CUDA texture object.
+ *
+ * \tparam         ValueType      Raw data type
+ *
+ * \param[in,out]  deviceBuffer   Device buffer to store data in.
+ * \param[in,out]  deviceTexture  Device texture object to unbind.
+ */
+template<typename ValueType>
+void destroyParamLookupTable(DeviceBuffer<ValueType>* deviceBuffer, DeviceTexture& deviceTexture)
+{
+    if (!c_disableCudaTextures && deviceTexture && deviceBuffer)
+    {
+        cudaError_t stat = cudaDestroyTextureObject(deviceTexture);
+        GMX_RELEASE_ASSERT(
+                stat == cudaSuccess,
+                gmx::formatString(
+                        "cudaDestroyTextureObject on texture object failed (CUDA error: %s).",
+                        cudaGetErrorName(stat))
+                        .c_str());
+    }
+    freeDeviceBuffer(deviceBuffer);
+}
+
 #endif