ch17

[book_gpu.git] / BookGPU / Chapters / chapter10 / updateBasis.cu

diff --git a/BookGPU/Chapters/chapter10/updateBasis.cu b/BookGPU/Chapters/chapter10/updateBasis.cu
index 894667a8dcb03298c0886007bcf671ca9bb85e2a..a699a6cab1377aef5024579459680012488aef95 100644 (file)
--- a/BookGPU/Chapters/chapter10/updateBasis.cu
+++ b/BookGPU/Chapters/chapter10/updateBasis.cu
@@ -1,25 +1,25 @@
-extern __shared__ volatile double sData[];
+extern __shared__ volatile double sdata[];
 __global__ void
-updateBasisKernel(int m, uint l, double d_l, double *B, uint pitch_B, double *d)
-{        
+updateBasisKernel(int m, uint l, double d_l, double *B,
+                  uint pitch_B, double *d) {        
        uint bId = blockIdx.x, tId = threadIdx.x;
        uint colStart = bId*pitch_B;    
        double Bij, d_i, B2ij;
 
-       // First thread load Blj so it can be 
-       // broadcasted via shared memory to each threads
-       if(tId == 0)
-               sdata[0] = B[colStart+leave] / d_l;                                     
+       // First thread loads Blj so it can be 
+       // broadcast via shared memory to each thread
+       if (tId == 0)
+               sdata[0] = B[colStart+l] / d_l;                                 
        __syncthreads();                
        
-       // Each thread proccess mutiple elements
-       while(tId < m){         
+       // Each thread proccesses multiple elements
+       while (tId < m) {               
                // Load di and Bij
                d_i = d[tId];
                Bij = B[colStart+tId];
                // Update Bij
                B2ij = sdata[0];
-               if(tId != q){
+               if (tId != q) {
                        B2ij *= -d_i;
                        B2ij += Bij
                }
@@ -28,4 +28,4 @@ updateBasisKernel(int m, uint l, double d_l, double *B, uint pitch_B, double *d)
 
                tId += blockDim.x;                                      
        }
-}
\ No newline at end of file
+}