   //printf("Data created");
 
   gettimeofday(&start, NULL);
-  unsigned char* region = grow_region_gpu(data);
-  //unsigned char* region = grow_region_serial(data);
+  //unsigned char* region = grow_region_gpu(data);
+  unsigned char* region = grow_region_serial(data);
   gettimeofday(&end, NULL);
   print_time(start, end);
 
   //printf("Region grown");
 
   gettimeofday(&start, NULL);
-  unsigned char* image = raycast_gpu(data, region);
+  unsigned char* image = raycast_gpu_texture(data, region);
   //unsigned char* image = raycast_serial(data, region);
   gettimeofday(&end, NULL);
   print_time(start, end);

#define DATA_DIM 512
// image is 2D, total size is IMAGE_DIM x IMAGE_DIM
#define IMAGE_DIM 512


// Stack for the serial region growing



unsigned char* raycast_gpu_texture(unsigned char* data, unsigned char* region){
  unsigned char* image = (unsigned char*)malloc(sizeof(unsigned char)*IMAGE_DIM*IMAGE_DIM);

  //unsigned char *data_device;
  //unsigned char *region_device;
  unsigned char* image_device;
  //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**)&image_device,DATA_DIM*DATA_DIM*sizeof(unsigned char));
  //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( image_device, image, DATA_DIM*DATA_DIM*sizeof(unsigned char), cudaMemcpyHostToDevice);



  cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32,0,0,0,cudaChannelFormatKindFloat);

  cudaExtent extent = make_cudaExtent(512,512,512);

  cudaArray* data_array;
  cudaArray* region_array;

  cudaMalloc3DArray(&data_array, &channelDesc, extent, 0);
  cudaMalloc3DArray(&region_array, &channelDesc, extent, 0);

  cudaMemcpy3DParms copyDataParams    = {0};
  cudaMemcpy3DParms copyRegionParams  = {0};
  copyDataParams.srcPtr     = make_cudaPitchedPtr(data, 512*sizeof(unsigned char), 512, 512);
  copyRegionParams.srcPtr   = make_cudaPitchedPtr(region, 512*sizeof(unsigned char), 512, 512);
  copyDataParams.dstArray   = data_array;
  copyRegionParams.dstArray = data_array;
  copyDataParams.extent     = extent;
  copyRegionParams.extent   = extent;
  copyDataParams.kind       = cudaMemcpyHostToDevice;
  copyRegionParams.kind     = cudaMemcpyHostToDevice;

  cudaMemcpy3D(&copyDataParams);
  cudaMemcpy3D(&copyRegionParams);

  cudaBindTextureToArray(data_texture, data_array);
  cudaBindTextureToArray(region_texture, region_array);





  dim3 dimBlock( 32, 32 );
  dim3 dimGrid( 16, 16 );

  raycast_kernel_texture<<<dimGrid, dimBlock>>>(image_device);

  cudaMemcpy( image, image_device, DATA_DIM*DATA_DIM*sizeof(unsigned char), cudaMemcpyDeviceToHost);
  cudaFree(image_device);

  return image;
}


