creating branch for adding support for CUDA 3.x and Fermi

[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6829]

1 files changed, 290 insertions, 0 deletions

diff --git a/benchmarks/CUDA/LPS/laplace3d.cu b/benchmarks/CUDA/LPS/laplace3d.cu
new file mode 100644
index 0000000..e7564dd
--- /dev/null
+++ b/benchmarks/CUDA/LPS/laplace3d.cu
@@ -0,0 +1,290 @@
+//
+// Program to solve Laplace equation on a regular 3D grid
+//
+
+#include <stdlib.h>
+#include <stdio.h>
+#include <string.h>
+#include <math.h>
+#include <cutil.h>
+
+////////////////////////////////////////////////////////////////////////
+// define kernel block size
+////////////////////////////////////////////////////////////////////////
+
+#define BLOCK_X 32
+#define BLOCK_Y 4
+
+////////////////////////////////////////////////////////////////////////
+// include kernel function
+////////////////////////////////////////////////////////////////////////
+
+#include <laplace3d_kernel.cu>
+
+////////////////////////////////////////////////////////////////////////
+// declaration, forward
+////////////////////////////////////////////////////////////////////////
+
+extern "C" 
+void Gold_laplace3d(int NX, int NY, int NZ, float* h_u1, float* h_u2);
+
+void printHelp(void);
+
+////////////////////////////////////////////////////////////////////////
+// Main program
+////////////////////////////////////////////////////////////////////////
+
+int main(int argc, char **argv){
+
+  // 'h_' prefix - CPU (host) memory space
+
+  int    NX, NY, NZ, REPEAT, bx, by, i, j, k, ind, pitch;
+  size_t pitch_bytes;
+  float  *h_u1, *h_u2, *h_u3, *h_foo, err;
+
+  unsigned int   hTimer;
+
+  // 'd_' prefix - GPU (device) memory space
+
+  float  *d_u1, *d_u2, *d_foo;
+
+  // check command line inputs
+
+  if(cutCheckCmdLineFlag( argc, (const char**)argv, "help")) {
+    printHelp();
+    return 1;
+  }
+
+  if( cutGetCmdLineArgumenti( argc, (const char**)argv, "nx", &NX) ) {
+    if( NX <= 99 ) {
+      printf("Illegal argument - nx must be greater than 99\n");
+      return -1;
+    }
+  }
+  else
+    NX = 100;
+
+  if( cutGetCmdLineArgumenti( argc, (const char**)argv, "ny", &NY) ) {
+    if( NY <= 99 ) {
+      printf("Illegal argument - ny must be greater than 99\n");
+      return -1;
+    }
+  }
+  else
+    NY = 100;
+
+  if( cutGetCmdLineArgumenti( argc, (const char**)argv, "nz", &NZ) ) {
+    if( NZ <= 99 ) {
+      printf("Illegal argument - nz must be greater than 99\n");
+      return -1;
+    }
+  }
+  else
+    NZ = 100;
+
+  if( cutGetCmdLineArgumenti( argc, (const char**)argv, "repeat", &REPEAT) ) {
+    if( REPEAT <= 0 ) {
+      printf("Illegal argument - repeat must be greater than zero\n");
+      return -1;
+    }
+  }
+  else
+    REPEAT = 1;
+
+  printf("\nGrid dimensions: %d x %d x %d\n", NX, NY, NZ);
+
+  // initialise card and timer
+  int deviceCount;                                                         
+  CUDA_SAFE_CALL_NO_SYNC(cudaGetDeviceCount(&deviceCount));                
+  if (deviceCount == 0) {                                                  
+      fprintf(stderr, "There is no device.\n");                            
+      exit(EXIT_FAILURE);                                                  
+  }                                                                        
+  int dev;                                                                 
+  for (dev = 0; dev < deviceCount; ++dev) {                                
+      cudaDeviceProp deviceProp;                                           
+      CUDA_SAFE_CALL_NO_SYNC(cudaGetDeviceProperties(&deviceProp, dev));   
+      if (deviceProp.major >= 1)                                           
+          break;                                                           
+  }                                                                        
+  if (dev == deviceCount) {                                                
+      fprintf(stderr, "There is no device supporting CUDA.\n");            
+      exit(EXIT_FAILURE);                                                  
+  }                                                                        
+  else                                                                     
+      CUDA_SAFE_CALL(cudaSetDevice(dev));  
+  CUT_SAFE_CALL( cutCreateTimer(&hTimer) );
+ 
+  // allocate memory for arrays
+
+  h_u1 = (float *)malloc(sizeof(float)*NX*NY*NZ);
+  h_u2 = (float *)malloc(sizeof(float)*NX*NY*NZ);
+  h_u3 = (float *)malloc(sizeof(float)*NX*NY*NZ);
+  CUDA_SAFE_CALL( cudaMallocPitch((void **)&d_u1, &pitch_bytes, sizeof(float)*NX, NY*NZ) );
+  CUDA_SAFE_CALL( cudaMallocPitch((void **)&d_u2, &pitch_bytes, sizeof(float)*NX, NY*NZ) );
+
+  pitch = pitch_bytes/sizeof(float);
+
+  // initialise u1
+    
+  for (k=0; k<NZ; k++) {
+    for (j=0; j<NY; j++) {
+      for (i=0; i<NX; i++) {
+        ind = i + j*NX + k*NX*NY;
+
+        if (i==0 || i==NX-1 || j==0 || j==NY-1|| k==0 || k==NZ-1)
+          h_u1[ind] = 1.0f;           // Dirichlet b.c.'s
+        else
+          h_u1[ind] = 0.0f;
+      }
+    }
+  }
+
+  // copy u1 to device
+
+  CUT_SAFE_CALL(cutStartTimer(hTimer));
+  CUDA_SAFE_CALL( cudaMemcpy2D(d_u1, pitch_bytes,
+                               h_u1, sizeof(float)*NX,
+                               sizeof(float)*NX, NY*NZ,
+                               cudaMemcpyHostToDevice) );
+  CUDA_SAFE_CALL( cudaThreadSynchronize() );
+  CUT_SAFE_CALL(cutStopTimer(hTimer));
+  printf("\nCopy u1 to device: %f (ms) \n", cutGetTimerValue(hTimer));
+  CUT_SAFE_CALL( cutResetTimer(hTimer) );
+
+  // Set up the execution configuration
+
+  bx = 1 + (NX-1)/BLOCK_X;
+  by = 1 + (NY-1)/BLOCK_Y;
+
+  dim3 dimGrid(bx,by);
+  dim3 dimBlock(BLOCK_X,BLOCK_Y);
+
+  printf("\n dimGrid  = %d %d %d \n",dimGrid.x,dimGrid.y,dimGrid.z);
+  printf(" dimBlock = %d %d %d \n",dimBlock.x,dimBlock.y,dimBlock.z);
+
+  // Execute GPU kernel
+
+  CUDA_SAFE_CALL( cudaThreadSynchronize() );
+  CUT_SAFE_CALL( cutResetTimer(hTimer) );
+  CUT_SAFE_CALL( cutStartTimer(hTimer) );
+
+  for (i = 1; i <= REPEAT; ++i) {
+    GPU_laplace3d<<<dimGrid, dimBlock>>>(NX, NY, NZ, pitch, d_u1, d_u2);
+    d_foo = d_u1; d_u1 = d_u2; d_u2 = d_foo;   // swap d_u1 and d_u3
+
+    CUDA_SAFE_CALL( cudaThreadSynchronize() );
+    CUT_CHECK_ERROR("GPU_laplace3d execution failed\n");
+  }
+
+  CUT_SAFE_CALL( cutStopTimer(hTimer) );
+  printf("\n%dx GPU_laplace3d: %f (ms) \n", REPEAT, cutGetTimerValue(hTimer));
+
+  CUT_SAFE_CALL( cutResetTimer(hTimer) );
+
+  // Read back GPU results
+
+  CUT_SAFE_CALL( cutStartTimer(hTimer) );
+  CUDA_SAFE_CALL( cudaMemcpy2D(h_u2, sizeof(float)*NX,
+                               d_u1, pitch_bytes,
+                               sizeof(float)*NX, NY*NZ,
+                               cudaMemcpyDeviceToHost) );
+  CUT_SAFE_CALL( cutStopTimer(hTimer) );
+  printf("\nCopy u2 to host: %f (ms) \n", cutGetTimerValue(hTimer));
+  CUT_SAFE_CALL( cutResetTimer(hTimer) );
+
+
+  // print out corner of array
+
+  /*
+  for (k=0; k<3; k++) {
+    for (j=0; j<8; j++) {
+      for (i=0; i<8; i++) {
+        ind = i + j*NX + k*NX*NY;
+        printf(" %5.2f ", h_u2[ind]);
+      }
+      printf("\n");
+    }
+    printf("\n");
+  }
+  */
+
+  // Gold treatment
+
+  CUT_SAFE_CALL( cutResetTimer(hTimer) );
+  CUT_SAFE_CALL( cutStartTimer(hTimer) );
+
+  for (int i = 1; i <= REPEAT; ++i) {
+    Gold_laplace3d(NX, NY, NZ, h_u1, h_u3);
+    h_foo = h_u1; h_u1 = h_u3; h_u3 = h_foo;   // swap h_u1 and h_u3
+  }
+
+  CUT_SAFE_CALL( cutStopTimer(hTimer) );
+  printf("\n%dx Gold_laplace3d: %f (ms) \n \n", REPEAT, cutGetTimerValue(hTimer));
+
+  // print out corner of array
+
+  /*
+  for (k=0; k<3; k++) {
+    for (j=0; j<8; j++) {
+      for (i=0; i<8; i++) {
+        ind = i + j*NX + k*NX*NY;
+        printf(" %5.2f ", h_u1[ind]);
+      }
+      printf("\n");
+    }
+    printf("\n");
+  }
+  */
+
+  // error check
+
+  err = 0.0;
+
+  for (k=0; k<NZ; k++) {
+    for (j=0; j<NY; j++) {
+      for (i=0; i<NX; i++) {
+        ind = i + j*NX + k*NX*NY;
+        err += (h_u1[ind]-h_u2[ind])*(h_u1[ind]-h_u2[ind]);
+      }
+    }
+  }
+
+  printf("\n rms error = %f \n",sqrt(err/ (float)(NX*NY*NZ)));
+
+ // Release GPU and CPU memory
+  printf("CUDA_SAFE_CALL( cudaFree(d_u1) );\n"); fflush(stdout);
+  CUDA_SAFE_CALL( cudaFree(d_u1) );
+  printf("CUDA_SAFE_CALL( cudaFree(d_u2) );\n"); fflush(stdout);
+  CUDA_SAFE_CALL( cudaFree(d_u2) );
+  printf("free(h_u1);\n"); fflush(stdout);
+  free(h_u1);
+  printf("free(h_u2);\n"); fflush(stdout);
+  free(h_u2);
+  printf("free(h_u3);\n"); fflush(stdout);
+  free(h_u3);
+
+  CUT_SAFE_CALL( cutDeleteTimer(hTimer) );
+  CUT_EXIT(argc, argv);
+}
+
+
+///////////////////////////////////////////////////////////////////////////
+//Print help screen
+///////////////////////////////////////////////////////////////////////////
+void printHelp(void)
+{
+  printf("Usage:  laplace3d [OPTION]...\n");
+  printf("6-point stencil 3D Laplace test \n");
+  printf("\n");
+  printf("Example: run 100 iterations on a 256x128x128 grid\n");
+  printf("./laplace3d --nx=256 --ny=128 --nz=128 --repeat=100\n");
+
+  printf("\n");
+  printf("Options:\n");
+  printf("--help\t\t\tDisplay this help menu\n");
+  printf("--nx=[SIZE]\t\tGrid width\n");
+  printf("--ny=[SIZE]\t\tGrid height\n");
+  printf("--nz=[SIZE]\t\tGrid depth\n");
+  printf("--repeat=[COUNT]\tNumber of repetitions\n");
+}