Increased block sizes on Fermi

2bce659f · Peter Eastman · 75f74def · 2bce659f · 2bce659f · 2bce659f
Commit 2bce659f authored Aug 11, 2010 by Peter Eastman
6 changed files
--- a/plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaElectrostatic.cu
+++ b/plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaElectrostatic.cu
@@ -784,7 +784,7 @@ void cudaComputeAmoebaElectrostatic( amoebaGpuContext amoebaGpu )
    if( threadsPerBlock == 0 ){
        unsigned int maxThreads;
        if (gpu->sm_version >= SM_20)
-            maxThreads = 256;
+            maxThreads = 384;
        else if (gpu->sm_version >= SM_12)
            maxThreads = 128;
        else

--- a/plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaElectrostatic.h
+++ b/plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaElectrostatic.h
@@ -28,7 +28,7 @@

 __global__ 
 #if (__CUDA_ARCH__ >= 200)
-__launch_bounds__(256, 1)
+__launch_bounds__(384, 1)
 #elif (__CUDA_ARCH__ >= 130)
 __launch_bounds__(128, 1)
 #else

--- a/plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaKirkwood.cu
+++ b/plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaKirkwood.cu
@@ -1886,7 +1886,7 @@ void kCalculateAmoebaKirkwood( amoebaGpuContext amoebaGpu )
    if( threadsPerBlock == 0 ){
        unsigned int maxThreads;
        if (gpu->sm_version >= SM_20)
-            maxThreads = 256;
+            maxThreads = 512;
        else if (gpu->sm_version >= SM_12)
            maxThreads = 128;
        else

--- a/plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaKirkwood.h
+++ b/plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaKirkwood.h
@@ -28,7 +28,7 @@

 __global__
 #if (__CUDA_ARCH__ >= 200)
-__launch_bounds__(256, 1)
+__launch_bounds__(512, 1)
 #elif (__CUDA_ARCH__ >= 130)
 __launch_bounds__(128, 1)
 #else

--- a/plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaKirkwoodEDiff.cu
+++ b/plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaKirkwoodEDiff.cu
@@ -1054,7 +1054,7 @@ void kCalculateAmoebaKirkwoodEDiff( amoebaGpuContext amoebaGpu )
    if( threadsPerBlock == 0 ){
        unsigned int maxThreads;
        if (gpu->sm_version >= SM_20)
-            maxThreads = 192;
+            maxThreads = 384;
        else if (gpu->sm_version >= SM_12)
            maxThreads = 96;
        else

--- a/plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaKirkwoodEDiff.h
+++ b/plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaKirkwoodEDiff.h
@@ -28,7 +28,7 @@

 __global__
 #if (__CUDA_ARCH__ >= 200)
-__launch_bounds__(192, 1)
+__launch_bounds__(384, 1)
 #elif (__CUDA_ARCH__ >= 130)
 __launch_bounds__(96, 1)
 #else