changed constant from K to KMAX
authorLawrence Cayton <lcayton@tuebingen.mpg.de>
Mon, 2 May 2011 16:19:47 +0000 (18:19 +0200)
committerLawrence Cayton <lcayton@tuebingen.mpg.de>
Mon, 2 May 2011 16:19:47 +0000 (18:19 +0200)
brute.cu
defs.h
driver.cu
kernels.cu
rbc.cu

index f40d70a..c152ad5 100644 (file)
--- a/brute.cu
+++ b/brute.cu
@@ -75,8 +75,8 @@ void bruteK(matrix x, matrix q, intMatrix NNs, matrix NNdists){
   
   dx.r=x.r; dx.pr=x.pr; dx.c=x.c; dx.pc=x.pc; dx.ld=x.ld;
   dq.r=q.r; dq.pr=q.pr; dq.c=q.c; dq.pc=q.pc; dq.ld=q.ld;
-  dNNdists.r=q.r; dNNdists.pr=q.pr; dNNdists.c=K; dNNdists.pc=K; dNNdists.ld=dNNdists.pc;
-  dMinIDs.r=q.r; dMinIDs.pr=q.pr; dMinIDs.c=K; dMinIDs.pc=K; dMinIDs.ld=dMinIDs.pc;
+  dNNdists.r=q.r; dNNdists.pr=q.pr; dNNdists.c=KMAX; dNNdists.pc=KMAX; dNNdists.ld=dNNdists.pc;
+  dMinIDs.r=q.r; dMinIDs.pr=q.pr; dMinIDs.c=KMAX; dMinIDs.pc=KMAX; dMinIDs.ld=dMinIDs.pc;
 
   checkErr( cudaMalloc((void**)&dNNdists.mat, dNNdists.pc*dNNdists.pr*sizeof(*dNNdists.mat)) );
   checkErr( cudaMalloc((void**)&dMinIDs.mat, dMinIDs.pc*dMinIDs.pr*sizeof(*dMinIDs.mat)) );
@@ -143,7 +143,7 @@ void bruteCPU(matrix X, matrix Q, unint *NNs){
 /*     for( j=0; j<x.r; j++) */
 /*       d[i][j] = distVec( q, x, i, j ); */
 /*     gsl_sort_float_index(t[i], d[i], 1, x.r); */
-/*     for ( j=0; j<K; j++) */
+/*     for ( j=0; j<KMAX; j++) */
 /*       NNs.mat[IDX( i, j, NNs.ld )] = t[i][j]; */
 /*   } */
 
diff --git a/defs.h b/defs.h
index 53e87d9..91bd4b3 100644 (file)
--- a/defs.h
+++ b/defs.h
@@ -6,7 +6,7 @@
 
 #define FLOAT_TOL 1e-7
 
-#define K 32 //for k-NN.  Do not change!
+#define KMAX 32 //for k-NN.  Do not change!
 #define BLOCK_SIZE 16 //must be a power of 2 (current 
 // implementation of findRange requires a power of 4, in fact)
 
index c6b606c..50ce22c 100644 (file)
--- a/driver.cu
+++ b/driver.cu
@@ -40,11 +40,11 @@ int main(int argc, char**argv){
   parseInput(argc,argv);
   
   gettimeofday( &tvB, NULL );
-  printf("Using GPU #%d\n",deviceNum);
-  if(cudaSetDevice(deviceNum) != cudaSuccess){
-    printf("Unable to select device %d.. exiting. \n",deviceNum);
-    exit(1);
-  }
+  printf("Using GPU #%d\n",deviceNum); 
+  if(cudaSetDevice(deviceNum) != cudaSuccess){ 
+    printf("Unable to select device %d.. exiting. \n",deviceNum); 
+    exit(1); 
+  } 
   
   size_t memFree, memTot;
   cudaMemGetInfo(&memFree, &memTot);
@@ -71,8 +71,8 @@ int main(int argc, char**argv){
 
 
   //Allocate space for NNs and dists
-  initIntMat( &nnsRBC, m, K );  //K is defined in defs.h
-  initMat( &distsRBC, m, K );
+  initIntMat( &nnsRBC, m, KMAX );  //KMAX is defined in defs.h
+  initMat( &distsRBC, m, KMAX );
   nnsRBC.mat = (unint*)calloc( sizeOfIntMat(nnsRBC), sizeof(*nnsRBC.mat) );
   distsRBC.mat = (real*)calloc( sizeOfMat(distsRBC), sizeof(*distsRBC.mat) );
 
@@ -92,9 +92,9 @@ int main(int argc, char**argv){
   if( runBrute ){
     intMatrix nnsBrute;
     matrix distsBrute;
-    initIntMat( &nnsBrute, m, K );
+    initIntMat( &nnsBrute, m, KMAX );
     nnsBrute.mat = (unint*)calloc( sizeOfIntMat(nnsBrute), sizeof(*nnsBrute.mat) );
-    initMat( &distsBrute, m, K );
+    initMat( &distsBrute, m, KMAX );
     distsBrute.mat = (real*)calloc( sizeOfMat(distsBrute), sizeof(*distsBrute.mat) );
     
     printf("running k-brute force..\n");
@@ -300,7 +300,7 @@ void evalKNNerror(matrix x, matrix q, intMatrix NNs){
   NNsB.r=q.r; NNsB.pr=q.pr; NNsB.c=NNsB.pc=32; NNsB.ld=NNsB.pc;
   NNsB.mat = (unint*)calloc( NNsB.pr*NNsB.pc, sizeof(*NNsB.mat) );
   matrix distsBrute;
-  distsBrute.r=q.r; distsBrute.pr=q.pr; distsBrute.c=distsBrute.pc=K; distsBrute.ld=distsBrute.pc;
+  distsBrute.r=q.r; distsBrute.pr=q.pr; distsBrute.c=distsBrute.pc=KMAX; distsBrute.ld=distsBrute.pc;
   distsBrute.mat = (real*)calloc( distsBrute.pr*distsBrute.pc, sizeof(*distsBrute.mat) );
 
   gettimeofday(&tvB,NULL);
@@ -309,8 +309,8 @@ void evalKNNerror(matrix x, matrix q, intMatrix NNs){
 
    //calc overlap
   for(i=0; i<m; i++){
-    for(j=0; j<K; j++){
-      for(k=0; k<K; k++){
+    for(j=0; j<KMAX; j++){
+      for(k=0; k<KMAX; k++){
        ol[i] += ( NNs.mat[IDX(i, j, NNs.ld)] == NNsB.mat[IDX(i, k, NNsB.ld)] );
       }
     }
@@ -326,7 +326,7 @@ void evalKNNerror(matrix x, matrix q, intMatrix NNs){
   for(i=0;i<m;i++) {
     var += (((double)ol[i])-mean)*(((double)ol[i])-mean)/((double)m);
   }
-  printf("\tavg overlap = %6.4f/%d; std dev = %6.4f \n", mean, K, sqrt(var));
+  printf("\tavg overlap = %6.4f/%d; std dev = %6.4f \n", mean, KMAX, sqrt(var));
 
   FILE* fp;
   if(outFile){
@@ -336,7 +336,7 @@ void evalKNNerror(matrix x, matrix q, intMatrix NNs){
 
   real *ranges = (real*)calloc(q.pr,sizeof(*ranges));
   for(i=0;i<q.r;i++){
-    ranges[i] = distVec(q,x,i,NNs.mat[IDX(i, K-1, NNs.ld)]);
+    ranges[i] = distVec(q,x,i,NNs.mat[IDX(i, KMAX-1, NNs.ld)]);
   }
   
   unint *cnts = (unint*)calloc(q.pr,sizeof(*cnts));
@@ -378,13 +378,13 @@ void writeNeighbs(char *file, char *filetxt, intMatrix NNs, matrix dNNs){
     }
     
     for( i=0; i<m; i++ ){
-      for( j=0; j<K; j++ )
+      for( j=0; j<KMAX; j++ )
        fprintf( fp, "%u ", NNs.mat[IDX( i, j, NNs.ld )] );
       fprintf(fp, "\n");
     }
     
     for( i=0; i<m; i++ ){
-      for( j=0; j<K; j++ )
+      for( j=0; j<KMAX; j++ )
        fprintf( fp, "%f ", dNNs.mat[IDX( i, j, dNNs.ld )]); 
       fprintf(fp, "\n");
     }
@@ -401,9 +401,9 @@ void writeNeighbs(char *file, char *filetxt, intMatrix NNs, matrix dNNs){
     }
     
     for( i=0; i<m; i++ )
-      fwrite( &NNs.mat[IDX( i, 0, NNs.ld )], sizeof(*NNs.mat), K, fp );
+      fwrite( &NNs.mat[IDX( i, 0, NNs.ld )], sizeof(*NNs.mat), KMAX, fp );
     for( i=0; i<m; i++ )
-      fwrite( &dNNs.mat[IDX( i, 0, dNNs.ld )], sizeof(*dNNs.mat), K, fp );
+      fwrite( &dNNs.mat[IDX( i, 0, dNNs.ld )], sizeof(*dNNs.mat), KMAX, fp );
     
     fclose(fp);
   }
index bdef21e..e6f6adf 100644 (file)
@@ -99,8 +99,8 @@ __global__ void planKNNKernel(const matrix Q, const unint *qMap, const matrix X,
   unint i,j,k;
   unint groupIts;
   
-  __shared__ real dNN[BLOCK_SIZE][K+BLOCK_SIZE];
-  __shared__ unint idNN[BLOCK_SIZE][K+BLOCK_SIZE];
+  __shared__ real dNN[BLOCK_SIZE][KMAX+BLOCK_SIZE];
+  __shared__ unint idNN[BLOCK_SIZE][KMAX+BLOCK_SIZE];
 
   __shared__ real Xs[BLOCK_SIZE][BLOCK_SIZE];
   __shared__ real Qs[BLOCK_SIZE][BLOCK_SIZE];
@@ -238,8 +238,8 @@ __global__ void knnKernel(const matrix Q, unint numDone, const matrix X, matrix
   __shared__ real Xs[BLOCK_SIZE][BLOCK_SIZE];
   __shared__ real Qs[BLOCK_SIZE][BLOCK_SIZE];
   
-  __shared__ real dNN[BLOCK_SIZE][K+BLOCK_SIZE];
-  __shared__ unint idNN[BLOCK_SIZE][K+BLOCK_SIZE];
+  __shared__ real dNN[BLOCK_SIZE][KMAX+BLOCK_SIZE];
+  __shared__ unint idNN[BLOCK_SIZE][KMAX+BLOCK_SIZE];
 
   dNN[offQ][offX] = MAX_REAL;
   dNN[offQ][offX+16] = MAX_REAL;
@@ -572,45 +572,45 @@ __device__ void sort16off(real x[][48], unint xi[][48]){
   int j = threadIdx.y;
 
   if(i%2==0)
-    mmGateI( x[j]+K+i, x[j]+K+i+1, xi[j]+K+i, xi[j]+K+i+1 );
+    mmGateI( x[j]+KMAX+i, x[j]+KMAX+i+1, xi[j]+KMAX+i, xi[j]+KMAX+i+1 );
   __syncthreads();
 
   if(i%4<2)
-    mmGateI( x[j]+K+i, x[j]+K+i+2, xi[j]+K+i, xi[j]+K+i+2 );
+    mmGateI( x[j]+KMAX+i, x[j]+KMAX+i+2, xi[j]+KMAX+i, xi[j]+KMAX+i+2 );
   __syncthreads();
 
   if(i%4==1)
-    mmGateI( x[j]+K+i, x[j]+K+i+1, xi[j]+K+i, xi[j]+K+i+1 );
+    mmGateI( x[j]+KMAX+i, x[j]+KMAX+i+1, xi[j]+KMAX+i, xi[j]+KMAX+i+1 );
   __syncthreads();
   
   if(i%8<4)
-    mmGateI( x[j]+K+i, x[j]+K+i+4, xi[j]+K+i, xi[j]+K+i+4 );
+    mmGateI( x[j]+KMAX+i, x[j]+KMAX+i+4, xi[j]+KMAX+i, xi[j]+KMAX+i+4 );
   __syncthreads();
   
   if(i%8==2 || i%8==3)
-    mmGateI( x[j]+K+i, x[j]+K+i+2, xi[j]+K+i, xi[j]+K+i+2 );
+    mmGateI( x[j]+KMAX+i, x[j]+KMAX+i+2, xi[j]+KMAX+i, xi[j]+KMAX+i+2 );
   __syncthreads();
 
   if( i%2 && i%8 != 7 ) 
-    mmGateI( x[j]+K+i, x[j]+K+i+1, xi[j]+K+i, xi[j]+K+i+1 );
+    mmGateI( x[j]+KMAX+i, x[j]+KMAX+i+1, xi[j]+KMAX+i, xi[j]+KMAX+i+1 );
   __syncthreads();
   
   //0-7; 8-15 now sorted.  merge time.
   if( i<8)
-    mmGateI( x[j]+K+i, x[j]+K+i+8, xi[j]+K+i, xi[j]+K+i+8 );
+    mmGateI( x[j]+KMAX+i, x[j]+KMAX+i+8, xi[j]+KMAX+i, xi[j]+KMAX+i+8 );
   __syncthreads();
   
   if( i>3 && i<8 )
-    mmGateI( x[j]+K+i, x[j]+K+i+4, xi[j]+K+i, xi[j]+K+i+4 );
+    mmGateI( x[j]+KMAX+i, x[j]+KMAX+i+4, xi[j]+KMAX+i, xi[j]+KMAX+i+4 );
   __syncthreads();
   
   int os = (i/2)*4+2 + i%2;
   if(i<6)
-    mmGateI( x[j]+K+os, x[j]+K+os+2, xi[j]+K+os, xi[j]+K+os+2 );
+    mmGateI( x[j]+KMAX+os, x[j]+KMAX+os+2, xi[j]+KMAX+os, xi[j]+KMAX+os+2 );
   __syncthreads();
   
   if( i%2 && i<15)
-    mmGateI( x[j]+K+i, x[j]+K+i+1, xi[j]+K+i, xi[j]+K+i+1 );
+    mmGateI( x[j]+KMAX+i, x[j]+KMAX+i+1, xi[j]+KMAX+i, xi[j]+KMAX+i+1 );
 }
 
 //min-max gate: it sets the minimum of x and y into x, the maximum into y, and 
diff --git a/rbc.cu b/rbc.cu
index c0494e6..32f389e 100644 (file)
--- a/rbc.cu
+++ b/rbc.cu
@@ -331,15 +331,15 @@ void computeNNs(matrix dx, intMatrix dxMap, matrix dq, unint *dqMap, compPlan dc
 void computeKNNs(matrix dx, intMatrix dxMap, matrix dq, unint *dqMap, compPlan dcP, intMatrix NNs, matrix NNdists, unint compLength){
   matrix dNNdists;
   intMatrix dMinIDs;
-  dNNdists.r=compLength; dNNdists.pr=compLength; dNNdists.c=K; dNNdists.pc=K; dNNdists.ld=dNNdists.pc;
-  dMinIDs.r=compLength; dMinIDs.pr=compLength; dMinIDs.c=K; dMinIDs.pc=K; dMinIDs.ld=dMinIDs.pc;
+  dNNdists.r=compLength; dNNdists.pr=compLength; dNNdists.c=KMAX; dNNdists.pc=KMAX; dNNdists.ld=dNNdists.pc;
+  dMinIDs.r=compLength; dMinIDs.pr=compLength; dMinIDs.c=KMAX; dMinIDs.pc=KMAX; dMinIDs.ld=dMinIDs.pc;
 
   checkErr( cudaMalloc((void**)&dNNdists.mat,dNNdists.pr*dNNdists.pc*sizeof(*dNNdists.mat)) );
   checkErr( cudaMalloc((void**)&dMinIDs.mat,dMinIDs.pr*dMinIDs.pc*sizeof(*dMinIDs.mat)) );
 
   planKNNWrap(dq, dqMap, dx, dxMap, dNNdists, dMinIDs, dcP, compLength);
-  cudaMemcpy( NNs.mat, dMinIDs.mat, dq.r*K*sizeof(*NNs.mat), cudaMemcpyDeviceToHost );
-  cudaMemcpy( NNdists.mat, dNNdists.mat, dq.r*K*sizeof(*NNdists.mat), cudaMemcpyDeviceToHost );
+  cudaMemcpy( NNs.mat, dMinIDs.mat, dq.r*KMAX*sizeof(*NNs.mat), cudaMemcpyDeviceToHost );
+  cudaMemcpy( NNdists.mat, dNNdists.mat, dq.r*KMAX*sizeof(*NNdists.mat), cudaMemcpyDeviceToHost );
 
   cudaFree(dNNdists.mat);
   cudaFree(dMinIDs.mat);