   // f(I(xi)-I(x)) = 4-|I(xi)-I(x)| if greater than 0, 0 otherwise
   // horizontal filter
 
-
   // Current pixel being filtered is middle of our set (4 back, in orginal its 3 for some reason)
   //Note this isn't truely the center since original uses 8 vectore resisters
-  float val_curr = *(D_copy+idx);
+  float val_curr = D_shared[ut][vt];
 
   float weight_sum0 = 0;
   float weight_sum = 0;
   float factor_sum = 0;
 
-  float weight_sum2 = 0;
-  float factor_sum2 = 0;
-
   for(int32_t i=0; i < 8; i++){
-    weight_sum0 = 4.0f - fabs(*(D_copy+idx+(i-4))-val_curr);
+    weight_sum0 = 4.0f - fabs(D_shared[ut+(i-4)][vt]-val_curr);
     weight_sum0 = max(0.0f, weight_sum0);
     weight_sum += weight_sum0;
-    factor_sum += *(D_copy+idx+(i-4))*weight_sum0;
-  }
-
-  for(int32_t i=0; i < 8; i++){
-    weight_sum0 = 4.0f - fabs(D_copy_S[ut+(i-4)][vt]-val_curr);
-    weight_sum0 = max(0.0f, weight_sum0);
-    weight_sum2 += weight_sum0;
-    factor_sum2 += D_copy_S[ut+(i-4)][vt]*weight_sum0;
-  }
-
-  if(blockIdx.x == 0 && blockIdx.y == 0){
-      if(weight_sum != weight_sum2){
-          printf("Shared error %.10f - %.10f | %d - %d\n", weight_sum, weight_sum2,threadIdx.x,threadIdx.y);
-      }
-  }
-  
-  if (weight_sum2>0) {
-      float d = factor_sum2/weight_sum2;
-      if (d>=0) *(D_horiz+idx) = d;
+    factor_sum += D_shared[ut+(i-4)][vt]*weight_sum0;
+  }
+
+  if (weight_sum>0) {
+      float d = factor_sum/weight_sum;
+      if (d>=0) *(D+idx) = d;
   }
   
   __syncthreads();
-
+  //Update shared memory
   if(threadIdx.x == blockDim.x-1){
-      D_copy_S[ut+1][vt] = D_horiz[idx+1];
-      D_copy_S[ut+2][vt] = D_horiz[idx+2];
-      D_copy_S[ut+3][vt] = D_horiz[idx+3];
-      D_copy_S[ut+4][vt] = D_horiz[idx+4];
+      D_shared[ut+1][vt] = D[idx+1];
+      D_shared[ut+2][vt] = D[idx+2];
+      D_shared[ut+3][vt] = D[idx+3];
+      //D_shared[ut+4][vt] = D[idx+4];
   }
   if(threadIdx.x == 0){
-      D_copy_S[ut-4][vt] = D_horiz[idx-4];
-      D_copy_S[ut-3][vt] = D_horiz[idx-3];
-      D_copy_S[ut-2][vt] = D_horiz[idx-2];
-      D_copy_S[ut-1][vt] = D_horiz[idx-1];
+      D_shared[ut-4][vt] = D[idx-4];
+      D_shared[ut-3][vt] = D[idx-3];
+      D_shared[ut-2][vt] = D[idx-2];
+      D_shared[ut-1][vt] = D[idx-1];
   }
   if(threadIdx.y == 0){
-      D_copy_S[ut][vt-4] = D_horiz[(v0-4)*D_width+u0];
-      D_copy_S[ut][vt-3] = D_horiz[(v0-3)*D_width+u0];
-      D_copy_S[ut][vt-2] = D_horiz[(v0-2)*D_width+u0];
-      D_copy_S[ut][vt-1] = D_horiz[(v0-1)*D_width+u0];
+      D_shared[ut][vt-4] = D[(v0-4)*D_width+u0];
+      D_shared[ut][vt-3] = D[(v0-3)*D_width+u0];
+      D_shared[ut][vt-2] = D[(v0-2)*D_width+u0];
+      D_shared[ut][vt-1] = D[(v0-1)*D_width+u0];
   }
   if(threadIdx.y == blockDim.y-1){
-      D_copy_S[ut][vt+1] = D_horiz[(v0+1)*D_width+u0];
-      D_copy_S[ut][vt+2] = D_horiz[(v0+2)*D_width+u0];
-      D_copy_S[ut][vt+3] = D_horiz[(v0+3)*D_width+u0];
-      D_copy_S[ut][vt+4] = D_horiz[(v0+4)*D_width+u0];
+      D_shared[ut][vt+1] = D[(v0+1)*D_width+u0];
+      D_shared[ut][vt+2] = D[(v0+2)*D_width+u0];
+      D_shared[ut][vt+3] = D[(v0+3)*D_width+u0];
+      //D_shared[ut][vt+4] = D[(v0+4)*D_width+u0];
   }
 
   if(D[idx] < 0){
-      D_copy_S[ut][vt] = -10;
+      D_shared[ut][vt] = -10;
   }else{
-      D_copy_S[ut][vt] = D_horiz[idx];
+      D_shared[ut][vt] = D[idx];
   }
 
   __syncthreads();
 
   // vertical filter
   // set pixel of interest
-  val_curr = *(D_horiz+idx);
+  val_curr = D_shared[ut][vt];
 
   weight_sum0 = 0;
   weight_sum = 0;
   factor_sum = 0;
 
   for(int32_t i=0; i < 8; i++){
-    weight_sum0 = 4.0f - fabs(*(D_horiz+(v0+(i-4))*D_width+u0)-val_curr);
+    weight_sum0 = 4.0f - fabs(D_shared[ut][vt+(i-4)]-val_curr);
     weight_sum0 = max(0.0f, weight_sum0);
     weight_sum += weight_sum0;
-    factor_sum += *(D_horiz+(v0+(i-4))*D_width+u0)*weight_sum0;
-  }
-  
-  weight_sum2 = 0;
-  factor_sum2 = 0;
-
-  for(int32_t i=0; i < 8; i++){
-    weight_sum0 = 4.0f - fabs(D_copy_S[ut][vt+(i-4)]-val_curr);
-    weight_sum0 = max(0.0f, weight_sum0);
-    weight_sum2 += weight_sum0;
-    factor_sum2 += D_copy_S[ut][vt+(i-4)]*weight_sum0;
-  }
-
-  if(blockIdx.x == 0 && blockIdx.y == 0){
-      if(weight_sum != weight_sum2){
-          printf("Shared error in Y %.10f - %.10f | %d - %d\n", weight_sum, weight_sum2,threadIdx.x,threadIdx.y);
-      }
+    factor_sum += D_shared[ut][vt+(i-4)]*weight_sum0;
   }
 
   if (weight_sum>0) {

}

// implements approximation to 8x8 bilateral filtering
__global__ void adaptiveMeanGPU8 (float* D, float* D_copy, float* D_horiz, int32_t D_width, int32_t D_height) {
  
  // Global coordinates and Pixel id
  uint32_t u0 = blockDim.x*blockIdx.x + threadIdx.x + 4;


    // CUDA copy over needed memory information
    // disparity_grid and respective copies
    float* d_D;

    // Allocate on global memory and copy
    cudaMalloc((void**) &d_D, width*height*sizeof(float));
    cudaMemcpy(d_D, D, width*height*sizeof(float), cudaMemcpyHostToDevice);

    //Kernel go!
    adaptiveMeanGPU8<<<DimGrid, DimBlock>>>(d_D, width, height);

    // Sync after the kernel is launched
    cudaDeviceSynchronize();
