Added text files; minor cleanup
authorLawrence Cayton <lcayton@tuebingen.mpg.de>
Thu, 12 Aug 2010 15:32:14 +0000 (17:32 +0200)
committerLawrence Cayton <lcayton@tuebingen.mpg.de>
Thu, 12 Aug 2010 15:32:14 +0000 (17:32 +0200)
brute.cu
kernels.cu
rbc.cu
readme.txt [new file with mode: 0644]
sKernel.cu
todo.txt [new file with mode: 0644]

index 388ee26..ca55c54 100644 (file)
--- a/brute.cu
+++ b/brute.cu
@@ -13,7 +13,7 @@
 
 void bruteRangeCount(matrix x, matrix q, real *ranges, int *cnts){
   matrix dx, dq;
-  realdranges;
+  real *dranges;
   int *dcnts;
   
   copyAndMove(&dx, &x);
@@ -60,9 +60,9 @@ void bruteSearch(matrix x, matrix q, int *NNs){
   cudaFree(dMinIDs);
   cudaFree(dx.mat);
   cudaFree(dq.mat);
-
 }
 
+
 void bruteCPU(matrix X, matrix Q, int *NNs){
   real *dtoNNs; 
   real temp;
index f015d3d..04cd252 100644 (file)
@@ -191,14 +191,6 @@ __global__ void nnKernel(const matrix Q, const matrix X, real *dMins, int *dMinI
 }
 
 
-
-
-
-
-
-
-
-
 __device__ void dist1Kernel(const matrix Q, const matrix X, matrix D){
   int c, i, j;
 
@@ -234,6 +226,7 @@ __device__ void dist1Kernel(const matrix Q, const matrix X, matrix D){
 }
 
 
+
 __global__ void findRangeKernel(matrix D, real *ranges, int cntWant){
   
   int row = blockIdx.y*(BLOCK_SIZE/4)+threadIdx.y;
diff --git a/rbc.cu b/rbc.cu
index db3f339..126450c 100644 (file)
--- a/rbc.cu
+++ b/rbc.cu
@@ -43,9 +43,6 @@ void rbc(matrix x, matrix q, int numReps, int s, int *NNs){
 
   qID = (int*)calloc(PAD(m+(BLOCK_SIZE-1)*pnr),sizeof(*qID));
 
-  for(i=0;i<m;i++)
-    qID[i]=i;
-
   cM.mat = (char*)calloc(pnr*pnr,sizeof(*cM.mat));
   cM.r=numReps; cM.c=numReps; cM.pr=pnr; cM.pc=pnr; cM.ld=cM.pc;
 
@@ -290,8 +287,7 @@ void buildQMap(matrix q, int *qID, int *repIDs, int numReps, int *compLength){
   int n=q.r;
   int i;
   int *gS; //groupSize
-  int *yID;
-
+  
   gS = (int*)calloc(numReps+1,sizeof(*gS));
   
   for( i=0; i<n; i++ )
@@ -304,21 +300,14 @@ void buildQMap(matrix q, int *qID, int *repIDs, int numReps, int *compLength){
   
   *compLength = gS[numReps];
   
-  yID = (int*)calloc((*compLength),sizeof(*yID));
   for( i=0; i<(*compLength); i++ )
-    yID[i]=DUMMY_IDX;
+    qID[i]=DUMMY_IDX;
   
-
   for( i=0; i<n; i++ ){
-    yID[gS[repIDs[i]]]=qID[i];
+    qID[gS[repIDs[i]]]=i;
     gS[repIDs[i]]++;
   }
 
-  for( i=0; i<(*compLength); i++ ){
-    qID[i]=yID[i];
-  }
-  
-  free(yID);
   free(gS);
 }
 
diff --git a/readme.txt b/readme.txt
new file mode 100644 (file)
index 0000000..9a792b0
--- /dev/null
@@ -0,0 +1,73 @@
+***Random Ball Cover (RBC)***
+Lawrence Cayton
+lcayton@tuebingen.mpg.de
+
+(C) Copyright 2010, Lawrence Cayton
+This program is free software: you can redistribute it and/or modify 
+it under the terms of the GNU General Public License as published by
+the Free Software Foundation, either version 3 of the License, or
+(at your option) any later version.
+
+This program is distributed in the hope that it will be useful,
+but WITHOUT ANY WARRANTY; without even the implied warranty of
+MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+GNU General Public License for more details.
+
+You should have received a copy of the GNU General Public License
+along with this program.  If not, see <http://www.gnu.org/licenses/>.
+
+---------------------------------------------------------------------
+
+This is a C and CUDA implementation of the Random Ball Cover data 
+structure for nearest neighbor search.  
+
+Notes on the code:
+* The code requires that the entire DB and query set fit into the 
+device memory.  
+
+* For the most part, device variables (ie arrays residing on the GPU)
+begin with a lowercase d.  For example, the device version of the 
+DB variable x is dx.  
+
+* The computePlan code is a bit more complex than is needed for the 
+version of the RBC search algorithm described in the paper.  The 
+search algorithm described in the paper has two steps: (1) Find the 
+closest representative to the query.  (2) Explore the points owned
+by that representative (ie the s-closest points to the representative
+in the DB).  The computePlan code is more complex to make it easy
+to try out other options.  For example, one could search the points
+owned by the two closest representatives to the query instead.  This
+would require only minor changes to the code.
+
+* Currently the software works only in single precision.  If you wish
+to switch to double precision, you must edit the defs.h file.  Simply 
+uncomment the lines
+
+typedef double real;
+#define MAX_REAL DBL_MAX
+
+and comment out the lines
+
+typedef float real;
+#define MAX_REAL FLT_MAX
+
+Then, you must do a 
+
+make clean
+
+followed by another make.
+
+* This software has been tested on the following graphics cards:
+NVIDIA GTX 285
+NVIDIA Tesla c2050.
+
+* This sotware has been tested under the following software setup:
+Ubuntu 9.10 (linux)
+gcc 4.4
+cuda 3.1
+
+Please share your experience getting it to work under Windows and
+Mac OSX!
+
+
index 12342a8..e626e2e 100644 (file)
@@ -47,8 +47,7 @@ __global__ void sumKernel(charMatrix in, intMatrix sum, intMatrix sumaux, int n)
       t = ssum[ off*(2*id+1)-1 ];
       ssum[ off*(2*id+1)-1 ] = ssum[ off*(2*id+2)-1 ];
       ssum[ off*(2*id+2)-1 ] += t;
-    }
-    
+    }   
   }
 
   __syncthreads();
diff --git a/todo.txt b/todo.txt
new file mode 100644 (file)
index 0000000..8fb52c8
--- /dev/null
+++ b/todo.txt
@@ -0,0 +1,12 @@
+- fix problems related to limitations on number of kernel launches
+- Make sure that the code works with r,s that are not powers of 2
+
+- Currently, the scan Kernel is limited to 1024*1024.
+
+- consider putting iteration limit on findRangeKernel-- in fact this
+  is necessary for the case when there are multiple points at the 
+  same distance
+
+- findRangeKernel will be called with too few rows to make use of the
+  processor.  So this is another reason to break it down so that 
+  each block contains only one or a few rows.