updated text files
[RBC.git] / kernelWrap.cu
index 292d941..d46f2cb 100644 (file)
@@ -1,3 +1,7 @@
+/* This file is part of the Random Ball Cover (RBC) library.
+ * (C) Copyright 2010, Lawrence Cayton [lcayton@tuebingen.mpg.de]
+ */
+
 #ifndef KERNELWRAP_CU
 #define KERNELWRAP_CU
 
 #include "kernels.h"
 #include "defs.h"
 
-void dist1Wrap(matrix dq, matrix dx, matrix dD){
+void dist1Wrap(const matrix dq, const matrix dx, matrix dD){
   dim3 block(BLOCK_SIZE,BLOCK_SIZE);
   dim3 grid;
   
-  int todoX, todoY, numDoneX, numDoneY;
+  unint todoX, todoY, numDoneX, numDoneY;
 
   numDoneX = 0;
   while ( numDoneX < dx.pr ){
-    todoX = min( dx.pr - numDoneX, MAX_BS*BLOCK_SIZE );
+    todoX = MIN( dx.pr - numDoneX, MAX_BS*BLOCK_SIZE );
     grid.x = todoX/BLOCK_SIZE;
     numDoneY = 0;
     while( numDoneY < dq.pr ){
-      todoY = min( dq.pr - numDoneY, MAX_BS*BLOCK_SIZE );
+      todoY = MIN( dq.pr - numDoneY, MAX_BS*BLOCK_SIZE );
       grid.y = todoY/BLOCK_SIZE;
       dist1Kernel<<<grid,block>>>(dq, numDoneY, dx, numDoneX, dD);
       numDoneY += todoY;
@@ -30,29 +34,35 @@ void dist1Wrap(matrix dq, matrix dx, matrix dD){
 }
 
 
-void findRangeWrap(matrix dD, real *dranges, int cntWant){
+void findRangeWrap(const matrix dD, real *dranges, unint cntWant){
   dim3 block(4*BLOCK_SIZE,BLOCK_SIZE/4);
   dim3 grid(1,4*(dD.pr/BLOCK_SIZE));
-
-  findRangeKernel<<<grid,block>>>(dD,dranges,cntWant);
-
+  unint numDone, todo;
   
+  numDone=0;
+  while( numDone < dD.pr ){
+    todo = MIN ( dD.pr - numDone, MAX_BS*BLOCK_SIZE/4 );
+    grid.y = 4*(todo/BLOCK_SIZE);
+    findRangeKernel<<<grid,block>>>(dD, numDone, dranges, cntWant);
+    numDone += todo;
+  }
   cudaThreadSynchronize();
 }
 
-void rangeSearchWrap(matrix dD, real *dranges, charMatrix dir){
+
+void rangeSearchWrap(const matrix dD, const real *dranges, charMatrix dir){
   dim3 block(BLOCK_SIZE,BLOCK_SIZE);
-  dim3 grid(dD.pc/BLOCK_SIZE,dD.pr/BLOCK_SIZE);
+  dim3 grid;
 
-  int todoX, todoY, numDoneX, numDoneY;
+  unint todoX, todoY, numDoneX, numDoneY;
   
   numDoneX = 0;
   while ( numDoneX < dD.pc ){
-    todoX = min( dD.pc - numDoneX, MAX_BS*BLOCK_SIZE );
+    todoX = MIN( dD.pc - numDoneX, MAX_BS*BLOCK_SIZE );
     grid.x = todoX/BLOCK_SIZE;
     numDoneY = 0;
     while( numDoneY < dD.pr ){
-      todoY = min( dD.pr - numDoneY, MAX_BS*BLOCK_SIZE );
+      todoY = MIN( dD.pr - numDoneY, MAX_BS*BLOCK_SIZE );
       grid.y = todoY/BLOCK_SIZE;
       rangeSearchKernel<<<grid,block>>>(dD, numDoneX, numDoneY, dranges, dir);
       numDoneY += todoY;
@@ -63,32 +73,94 @@ void rangeSearchWrap(matrix dD, real *dranges, charMatrix dir){
   cudaThreadSynchronize();
 }
 
-void nnWrap(const matrix dx, const matrix dy, real *dMins, int *dMinIDs){
-  dim3 dimBlock(BLOCK_SIZE,BLOCK_SIZE);
-  dim3 dimGrid;
+void nnWrap(const matrix dq, const matrix dx, real *dMins, unint *dMinIDs){
+  dim3 block(BLOCK_SIZE,BLOCK_SIZE);
+  dim3 grid;
+  unint numDone, todo;
+  
+  grid.x = 1;
+
+  numDone = 0;
+  while( numDone < dq.pr ){
+    todo = MIN( dq.pr - numDone, MAX_BS*BLOCK_SIZE );
+    grid.y = todo/BLOCK_SIZE;
+    nnKernel<<<grid,block>>>(dq,numDone,dx,dMins,dMinIDs);
+    numDone += todo;
+  }
+  cudaThreadSynchronize();
+
+}
+
+
+void knnWrap(const matrix dq, const matrix dx, matrix dMins, intMatrix dMinIDs){
+  dim3 block(BLOCK_SIZE,BLOCK_SIZE);
+  dim3 grid;
+  unint numDone, todo;
   
-  dimGrid.x = 1;
-  dimGrid.y = dx.pr/dimBlock.y + (dx.pr%dimBlock.y==0 ? 0 : 1);
-  nnKernel<<<dimGrid,dimBlock>>>(dx,dy,dMins,dMinIDs);
+  grid.x = 1;
+
+  numDone = 0;
+  while( numDone < dq.pr ){
+    todo = MIN( dq.pr - numDone, MAX_BS*BLOCK_SIZE );
+    grid.y = todo/BLOCK_SIZE;
+    knnKernel<<<grid,block>>>(dq,numDone,dx,dMins,dMinIDs);
+    numDone += todo;
+  }
   cudaThreadSynchronize();
+
 }
 
 
-void rangeCountWrap(const matrix dq, const matrix dx, real *dranges, int *dcounts){
+void planNNWrap(const matrix dq, const unint *dqMap, const matrix dx, const intMatrix dxMap, real *dMins, unint *dMinIDs, compPlan dcP, unint compLength){
   dim3 block(BLOCK_SIZE,BLOCK_SIZE);
-  dim3 grid(1,dq.pr/BLOCK_SIZE);
+  dim3 grid;
+  unint todo;
+
+  grid.x = 1;
+  unint numDone = 0;
+  while( numDone<compLength ){
+    todo = MIN( (compLength-numDone) , MAX_BS*BLOCK_SIZE );
+    grid.y = todo/BLOCK_SIZE;
+    planNNKernel<<<grid,block>>>(dq,dqMap,dx,dxMap,dMins,dMinIDs,dcP,numDone);
+    numDone += todo;
+  }
+  cudaThreadSynchronize();
+}
+
 
-  rangeCountKernel<<<grid,block>>>(dq,dx,dranges,dcounts);
+void planKNNWrap(const matrix dq, const unint *dqMap, const matrix dx, const intMatrix dxMap, matrix dMins, intMatrix dMinIDs, compPlan dcP, unint compLength){
+  dim3 block(BLOCK_SIZE,BLOCK_SIZE);
+  dim3 grid;
+  unint todo;
+
+  grid.x = 1;
+  unint numDone = 0;
+  while( numDone<compLength ){
+    todo = MIN( (compLength-numDone) , MAX_BS*BLOCK_SIZE );
+    grid.y = todo/BLOCK_SIZE;
+    planKNNKernel<<<grid,block>>>(dq,dqMap,dx,dxMap,dMins,dMinIDs,dcP,numDone);
+    numDone += todo;
+  }
   cudaThreadSynchronize();
 }
 
 
-/*NOTE: can be deleted */
-void pruneWrap(charMatrix dcM, matrix dD, real *dradiiX, real *dradiiQ){
+
+void rangeCountWrap(const matrix dq, const matrix dx, real *dranges, unint *dcounts){
   dim3 block(BLOCK_SIZE,BLOCK_SIZE);
-  dim3 grid(dD.pr/BLOCK_SIZE,dD.pc/BLOCK_SIZE);
-  
-  pruneKernel<<<grid,block>>>(dD,dradiiX,dradiiQ,dcM);
+  dim3 grid;
+  unint numDone, todo;
+
+  grid.x=1;
+
+  numDone = 0;
+  while( numDone < dq.pr ){
+    todo = MIN( dq.pr - numDone, MAX_BS*BLOCK_SIZE );
+    grid.y = todo/BLOCK_SIZE;
+    rangeCountKernel<<<grid,block>>>(dq,numDone,dx,dranges,dcounts);
+    numDone += todo;
+  }
   cudaThreadSynchronize();
 }
+
 #endif