Merge branch release-4-6 into master

[alexxy/gromacs.git] / src / gromacs / mdlib / nbnxn_cuda / nbnxn_cuda_kernel_utils.cuh

diff --git a/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda_kernel_utils.cuh b/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda_kernel_utils.cuh
index de0acc0efc313987134ad95fc6e3bb52af2e905e..ce219a2a82bc6a31f212b5fbb23ec7063b3cd324 100644 (file)
--- a/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda_kernel_utils.cuh
+++ b/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda_kernel_utils.cuh
@@ -62,10 +62,26 @@ float interpolate_coulomb_force_r(float r, float scale)
     float   fract2 = normalized - index;
     float   fract1 = 1.0f - fract2;
 
-    return  fract1 * tex1Dfetch(tex_coulomb_tab, index)
-            + fract2 * tex1Dfetch(tex_coulomb_tab, index + 1);
+    return  fract1 * tex1Dfetch(coulomb_tab_texref, index)
+            + fract2 * tex1Dfetch(coulomb_tab_texref, index + 1);
 }
 
+#ifdef TEXOBJ_SUPPORTED
+static inline __device__
+float interpolate_coulomb_force_r(cudaTextureObject_t texobj_coulomb_tab,
+                                  float r, float scale)
+{
+    float   normalized = scale * r;
+    int     index = (int) normalized;
+    float   fract2 = normalized - index;
+    float   fract1 = 1.0f - fract2;
+
+    return  fract1 * tex1Dfetch<float>(texobj_coulomb_tab, index) +
+            fract2 * tex1Dfetch<float>(texobj_coulomb_tab, index + 1);
+}
+#endif
+
+
 /*! Calculate analytical Ewald correction term. */
 static inline __device__
 float pmecorrF(float z2)