Basic cleanup
authorLawrence Cayton <lcayton@tuebingen.mpg.de>
Fri, 13 Aug 2010 08:58:36 +0000 (10:58 +0200)
committerLawrence Cayton <lcayton@tuebingen.mpg.de>
Fri, 13 Aug 2010 08:58:36 +0000 (10:58 +0200)
defs.h
kernelWrap.cu
kernels.cu
kernels.h
rbc.cu
sKernel.cu

diff --git a/defs.h b/defs.h
index a517024..158fccb 100644 (file)
--- a/defs.h
+++ b/defs.h
@@ -7,7 +7,7 @@
 #define BLOCK_SIZE 16 //must be a power of 2
 
 #define MAX_BS 65535 //max block size
-#define SCAN_WIDTH 1024 
+#define SCAN_WIDTH 1024
 
 // Format that the data is manipulated in:
 typedef float real;
index a2f59b6..b692f7b 100644 (file)
@@ -8,9 +8,24 @@
 
 void dist1Wrap(matrix dq, matrix dx, matrix dD){
   dim3 block(BLOCK_SIZE,BLOCK_SIZE);
-  dim3 grid(dx.pr/BLOCK_SIZE,dq.pr/BLOCK_SIZE);
+  dim3 grid;
+  
+  int 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;
+  }
   
-  dist1Kernel<<<grid,block>>>(dq,dx,dD);
   cudaThreadSynchronize();
 }
 
@@ -20,14 +35,31 @@ void findRangeWrap(matrix dD, real *dranges, int cntWant){
   dim3 grid(1,4*(dD.pr/BLOCK_SIZE));
 
   findRangeKernel<<<grid,block>>>(dD,dranges,cntWant);
+
+  
   cudaThreadSynchronize();
 }
 
 void rangeSearchWrap(matrix dD, real *dranges, charMatrix dir){
   dim3 block(BLOCK_SIZE,BLOCK_SIZE);
   dim3 grid(dD.pc/BLOCK_SIZE,dD.pr/BLOCK_SIZE);
+
+  int todoX, todoY, numDoneX, numDoneY;
   
-  rangeSearchKernel<<<grid,block>>>(dD,dranges,dir);
+  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;
+  }
+
   cudaThreadSynchronize();
 }
 
index 04cd252..d3fc003 100644 (file)
@@ -191,12 +191,12 @@ __global__ void nnKernel(const matrix Q, const matrix X, real *dMins, int *dMinI
 }
 
 
-__device__ void dist1Kernel(const matrix Q, const matrix X, matrix D){
+__device__ void dist1Kernel(const matrix Q, int qStart, const matrix X, int xStart, matrix D){
   int c, i, j;
 
-  int qB = blockIdx.y*BLOCK_SIZE;
+  int qB = blockIdx.y*BLOCK_SIZE + qStart;
   int q  = threadIdx.y;
-  int xB = blockIdx.x*BLOCK_SIZE;
+  int xB = blockIdx.x*BLOCK_SIZE + xStart;
   int x = threadIdx.x;
 
   real ans=0;
@@ -323,9 +323,9 @@ __global__ void findRangeKernel(matrix D, real *ranges, int cntWant){
 }
 
 
-__global__ void rangeSearchKernel(matrix D, real *ranges, charMatrix ir){
-  int col = blockIdx.x*BLOCK_SIZE + threadIdx.x;
-  int row = blockIdx.y*BLOCK_SIZE + threadIdx.y;
+__global__ void rangeSearchKernel(matrix D, int xOff, int yOff, real *ranges, charMatrix ir){
+  int col = blockIdx.x*BLOCK_SIZE + threadIdx.x + xOff;
+  int row = blockIdx.y*BLOCK_SIZE + threadIdx.y + yOff;
 
   ir.mat[IDX( row, col, ir.ld )] = D.mat[IDX( row, col, D.ld )] < ranges[row];
 
index c29b174..cf26734 100644 (file)
--- a/kernels.h
+++ b/kernels.h
@@ -4,12 +4,12 @@
 #include "defs.h"
 __global__ void planNNKernel(const matrix,const matrix,real*,int*,compPlan,intMatrix,int*,int);
 __global__ void pruneKernel(const matrix,const real*,const real*,charMatrix);
-__global__ void dist1Kernel(const matrix,const matrix,matrix);
+__global__ void dist1Kernel(const matrix,int,const matrix,int,matrix);
 __device__ matrix getSubMat(matrix,int,int);
 __global__ void nnKernel(const matrix,const matrix,real*,int*);
 
 __global__ void findRangeKernel(matrix,real*,int);
-__global__ void rangeSearchKernel(matrix,real*,charMatrix);
+__global__ void rangeSearchKernel(matrix,int,int,real*,charMatrix);
 __global__ void rangeCountKernel(const matrix,const matrix,real*,int*);
 
 #endif
diff --git a/rbc.cu b/rbc.cu
index 126450c..95aedd8 100644 (file)
--- a/rbc.cu
+++ b/rbc.cu
@@ -175,7 +175,7 @@ void build(const matrix dx, const matrix dr, intMatrix xmap, int *counts, int s)
     distSubMat(dr, dx, dD, row, pip); //compute the distance matrix
     findRangeWrap(dD, dranges, s);  //find an appropriate range
     rangeSearchWrap(dD, dranges, dir); //set binary vector for points in range
-    
+    printf("after range search\n");
     sumWrap(dir, dSums);  //This and the next call perform the parallel compaction.
     buildMapWrap(xmap, dir, dSums, row);
     getCountsWrap(dCnts,dir,dSums);  //How many points are assigned to each rep?  It is not 
index e626e2e..d643927 100644 (file)
@@ -129,6 +129,8 @@ void sumWrap(charMatrix in, intMatrix sum){
   
   sumKernel<<<grid,block>>>(in, sum, dAux, n);
   cudaThreadSynchronize();
+
+  //While(){}
   grid.x=1;
 
   sumKernelI<<<grid,block>>>(dAux, dAux, dDummy, numScans);