Fixed problems on compute level 1.2 GPUs

81cd9f5f · Peter Eastman · ff5f5d5a · 81cd9f5f · 81cd9f5f · 81cd9f5f
Commit 81cd9f5f authored Jan 18, 2011 by Peter Eastman
20 changed files
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -329,6 +329,7 @@ IF(CUDA_FOUND)
        SET(FLAGS "")
        # Note that cmake will insert semicolons between these item automatically...
        SET(FLAGS ${FLAGS} -gencode arch=compute_11,code=sm_11)
+        SET(FLAGS ${FLAGS} -gencode arch=compute_12,code=sm_12)
        SET(FLAGS ${FLAGS} -gencode arch=compute_13,code=sm_13)
        SET(FLAGS ${FLAGS} -gencode arch=compute_20,code=sm_20)
        SET(FLAGS ${FLAGS} -use_fast_math)

--- a/platforms/cuda/src/kernels/kBrownianUpdate.cu
+++ b/platforms/cuda/src/kernels/kBrownianUpdate.cu
@@ -54,7 +54,7 @@ void GetBrownianUpdateSim(gpuContext gpu)
 __global__ 
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(GF1XX_UPDATE_THREADS_PER_BLOCK, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(GT2XX_UPDATE_THREADS_PER_BLOCK, 1)
 #else
 __launch_bounds__(G8X_UPDATE_THREADS_PER_BLOCK, 1)
@@ -93,7 +93,7 @@ void kBrownianUpdatePart1(gpuContext gpu)
 __global__ 
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(GF1XX_UPDATE_THREADS_PER_BLOCK, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(GT2XX_UPDATE_THREADS_PER_BLOCK, 1)
 #else
 __launch_bounds__(G8X_UPDATE_THREADS_PER_BLOCK, 1)

--- a/platforms/cuda/src/kernels/kCCMA.cu
+++ b/platforms/cuda/src/kernels/kCCMA.cu
@@ -51,7 +51,7 @@ void GetCCMASim(gpuContext gpu)
 __global__ void
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(1024, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(512, 1)
 #else
 __launch_bounds__(256, 1)
@@ -76,7 +76,7 @@ kComputeCCMAConstraintDirections()
 __global__ void
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(1024, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(512, 1)
 #else
 __launch_bounds__(256, 1)
@@ -126,7 +126,7 @@ kComputeCCMAConstraintForces(float4* atomPositions, bool addOldPosition)
 __global__ void
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(1024, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(512, 1)
 #else
 __launch_bounds__(256, 1)
@@ -156,7 +156,7 @@ kMultiplyByCCMAConstraintMatrix()
 __global__ void
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(1024, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(512, 1)
 #else
 __launch_bounds__(256, 1)

--- a/platforms/cuda/src/kernels/kCalculateCDLJForces.h
+++ b/platforms/cuda/src/kernels/kCalculateCDLJForces.h
@@ -36,7 +36,7 @@
 __global__ 
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(GF1XX_NONBOND_THREADS_PER_BLOCK, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(GT2XX_NONBOND_THREADS_PER_BLOCK, 1)
 #else
 __launch_bounds__(G8X_NONBOND_THREADS_PER_BLOCK, 1)

--- a/platforms/cuda/src/kernels/kCalculateCDLJObcGbsaForces1.h
+++ b/platforms/cuda/src/kernels/kCalculateCDLJObcGbsaForces1.h
@@ -36,7 +36,7 @@
 __global__ 
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(GF1XX_NONBOND_THREADS_PER_BLOCK, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(GT2XX_NONBOND_THREADS_PER_BLOCK, 1)
 #else
 __launch_bounds__(G8X_NONBOND_THREADS_PER_BLOCK, 1)

--- a/platforms/cuda/src/kernels/kCalculateCMAPTorsionForces.cu
+++ b/platforms/cuda/src/kernels/kCalculateCMAPTorsionForces.cu
@@ -81,7 +81,7 @@ using namespace std;
 __global__
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(GF1XX_LOCALFORCES_THREADS_PER_BLOCK, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(GT2XX_LOCALFORCES_THREADS_PER_BLOCK, 1)
 #else
 __launch_bounds__(G8X_LOCALFORCES_THREADS_PER_BLOCK, 1)

--- a/platforms/cuda/src/kernels/kCalculateCustomAngleForces.cu
+++ b/platforms/cuda/src/kernels/kCalculateCustomAngleForces.cu
@@ -84,7 +84,7 @@ void SetCustomAngleGlobalParams(const vector<float>& paramValues)
 __global__ 
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(1024, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(512, 1)
 #else
 __launch_bounds__(256, 1)

--- a/platforms/cuda/src/kernels/kCalculateCustomBondForces.cu
+++ b/platforms/cuda/src/kernels/kCalculateCustomBondForces.cu
@@ -81,7 +81,7 @@ void SetCustomBondGlobalParams(const vector<float>& paramValues)
 __global__ 
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(1024, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(512, 1)
 #else
 __launch_bounds__(256, 1)

--- a/platforms/cuda/src/kernels/kCalculateCustomExternalForces.cu
+++ b/platforms/cuda/src/kernels/kCalculateCustomExternalForces.cu
@@ -85,7 +85,7 @@ void SetCustomExternalGlobalParams(const vector<float>& paramValues)
 __global__ 
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(1024, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(512, 1)
 #else
 __launch_bounds__(256, 1)

--- a/platforms/cuda/src/kernels/kCalculateCustomTorsionForces.cu
+++ b/platforms/cuda/src/kernels/kCalculateCustomTorsionForces.cu
@@ -121,7 +121,7 @@ void SetCustomTorsionGlobalParams(const vector<float>& paramValues)
 __global__ 
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(GF1XX_LOCALFORCES_THREADS_PER_BLOCK, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(GT2XX_LOCALFORCES_THREADS_PER_BLOCK, 1)
 #else
 __launch_bounds__(G8X_LOCALFORCES_THREADS_PER_BLOCK, 1)

--- a/platforms/cuda/src/kernels/kCalculateGBVIBornSum.h
+++ b/platforms/cuda/src/kernels/kCalculateGBVIBornSum.h
@@ -40,7 +40,7 @@
 __global__ 
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(GF1XX_NONBOND_THREADS_PER_BLOCK, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(GT2XX_NONBOND_THREADS_PER_BLOCK, 1)
 #else
 __launch_bounds__(G8X_NONBOND_THREADS_PER_BLOCK, 1)

--- a/platforms/cuda/src/kernels/kCalculateGBVIForces2.h
+++ b/platforms/cuda/src/kernels/kCalculateGBVIForces2.h
@@ -40,7 +40,7 @@
 __global__ void 
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(GF1XX_BORNFORCE2_THREADS_PER_BLOCK, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(GT2XX_BORNFORCE2_THREADS_PER_BLOCK, 1)
 #else
 __launch_bounds__(G8X_BORNFORCE2_THREADS_PER_BLOCK, 1)

--- a/platforms/cuda/src/kernels/kCalculateLocalForces.cu
+++ b/platforms/cuda/src/kernels/kCalculateLocalForces.cu
@@ -126,7 +126,7 @@ void GetCalculateLocalForcesSim(gpuContext gpu)
 __global__ 
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(GF1XX_LOCALFORCES_THREADS_PER_BLOCK, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(GT2XX_LOCALFORCES_THREADS_PER_BLOCK, 1)
 #else
 __launch_bounds__(G8X_LOCALFORCES_THREADS_PER_BLOCK, 1)

--- a/platforms/cuda/src/kernels/kCalculateObcGbsaBornSum.h
+++ b/platforms/cuda/src/kernels/kCalculateObcGbsaBornSum.h
@@ -33,7 +33,7 @@
 __global__ 
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(GF1XX_NONBOND_THREADS_PER_BLOCK, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(GT2XX_NONBOND_THREADS_PER_BLOCK, 1)
 #else
 __launch_bounds__(G8X_NONBOND_THREADS_PER_BLOCK, 1)

--- a/platforms/cuda/src/kernels/kCalculateObcGbsaForces2.h
+++ b/platforms/cuda/src/kernels/kCalculateObcGbsaForces2.h
@@ -33,7 +33,7 @@
 __global__ 
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(GF1XX_BORNFORCE2_THREADS_PER_BLOCK, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(GT2XX_BORNFORCE2_THREADS_PER_BLOCK, 1)
 #else
 __launch_bounds__(G8X_BORNFORCE2_THREADS_PER_BLOCK, 1)

--- a/platforms/cuda/src/kernels/kCalculatePME.cu
+++ b/platforms/cuda/src/kernels/kCalculatePME.cu
@@ -95,7 +95,7 @@ inline __host__ __device__ float4 make_float4(int3 a)
 __global__ 
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(1024, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(512, 1)
 #else
 __launch_bounds__(256, 1)
@@ -183,7 +183,7 @@ void kUpdateBsplines_kernel()
 __global__
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(1024, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(512, 1)
 #else
 __launch_bounds__(256, 1)
@@ -284,7 +284,7 @@ void kGridSpreadCharge_kernel()
 __global__ 
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(1024, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(512, 1)
 #else
 __launch_bounds__(256, 1)
@@ -325,7 +325,7 @@ void kReciprocalConvolution_kernel()
 __global__ 
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(1024, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(512, 1)
 #else
 __launch_bounds__(256, 1)

--- a/platforms/cuda/src/kernels/kForces.cu
+++ b/platforms/cuda/src/kernels/kForces.cu
@@ -120,7 +120,7 @@ void kClearEnergy(gpuContext gpu)
 __global__ 
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(GF1XX_THREADS_PER_BLOCK, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(GT2XX_THREADS_PER_BLOCK, 1)
 #else
 __launch_bounds__(G8X_THREADS_PER_BLOCK, 1)
@@ -231,7 +231,7 @@ void kReduceBornSumAndForces(gpuContext gpu)
 __global__ 
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(GF1XX_THREADS_PER_BLOCK, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(GT2XX_THREADS_PER_BLOCK, 1)
 #else
 __launch_bounds__(G8X_THREADS_PER_BLOCK, 1)
@@ -301,7 +301,7 @@ double kReduceEnergy(gpuContext gpu)
 __global__ 
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(GF1XX_UPDATE_THREADS_PER_BLOCK, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(GT2XX_UPDATE_THREADS_PER_BLOCK, 1)
 #else
 __launch_bounds__(G8X_UPDATE_THREADS_PER_BLOCK, 1)
@@ -367,7 +367,7 @@ void kReduceObcGbsaBornForces_kernel()
 __global__ 
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(GF1XX_UPDATE_THREADS_PER_BLOCK, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(GT2XX_UPDATE_THREADS_PER_BLOCK, 1)
 #else
 __launch_bounds__(G8X_UPDATE_THREADS_PER_BLOCK, 1)

--- a/platforms/cuda/src/kernels/kLangevinUpdate.cu
+++ b/platforms/cuda/src/kernels/kLangevinUpdate.cu
@@ -104,7 +104,7 @@ void kLangevinUpdatePart2(gpuContext gpu)
 __global__ 
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(GF1XX_UPDATE_THREADS_PER_BLOCK, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(GT2XX_UPDATE_THREADS_PER_BLOCK, 1)
 #else
 __launch_bounds__(G8X_UPDATE_THREADS_PER_BLOCK, 1)
@@ -173,7 +173,7 @@ void kSelectLangevinStepSize(gpuContext gpu, float maxTimeStep)
 __global__
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(GF1XX_UPDATE_THREADS_PER_BLOCK, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(GT2XX_UPDATE_THREADS_PER_BLOCK, 1)
 #else
 __launch_bounds__(G8X_UPDATE_THREADS_PER_BLOCK, 1)

--- a/platforms/cuda/src/kernels/kLangevinUpdate.h
+++ b/platforms/cuda/src/kernels/kLangevinUpdate.h
@@ -35,7 +35,7 @@
 __global__
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(GF1XX_UPDATE_THREADS_PER_BLOCK, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(GT2XX_UPDATE_THREADS_PER_BLOCK, 1)
 #else
 __launch_bounds__(G8X_UPDATE_THREADS_PER_BLOCK, 1)
@@ -113,7 +113,7 @@ void kLangevinUpdatePart1_kernel()
 __global__
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(GF1XX_UPDATE_THREADS_PER_BLOCK, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(GT2XX_UPDATE_THREADS_PER_BLOCK, 1)
 #else
 __launch_bounds__(G8X_UPDATE_THREADS_PER_BLOCK, 1)

--- a/platforms/cuda/src/kernels/kRandom.cu
+++ b/platforms/cuda/src/kernels/kRandom.cu
@@ -57,7 +57,7 @@ extern __shared__ float3 sRand[];
 __global__ 
 #if (__CUDA_ARCH__ >= 200)
 __launch_bounds__(GF1XX_RANDOM_THREADS_PER_BLOCK, 1)
-#elif (__CUDA_ARCH__ >= 130)
+#elif (__CUDA_ARCH__ >= 120)
 __launch_bounds__(GT2XX_RANDOM_THREADS_PER_BLOCK, 1)
 #else
 __launch_bounds__(G8X_RANDOM_THREADS_PER_BLOCK, 1)