3D interpolation example

openCL/interpolation.cl

@@ -1,36 +1,47 @@
+#define ABS(i) ( (i<0.0f) ? (-i):(i) )
 
 void __kernel interpolate(image3d_t volume,
-                          float* img,
+						  sampler_t sampler,
+						  global float* img,
                           int img_width,
                           int img_height,
-                          float3* point,
-                          float3* norm)
-                       {
+                          global float* point,
+                          global float3* norm)
+{                       
 int pos_x = get_global_id(0);
 int pos_y = get_global_id(1);
-if (pos_x>=img_width)||(pos_y>img_height)
+if ((pos_x>=img_width)||(pos_y>img_height))
     return;
-float center_x = get_global_id(0)/2.0f;
-float center_y = get_global_id(1)/2.0f;
 float3 n_norm = normalize(norm[0]);
-float3 u_norm, v_norm
+float3 u_norm, v_norm;
+float3 pos;
 float nx = n_norm.x,
       ny = n_norm.y,
       nz = n_norm.z;
-float ax = abs(nx),
-      ay = abs(ny),
-      az = abs(nz);
+float ax = ABS(nx),
+      ay = ABS(ny),
+      az = ABS(nz);
 
-if (ax>=az) && (ay>=az)       //z smallest
+if ((ax>=az) && (ay>=az))       //z smallest
 	u_norm = (float3)( -ny, nx, 0.0f);
-else if  (ax>=ay) && (az>=ay) //y smallest
+else if  ((ax>=ay) && (az>=ay)) //y smallest
     u_norm = (float3)( -nz, 0.0f, nx);
-else if  (ay>=ax) && (az>=ax) //x smallest
+else if  ((ay>=ax) && (az>=ax)) //x smallest
     u_norm = (float3)( 0.0f, -nz, ny);
-
-//define 2 largest components a,b,c
-//create  orthogonal vector -b a 0
-// create third vector by cross product
-//use them to
-
+v_norm = cross(n_norm,u_norm);
+// u_norm, v_norm, n_norm is a direct orthonormal ref
+float3 tx=(float3)(u_norm.x,v_norm.x,n_norm.x);
+float3 ty=(float3)(u_norm.y,v_norm.y,n_norm.y);
+float3 tz=(float3)(u_norm.z,v_norm.z,n_norm.z);
+//transposed version
+float3 pos_uvn = (float3)(2.0f*((float)pos_x/(float)img_width)-1.0f,
+			  	   	   	  2.0f*((float)pos_y/(float)img_height)-1.0f,
+			  	   	   	   0.0f);
+float4 pos_xyz =  (float4)(dot(tx,pos_uvn)+point[0],
+						dot(ty,pos_uvn)+point[1],
+						dot(tz,pos_uvn)+point[2],
+						0.0f);
+
+float4 res = read_imagef(volume, sampler, pos_xyz);
+img[pos_x+img_width*pos_y] = res.x;
 }

openCL/interpolation.py

@@ -2,25 +2,42 @@
 import pyopencl,pyopencl.array
 import numpy
 ctx = pyopencl.create_some_context()
-queue=pyopencl.CommandQueue(ctx)
+queue = pyopencl.CommandQueue(ctx, properties=pyopencl.command_queue_properties.PROFILING_ENABLE)
 x,y,z = numpy.ogrid[-10:10:0.1,-10:10:0.1,-10:10:0.1]
 r=numpy.sqrt(x*x+y*y+z*z)
 data = ((x * x - y * y + z * z) * numpy.exp(-r)).astype("float32")
 gpu_vol = pyopencl.image_from_array(ctx, data, 1)
-gpu_img = pyopencl.array.empty(queue, (200, 200), numpy.float32)
+shape = (200,200)
+img = numpy.empty(shape,dtype=numpy.float32)
+gpu_img = pyopencl.array.empty(queue, shape, numpy.float32)
 prg = open("interpolation.cl").read()
 sampler = pyopencl.Sampler(ctx,
-                           False, # Not normalized coordinates
+                           True, # normalized coordinates
                            pyopencl.addressing_mode.CLAMP_TO_EDGE,
                            pyopencl.filter_mode.LINEAR)
 
 prg = pyopencl.Program(ctx, prg).build()
-n = pyopencl.array.to_device(numpy.array([1,0,0],dtype=numpy.float32))
-c = pyopencl.array.to_device(numpy.array([100, 100, 100], dtype=numpy.float32))
-
-prg.interpolate(queue, (256, 256), (16, 16), gpu_vol, gpu_img.data,
-                numpy.int32(200), numpy.int32(200), n.data, c.data)
+n = pyopencl.array.to_device(queue, numpy.array([1, 1, 1], dtype=numpy.float32))
+c = pyopencl.array.to_device(queue, numpy.array([0.5, 0.5, 0.5], dtype=numpy.float32))
+prg.interpolate(queue, (256, 256), (16, 16), gpu_vol, sampler, gpu_img.data,
+                numpy.int32(200), numpy.int32(200), c.data, n.data)
 img = gpu_img.get()
+
+
+#timing:
+evt = []
+evt.append(pyopencl.enqueue_copy(queue, n.data, (2.0*numpy.random.random(3)-1).astype(numpy.float32)))
+evt.append(pyopencl.enqueue_copy(queue, c.data, numpy.random.random(3).astype(numpy.float32)))
+evt.append(prg.interpolate(queue, (256, 256), (16, 16), gpu_vol, sampler, gpu_img.data,
+                numpy.int32(shape[1]), numpy.int32(shape[0]), c.data, n.data))
+evt.append(pyopencl.enqueue_copy(queue, img, gpu_img.data))
+print("Timings: %.3fms %.3fms %.3fms %.3fms total: %.3fms" % (1e-6 * (evt[0].profile.end - evt[0].profile.start),
+                                1e-6 * (evt[1].profile.end - evt[1].profile.start),
+                                1e-6 * (evt[2].profile.end - evt[2].profile.start),
+                                1e-6 * (evt[3].profile.end - evt[3].profile.start),
+                                1e-6 * (evt[-1].profile.end - evt[0].profile.start)))
+
+
 from pylab import *
 imshow(img)
 show()