updated text files
[RBC.git] / kernelWrap.cu
index a2f59b6..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(dx.pr/BLOCK_SIZE,dq.pr/BLOCK_SIZE);
+  dim3 grid;
   
-  dist1Kernel<<<grid,block>>>(dq,dx,dD);
+  unint todoX, todoY, numDoneX, numDoneY;
+
+  numDoneX = 0;
+  while ( numDoneX < dx.pr ){
+    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 );
+      grid.y = todoY/BLOCK_SIZE;
+      dist1Kernel<<<grid,block>>>(dq, numDoneY, dx, numDoneX, dD);
+      numDoneY += todoY;
+    }
+    numDoneX += todoX;
+  }
+
   cudaThreadSynchronize();
 }
 
 
-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));
+  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(const matrix dD, const real *dranges, charMatrix dir){
+  dim3 block(BLOCK_SIZE,BLOCK_SIZE);
+  dim3 grid;
+
+  unint todoX, todoY, numDoneX, numDoneY;
+  
+  numDoneX = 0;
+  while ( numDoneX < dD.pc ){
+    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 );
+      grid.y = todoY/BLOCK_SIZE;
+      rangeSearchKernel<<<grid,block>>>(dD, numDoneX, numDoneY, dranges, dir);
+      numDoneY += todoY;
+    }
+    numDoneX += todoX;
+  }
 
-  findRangeKernel<<<grid,block>>>(dD,dranges,cntWant);
   cudaThreadSynchronize();
 }
 
-void rangeSearchWrap(matrix dD, real *dranges, charMatrix dir){
+void nnWrap(const matrix dq, const matrix dx, real *dMins, unint *dMinIDs){
   dim3 block(BLOCK_SIZE,BLOCK_SIZE);
-  dim3 grid(dD.pc/BLOCK_SIZE,dD.pr/BLOCK_SIZE);
+  dim3 grid;
+  unint numDone, todo;
   
-  rangeSearchKernel<<<grid,block>>>(dD,dranges,dir);
+  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 nnWrap(const matrix dx, const matrix dy, real *dMins, int *dMinIDs){
-  dim3 dimBlock(BLOCK_SIZE,BLOCK_SIZE);
-  dim3 dimGrid;
+
+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;
 
-  rangeCountKernel<<<grid,block>>>(dq,dx,dranges,dcounts);
+  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();
 }
 
 
-/*NOTE: can be deleted */
-void pruneWrap(charMatrix dcM, matrix dD, real *dradiiX, real *dradiiQ){
+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(dD.pr/BLOCK_SIZE,dD.pc/BLOCK_SIZE);
-  
-  pruneKernel<<<grid,block>>>(dD,dradiiX,dradiiQ,dcM);
+  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();
+}
+
+
+
+void rangeCountWrap(const matrix dq, const matrix dx, real *dranges, unint *dcounts){
+  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;
+    rangeCountKernel<<<grid,block>>>(dq,numDone,dx,dranges,dcounts);
+    numDone += todo;
+  }
   cudaThreadSynchronize();
 }
+
 #endif