 
 
 unsigned char* raycast_gpu_texture(unsigned char* data, unsigned char* region){
-  return NULL;
+  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;
 }
 
 

  //printf("Data created");

  gettimeofday(&start, NULL);
  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_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

texture<int, cudaTextureType3D, cudaReadModeElementType> data_texture;
texture<int, cudaTextureType3D, cudaReadModeElementType> region_texture;


// Stack for the serial region growing
