Removed support for NVIDIA CC 2.x devices (codename Fermi)

[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 3840c48a7ebfd8d1ceb3a69add50b0f7cb3773d3..e9809c8c0d2acc43e437eb1e2fea33dbfec1adbb 100644 (file)
--- a/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda_kernel_utils.cuh
+++ b/src/gromacs/mdlib/nbnxn_cuda/nbnxn_cuda_kernel_utils.cuh
@@ -507,9 +507,8 @@ void reduce_force_j_generic(float *f_buf, float3 *fout,
 }
 
 /*! Final j-force reduction; this implementation only with power of two
- *  array sizes and with sm >= 3.0
+ *  array sizes.
  */
-#if GMX_PTX_ARCH >= 300 || GMX_PTX_ARCH == 0
 static __forceinline__ __device__
 void reduce_force_j_warp_shfl(float3 f, float3 *fout,
                               int tidxi, int aidx,
@@ -539,7 +538,6 @@ void reduce_force_j_warp_shfl(float3 f, float3 *fout,
         atomicAdd((&fout[aidx].x) + tidxi, f.x);
     }
 }
-#endif
 
 /*! Final i-force reduction; this generic implementation works with
  *  arbitrary array sizes.
@@ -634,9 +632,8 @@ void reduce_force_i(float *f_buf, float3 *f,
 }
 
 /*! Final i-force reduction; this implementation works only with power of two
- *  array sizes and with sm >= 3.0
+ *  array sizes.
  */
-#if GMX_PTX_ARCH >= 300 || GMX_PTX_ARCH == 0
 static __forceinline__ __device__
 void reduce_force_i_warp_shfl(float3 fin, float3 *fout,
                               float *fshift_buf, bool bCalcFshift,
@@ -671,7 +668,6 @@ void reduce_force_i_warp_shfl(float3 fin, float3 *fout,
         }
     }
 }
-#endif
 
 /*! Energy reduction; this implementation works only with power of two
  *  array sizes.
@@ -711,9 +707,8 @@ void reduce_energy_pow2(volatile float *buf,
 }
 
 /*! Energy reduction; this implementation works only with power of two
- *  array sizes and with sm >= 3.0
+ *  array sizes.
  */
-#if GMX_PTX_ARCH >= 300 || GMX_PTX_ARCH == 0
 static __forceinline__ __device__
 void reduce_energy_warp_shfl(float E_lj, float E_el,
                              float *e_lj, float *e_el,
@@ -738,6 +733,5 @@ void reduce_energy_warp_shfl(float E_lj, float E_el,
         atomicAdd(e_el, E_el);
     }
 }
-#endif /* GMX_PTX_ARCH */
 
 #endif /* NBNXN_CUDA_KERNEL_UTILS_CUH */