updated NN functions so that they return the distances as well as the indices
authorLawrence Cayton <lcayton@tuebingen.mpg.de>
Fri, 4 Feb 2011 14:00:54 +0000 (15:00 +0100)
committerLawrence Cayton <lcayton@tuebingen.mpg.de>
Fri, 4 Feb 2011 14:00:54 +0000 (15:00 +0100)
brute.cu
brute.h
driver.cu
rbc.cu
rbc.h

index 7041f28..f40d70a 100644 (file)
--- a/brute.cu
+++ b/brute.cu
@@ -68,17 +68,17 @@ void bruteSearch(matrix x, matrix q, unint *NNs){
 }
 
 
-void bruteK(matrix x, matrix q, intMatrix NNs){
-  matrix dMins;
+void bruteK(matrix x, matrix q, intMatrix NNs, matrix NNdists){
+  matrix dNNdists;
   intMatrix dMinIDs;
   matrix dx, dq;
   
   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;
-  dMins.r=q.r; dMins.pr=q.pr; dMins.c=K; dMins.pc=K; dMins.ld=dMins.pc;
+  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;
 
-  checkErr( cudaMalloc((void**)&dMins.mat, dMins.pc*dMins.pr*sizeof(*dMins.mat)) );
+  checkErr( cudaMalloc((void**)&dNNdists.mat, dNNdists.pc*dNNdists.pr*sizeof(*dNNdists.mat)) );
   checkErr( cudaMalloc((void**)&dMinIDs.mat, dMinIDs.pc*dMinIDs.pr*sizeof(*dMinIDs.mat)) );
   checkErr( cudaMalloc((void**)&dx.mat, dx.pr*dx.pc*sizeof(*dx.mat)) );
   checkErr( cudaMalloc((void**)&dq.mat, dq.pr*dq.pc*sizeof(*dq.mat)) );
@@ -86,11 +86,12 @@ void bruteK(matrix x, matrix q, intMatrix NNs){
   cudaMemcpy(dx.mat,x.mat,x.pr*x.pc*sizeof(*dx.mat),cudaMemcpyHostToDevice);
   cudaMemcpy(dq.mat,q.mat,q.pr*q.pc*sizeof(*dq.mat),cudaMemcpyHostToDevice);
   
-  knnWrap(dq,dx,dMins,dMinIDs);
+  knnWrap(dq,dx,dNNdists,dMinIDs);
 
   cudaMemcpy(NNs.mat,dMinIDs.mat,NNs.pr*NNs.pc*sizeof(*NNs.mat),cudaMemcpyDeviceToHost);
-  
-  cudaFree(dMins.mat);
+  cudaMemcpy(NNdists.mat,dNNdists.mat,NNdists.pr*NNdists.pc*sizeof(*NNdists.mat),cudaMemcpyDeviceToHost);
+
+  cudaFree(dNNdists.mat);
   cudaFree(dMinIDs.mat);
   cudaFree(dx.mat);
   cudaFree(dq.mat);
diff --git a/brute.h b/brute.h
index d551aca..7d610ed 100644 (file)
--- a/brute.h
+++ b/brute.h
@@ -10,6 +10,6 @@
 void bruteRangeCount(matrix,matrix,real*,unint*);
 void bruteSearch(matrix,matrix,unint*);
 void bruteCPU(matrix,matrix,unint*);
-void bruteK(matrix,matrix,intMatrix);
+void bruteK(matrix,matrix,intMatrix,matrix);
 //void bruteKCPU(matrix,matrix,intMatrix);
 #endif
index ff05bb1..aee8ee5 100644 (file)
--- a/driver.cu
+++ b/driver.cu
@@ -27,9 +27,8 @@ unint deviceNum=0;
 int main(int argc, char**argv){
   real *data;
   matrix x, q;
-  unint *NNs;
-  intMatrix NNsK, kNNsRBC;
-  unint i;
+  intMatrix nnsBrute, nnsRBC;
+  matrix distsBrute, distsRBC;
   struct timeval tvB,tvE;
   cudaError_t cE;
   rbcStruct rbcS;
@@ -58,24 +57,27 @@ int main(int argc, char**argv){
   x.r = n; x.c = d; x.pr = PAD(n); x.pc = PAD(d); x.ld = x.pc;
   q.r = m; q.c = d; q.pr = PAD(m); q.pc = PAD(d); q.ld = q.pc;
 
-  NNs = (unint*)calloc( m, sizeof(*NNs) );
-  for(i=0; i<m; i++)
-    NNs[i]=DUMMY_IDX;
-  
+  //Load data 
   readData(dataFile, (n+m), d, data);
   orgData(data, (n+m), d, x, q);
   free(data);
 
-  NNsK.r=q.r; NNsK.pr=q.pr; NNsK.pc=NNsK.c=K; NNsK.ld=NNsK.pc;
-  kNNsRBC.r=q.r; kNNsRBC.pr=q.pr; kNNsRBC.pc=kNNsRBC.c=K; kNNsRBC.ld=kNNsRBC.pc;
-  kNNsRBC.mat = (unint*)calloc(kNNsRBC.pr*kNNsRBC.pc, sizeof(*kNNsRBC.mat));
-  NNsK.mat = (unint*)calloc(NNsK.pr*NNsK.pc, sizeof(*NNsK.mat));
+  //Allocate space for NNs and dists
+  nnsBrute.r=q.r; nnsBrute.pr=q.pr; nnsBrute.pc=nnsBrute.c=K; nnsBrute.ld=nnsBrute.pc;
+  nnsBrute.mat = (unint*)calloc(nnsBrute.pr*nnsBrute.pc, sizeof(*nnsBrute.mat));
+  nnsRBC.r=q.r; nnsRBC.pr=q.pr; nnsRBC.pc=nnsRBC.c=K; nnsRBC.ld=nnsRBC.pc;
+  nnsRBC.mat = (unint*)calloc(nnsRBC.pr*nnsRBC.pc, sizeof(*nnsRBC.mat));
   
-  /* printf("running k-brute force..\n"); */
-  /* gettimeofday(&tvB,NULL); */
-  /* bruteK(x,q,NNsK); */
-  /* gettimeofday(&tvE,NULL); */
-  /* printf("\t.. time elapsed = %6.4f \n",timeDiff(tvB,tvE)); */
+  distsBrute.r=q.r; distsBrute.pr=q.pr; distsBrute.pc=distsBrute.c=K; distsBrute.ld=distsBrute.pc;
+  distsBrute.mat = (real*)calloc(distsBrute.pr*distsBrute.pc, sizeof(*distsBrute.mat));
+  distsRBC.r=q.r; distsRBC.pr=q.pr; distsRBC.pc=distsRBC.c=K; distsRBC.ld=distsRBC.pc;
+  distsRBC.mat = (real*)calloc(distsRBC.pr*distsRBC.pc, sizeof(*distsRBC.mat));
+
+  printf("running k-brute force..\n");
+  gettimeofday(&tvB,NULL);
+  bruteK(x,q,nnsBrute,distsBrute);
+  gettimeofday(&tvE,NULL);
+  printf("\t.. time elapsed = %6.4f \n",timeDiff(tvB,tvE));
 
   printf("\nrunning rbc..\n");
   gettimeofday(&tvB,NULL);
@@ -85,7 +87,7 @@ int main(int argc, char**argv){
 
   //This finds the 32-NN; if you are only interested in the 1-NN, use queryRBC(..) instead
   gettimeofday(&tvB,NULL);
-  kqueryRBC(q, rbcS, kNNsRBC);
+  kqueryRBC(q, rbcS, nnsRBC, distsRBC);
   gettimeofday(&tvE,NULL);
   printf("\t.. query time for krbc = %6.4f \n",timeDiff(tvB,tvE));
   
@@ -97,13 +99,14 @@ int main(int argc, char**argv){
     printf("Execution failed; error type: %s \n", cudaGetErrorString(cE) );
   }
   
-  evalKNNerror(x,q,kNNsRBC);
+  evalKNNerror(x,q,nnsRBC);
   
   cudaThreadExit();
   
-  free(NNs);
-  free(NNsK.mat);
-  free(kNNsRBC.mat);
+  free(nnsBrute.mat);
+  free(nnsRBC.mat);
+  free(distsBrute.mat);
+  free(distsRBC.mat);
   free(x.mat);
   free(q.mat);
 }
@@ -283,9 +286,12 @@ void evalKNNerror(matrix x, matrix q, intMatrix NNs){
   intMatrix NNsB;
   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.mat = (real*)calloc( distsBrute.pr*distsBrute.pc, sizeof(*distsBrute.mat) );
+
   gettimeofday(&tvB,NULL);
-  bruteK(x,q,NNsB);
+  bruteK(x,q,NNsB,distsBrute);
   gettimeofday(&tvE,NULL);
 
    //calc overlap
@@ -343,4 +349,5 @@ void evalKNNerror(matrix x, matrix q, intMatrix NNs){
   free(cnts);
   free(ol);
   free(NNsB.mat);
+  free(distsBrute.mat);
 }
diff --git a/rbc.cu b/rbc.cu
index 61fad14..3237fe1 100644 (file)
--- a/rbc.cu
+++ b/rbc.cu
@@ -16,7 +16,7 @@
 #include "kernelWrap.h"
 #include "sKernelWrap.h"
 
-void queryRBC(const matrix q, const rbcStruct rbcS, unint *NNs){
+void queryRBC(const matrix q, const rbcStruct rbcS, unint *NNs, real* NNdists){
   unint m = q.r;
   unint numReps = rbcS.dr.r;
   unint compLength;
@@ -55,7 +55,7 @@ void queryRBC(const matrix q, const rbcStruct rbcS, unint *NNs){
   checkErr( cudaMalloc( (void**)&dqMap, compLength*sizeof(*dqMap) ) );
   cudaMemcpy( dqMap, qMap, compLength*sizeof(*dqMap), cudaMemcpyHostToDevice );
   
-  computeNNs(rbcS.dx, rbcS.dxMap, dq, dqMap, dcP, NNs, compLength);
+  computeNNs(rbcS.dx, rbcS.dxMap, dq, dqMap, dcP, NNs, NNdists, compLength);
   
   free(qMap);
   freeCompPlan(&dcP);
@@ -68,7 +68,7 @@ void queryRBC(const matrix q, const rbcStruct rbcS, unint *NNs){
 
 //This function is very similar to queryRBC, with a couple of basic changes to handle
 //k-nn.  
-void kqueryRBC(const matrix q, const rbcStruct rbcS, intMatrix NNs){
+void kqueryRBC(const matrix q, const rbcStruct rbcS, intMatrix NNs, matrix NNdists){
   unint m = q.r;
   unint numReps = rbcS.dr.r;
   unint compLength;
@@ -110,7 +110,7 @@ void kqueryRBC(const matrix q, const rbcStruct rbcS, intMatrix NNs){
   checkErr( cudaMalloc( (void**)&dqMap, compLength*sizeof(*dqMap) ) );
   cudaMemcpy( dqMap, qMap, compLength*sizeof(*dqMap), cudaMemcpyHostToDevice );
   
-  computeKNNs(rbcS.dx, rbcS.dxMap, dq, dqMap, dcP, NNs, compLength);
+  computeKNNs(rbcS.dx, rbcS.dxMap, dq, dqMap, dcP, NNs, NNdists, compLength);
 
   free(qMap);
   freeCompPlan(&dcP);
@@ -311,34 +311,36 @@ void fullIntersection(charMatrix cM){
 }
 
 
-void computeNNs(matrix dx, intMatrix dxMap, matrix dq, unint *dqMap, compPlan dcP, unint *NNs, unint compLength){
-  real *dMins;
+void computeNNs(matrix dx, intMatrix dxMap, matrix dq, unint *dqMap, compPlan dcP, unint *NNs, real *NNdists, unint compLength){
+  real *dNNdists;
   unint *dMinIDs;
   
-  checkErr( cudaMalloc((void**)&dMins,compLength*sizeof(*dMins)) );
+  checkErr( cudaMalloc((void**)&dNNdists,compLength*sizeof(*dNNdists)) );
   checkErr( cudaMalloc((void**)&dMinIDs,compLength*sizeof(*dMinIDs)) );
 
-  planNNWrap(dq, dqMap, dx, dxMap, dMins, dMinIDs, dcP, compLength);
-  cudaMemcpy( NNs, dMinIDs, dq.r*sizeof(*NNs), cudaMemcpyDeviceToHost);
-  
-  cudaFree(dMins);
+  planNNWrap(dq, dqMap, dx, dxMap, dNNdists, dMinIDs, dcP, compLength );
+  cudaMemcpy( NNs, dMinIDs, dq.r*sizeof(*NNs), cudaMemcpyDeviceToHost );
+  cudaMemcpy( NNdists, dNNdists, dq.r*sizeof(*dNNdists), cudaMemcpyDeviceToHost );
+
+  cudaFree(dNNdists);
   cudaFree(dMinIDs);
 }
 
 
-void computeKNNs(matrix dx, intMatrix dxMap, matrix dq, unint *dqMap, compPlan dcP, intMatrix NNs, unint compLength){
-  matrix dMins;
+void computeKNNs(matrix dx, intMatrix dxMap, matrix dq, unint *dqMap, compPlan dcP, intMatrix NNs, matrix NNdists, unint compLength){
+  matrix dNNdists;
   intMatrix dMinIDs;
-  dMins.r=compLength; dMins.pr=compLength; dMins.c=K; dMins.pc=K; dMins.ld=dMins.pc;
+  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;
 
-  checkErr( cudaMalloc((void**)&dMins.mat,dMins.pr*dMins.pc*sizeof(*dMins.mat)) );
+  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, dMins, dMinIDs, dcP, compLength);
-  cudaMemcpy( NNs.mat, dMinIDs.mat, dq.r*K*sizeof(*NNs.mat), cudaMemcpyDeviceToHost);
+  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 );
 
-  cudaFree(dMins.mat);
+  cudaFree(dNNdists.mat);
   cudaFree(dMinIDs.mat);
 }
 
diff --git a/rbc.h b/rbc.h
index eaf19a1..ad8669f 100644 (file)
--- a/rbc.h
+++ b/rbc.h
@@ -9,8 +9,8 @@
 
 
 void buildRBC(const matrix,rbcStruct*,unint, unint);
-void queryRBC(const matrix,const rbcStruct,unint*);
-void kqueryRBC(const matrix,const rbcStruct,intMatrix);
+void queryRBC(const matrix,const rbcStruct,unint*,real*);
+void kqueryRBC(const matrix,const rbcStruct,intMatrix,matrix);
 void destroyRBC(rbcStruct*);
 void distSubMat(matrix,matrix,matrix,unint,unint);
 void computeReps(matrix,matrix,unint*,real*);
@@ -21,8 +21,8 @@ void idIntersection(charMatrix);
 void fullIntersection(charMatrix);
 void initCompPlan(compPlan*,charMatrix,unint*,unint*,unint);
 void freeCompPlan(compPlan*);
-void computeNNs(matrix,intMatrix,matrix,unint*,compPlan,unint*,unint);
-void computeKNNs(matrix,intMatrix,matrix,unint*,compPlan,intMatrix,unint);
+void computeNNs(matrix,intMatrix,matrix,unint*,compPlan,unint*,real*,unint);
+void computeKNNs(matrix,intMatrix,matrix,unint*,compPlan,intMatrix,matrix,unint);
 void setupReps(matrix,rbcStruct*,int);
 
 #endif