   cudaMalloc( (void**)&data_device, DATA_DIM*DATA_DIM*DATA_DIM*sizeof(unsigned char));
   cudaMalloc( (void**)&region_device, DATA_DIM*DATA_DIM*DATA_DIM*sizeof(unsigned char));
   cudaMalloc( (void**)&finished_device, sizeof(int));
-  cudaMemcpy( data_device, data, DATA_DIM*DATA_DIM*DATA_DIM*sizeof(unsigned char), cudaMemcpyDeviceToHost);
-  cudaMemcpy( region_device, region, DATA_DIM*DATA_DIM*DATA_DIM*sizeof(unsigned char), cudaMemcpyDeviceToHost);
-  cudaMemcpy( finished_device, finished, sizeof(int), cudaMemcpyDeviceToHost);
-
-  dim3 dimBlock( 16, 1 );
-  dim3 dimGrid( 1, 1 );
+  cudaMemcpy( data_device, data, DATA_DIM*DATA_DIM*DATA_DIM*sizeof(unsigned char), cudaMemcpyHostToDevice);
+  cudaMemcpy( region_device, region, DATA_DIM*DATA_DIM*DATA_DIM*sizeof(unsigned char), cudaMemcpyHostToDevice);
+  cudaMemcpy( finished_device, finished, sizeof(int), cudaMemcpyHostToDevice);
+
+  dim3 dimBlock( 8, 8, 8 );
+  dim3 dimGrid( 64, 64, 64 );
 
   while(finished > 0){
 
-  region_grow_kernel<<<dimGrid, dimBlock>>>(data_device, region_device, finished_device);
+    region_grow_kernel<<<dimGrid, dimBlock>>>(data_device, region_device, finished_device);
+    cudaMemcpy( finished, finished_device, sizeof(int), cudaMemcpyDeviceToHost);
 
   }
 
   return region;
-  }*/
+  }
 
 
 unsigned char* grow_region_gpu_shared(unsigned char* data){

}


/*unsigned char* grow_region_gpu(unsigned char* data){
  unsigned char* region = (unsigned char*)calloc(sizeof(unsigned char), DATA_DIM*DATA_DIM*DATA_DIM);
  int* finished = (int*)calloc(sizeof(int), 1);

  int3 seed = {.x=50, .y=300, .z=300};
  region[seed.z *DATA_DIM*DATA_DIM + seed.y*DATA_DIM + seed.x] = 1;

  unsigned char *data_device;
  unsigned char *region_device;

}


__global__ void region_grow_kernel(unsigned char* data, unsigned char* region, int* finished){
  unsigned char* region = (unsigned char*)calloc(sizeof(unsigned char), DATA_DIM*DATA_DIM*DATA_DIM);

  int x = blockIdx.x * blockDim + threadIdx.x;
  int y = blockIdx.y * blockDim + threadIdx.y;
  int z = blockIdx.z * blockDim + threadIdx.z;

  int dx[6] = {-1,1,0,0,0,0};
  int dy[6] = {0,0,-1,1,0,0};
  int dz[6] = {0,0,0,0,-1,1};

  int3 pixel = {.x=x .y=y .z=z};
  if(region[pixel.z * DATA_DIM*DATA_DIM + pixel.y*DATA_DIM + pixel.x]==2){
    region[candidate.z * DATA_DIM*DATA_DIM + candidate.y*DATA_DIM + candidate.x] = 1;
    for(int n = 0; n < 6; n++){
      int3 candidate = pixel;
      candidate.x += dx[n];
      candidate.y += dy[n];
      candidate.z += dz[n];

      if(!inside(candidate)){
        continue;
      }

      if(region[candidate.z * DATA_DIM*DATA_DIM + candidate.y*DATA_DIM + candidate.x]){
        continue;
      }

      if(similar(data, pixel, candidate)){
        region[candidate.z * DATA_DIM*DATA_DIM + candidate.y*DATA_DIM + candidate.x] = 2;
        finished[0]=0;
      }
    }
  }
}


__global__ void region_grow_kernel_shared(unsigned char* data, unsigned char* region, int* finished){
