 #define DATA_DIM 512
 // image is 2D, total size is IMAGE_DIM x IMAGE_DIM
 #define IMAGE_DIM 512
+
+texture<int, cudaTextureType3D, cudaReadModeElementType> data_texture;
+texture<int, cudaTextureType3D, cudaReadModeElementType> region_texture;
 
 
 // Stack for the serial region growing

  image[(y+(IMAGE_DIM/2)) * IMAGE_DIM + (x+(IMAGE_DIM/2))] = color > 255 ? 255 : color;
}


__global__ void raycast_kernel_texture(unsigned char* image){

}


    cudaMemcpy( finished, finished_device, sizeof(int), cudaMemcpyDeviceToHost);
  }
  cudaMemcpy( region, region_device, DATA_DIM*DATA_DIM*DATA_DIM*sizeof(unsigned char), cudaMemcpyDeviceToHost);
  cudaFree(data_device);
  cudaFree(region_device);
  cudaFree(finished_device);
  return region;
}


unsigned char* grow_region_gpu_shared(unsigned char* data){
