 }
 
 // Indexing function (note the argument order)
-__global__ int index(int z, int y, int x){
+__device__ int index(int z, int y, int x){
   return z * DATA_DIM*DATA_DIM + y*DATA_DIM + x;
 }
 
 // Trilinear interpolation
-__global__ float value_at(float3 pos, unsigned char* data){
+__device__ float value_at(float3 pos, unsigned char* data){
   if(!inside(pos)){
     return 0;
   }

    color += value_at(pos, data)*(0.01 + r) ;       // Update the color based on data value, and if we're in the region
  }

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

  return c;
}

__device__ float3 normalize(float3 v){
  float l = sqrt(v.x*v.x + v.y*v.y + v.z*v.z);
  v.x /= l;
  v.y /= l;
