Changing SPHFunction's cl_kernel in favor for cl_program.

Renaming step to step_array to avoid overriding OpenCL's step function.

source/pysph/solver/cl_common.h

@@ -4,6 +4,7 @@
 #define CL_COMMON
 
 #define TRUE 1
+#define FALSE 0
 
 typedef struct {
   int sizes[3];

source/pysph/solver/cl_integrator.py

@@ -208,9 +208,9 @@ def step(self, calcs, dt):
                     step_buffer = pa.get_cl_buffer(k_prop)
                     tmp_buffer = pa.get_cl_buffer('tmpx')
 
-                    self.program.step(queue, (np,1,1), (1,1,1),
-                                      current_buffer, step_buffer,
-                                      tmp_buffer, cl_dt)
+                    self.program.step_array(queue, (np,1,1), (1,1,1),
+                                            current_buffer, step_buffer,
+                                            tmp_buffer, cl_dt)
 
                     cl.enqueue_copy_buffer(queue, src=tmp_buffer,
                                            dst=current_buffer)
@@ -256,9 +256,9 @@ def final_step(self, calc, dt):
             k1_buffer = pa.get_cl_buffer(k_prop)
             tmp_buffer = pa.get_cl_buffer('tmpx')
 
-            self.program.step(queue, (np,1,1), (1,1,1),
-                              initial_buffer, k1_buffer,
-                              tmp_buffer, cl_dt).wait()
+            self.program.step_array(queue, (np,1,1), (1,1,1),
+                                    initial_buffer, k1_buffer,
+                                    tmp_buffer, cl_dt).wait()
 
             cl.enqueue_copy_buffer(queue, src=tmp_buffer,
                                    dst=initial_buffer).wait()

source/pysph/solver/integrator.cl

@@ -1,8 +1,8 @@
 #include "cl_common.h"
 
 // c[gid] = a[gid] + h*b[gid]
-__kernel void step(__global REAL* a, __global REAL* b, __global REAL* c,
-		   REAL const h)
+__kernel void step_array(__global REAL* a, __global REAL* b, __global REAL* c,
+			 REAL const h)
 {
   unsigned int work_dim = get_work_dim();
   unsigned int gid = get_gid(work_dim);

source/pysph/sph/funcs/density_funcs.pyx

@@ -78,10 +78,12 @@ cdef class SPHRho(CSPHFunctionParticle):
         args.append( self.dest.get_cl_buffer('tmpx') )
 
         # Enqueue the OpenCL kernel for execution
-        self.cl_kernel(queue, self.global_sizes, self.local_sizes,
-                       numpy.int32(kernel.get_type()), numpy.int32(kernel.dim),
-                       numpy.int32(self.source.get_number_of_particles()),
-                       *args).wait()
+        self.cl_program.SPHRho(
+            queue, self.global_sizes, self.local_sizes,
+            numpy.int32(kernel.get_type()),
+            numpy.int32(kernel.dim),
+            numpy.int32(self.source.get_number_of_particles()),
+            *args).wait()
 
 ################################################################################
 # `SPHDensityRate` class.

source/pysph/sph/funcs/external_force.pyx

@@ -51,8 +51,9 @@ cdef class GravityForce(SPHFunction):
         gy = get_real(self.gy, self.dest.cl_precision)
         gz = get_real(self.gz, self.dest.cl_precision)
 
-        self.cl_kernel(queue, self.global_sizes, self.local_sizes,
-                       gx, gy, gz, tmpx, tmpy, tmpz).wait()
+        self.cl_program.GravityForce(
+            queue, self.global_sizes, self.local_sizes,
+            gx, gy, gz, tmpx, tmpy, tmpz).wait()
 
 #############################################################################
 # `VectorForce` class.
@@ -235,9 +236,10 @@ cdef class NBodyForce(SPHFunctionParticle):
 
         eps = get_real(self.eps, self.dest.cl_precision)
 
-        self.cl_kernel(queue, self.global_sizes, self.local_sizes,
-                       numpy.int32(self.source.get_number_of_particles()),
-                       numpy.int32(self.dest.name==self.source.name),
-                       eps, *self.args).wait()
+        self.cl_program.NBodyForce(
+            queue, self.global_sizes, self.local_sizes,
+            numpy.int32(self.source.get_number_of_particles()),
+            numpy.int32(self.dest.name==self.source.name),
+            eps, *self.args).wait()
 
 ###########################################################################

source/pysph/sph/funcs/position_funcs.pyx

@@ -51,7 +51,8 @@ cdef class PositionStepping(SPHFunction):
         d_v = self.dest.get_cl_buffer('v')
         d_w = self.dest.get_cl_buffer('w')
 
-        self.cl_kernel(queue, self.global_sizes, self.local_sizes, d_u, d_v,
-                       d_w, tag, tmpx, tmpy, tmpz)           
+        self.cl_program.PositionStepping(
+            queue, self.global_sizes, self.local_sizes, d_u, d_v,
+            d_w, tag, tmpx, tmpy, tmpz)
 
 ##########################################################################

source/pysph/sph/sph_calc.pyx

@@ -341,14 +341,9 @@ class CLCalc(SPHCalc):
         self.prog = cl.Program(self.context, prog_src_file).build(
             build_options)
 
-        # find by name the OpenCL kernel for this operation
-        for cl_kernel in self.prog.all_kernels():
-            if cl_kernel.function_name == self.cl_kernel_function_name:
-                break
-
         # set the OpenCL kernel for each of the SPHFunctions
         for func in self.funcs:
-            func.setup_cl(cl_kernel, self.context)
+            func.setup_cl(self.prog, self.context)
 
     def sph(self):
         """ Evaluate the contribution from the sources on the

source/pysph/sph/sph_func.pxd

@@ -34,8 +34,9 @@ cdef class SPHFunction:
     cdef public str cl_kernel_src_file
     cdef public str cl_kernel_function_name
 
-    # OpenCL kernel name
+    # OpenCL program and kernel
     cdef public object cl_kernel
+    cdef public object cl_program
 
     # OpenCL kernel launch parameters
     cdef public tuple global_sizes

source/pysph/sph/sph_func.pyx

@@ -130,6 +130,7 @@ cdef class SPHFunction:
 
         self.cl_kernel_src_file = ''
         self.cl_kernel = object()
+        self.cl_program = object()
         self.context = object()
         self.args = []
 
@@ -228,7 +229,7 @@ cdef class SPHFunction:
         """
         pass
 
-    def setup_cl(self, object cl_kernel, object context):
+    def setup_cl(self, object program, object context):
         """ OpenCL setup for the function.
 
         You may determine the OpenCL kernel launch parameters from within
@@ -240,7 +241,7 @@ cdef class SPHFunction:
         local_sizes = (1, 1, 1)
 
         """
-        self.cl_kernel = cl_kernel
+        self.cl_program = program
         self.context = context
 
         self.set_cl_kernel_args()
@@ -257,8 +258,8 @@ cdef class SPHFunction:
         self.args.append( self.dest.get_cl_buffer('tmpy') )
         self.args.append( self.dest.get_cl_buffer('tmpz') )
 
-    def set_cl_kernel(self, object kernel):
-        self.cl_kernel = kernel
+    def set_cl_program(self, object program):
+        self.cl_program = program
 
 ################################################################################
 # `SPHFunctionParticle` class.

source/pysph/sph/tests/test_external_force.py

@@ -55,12 +55,9 @@ def setUp(self):
             self.q = q = cl.CommandQueue(ctx)
             pysph_root = solver.get_pysph_root()
 
-            src = solver.cl_read(path.join(pysph_root,
-                                           "sph/funcs/external_force.cl"),
-                                 precision=self.precision)
-
-            #src = open(path.join(pysph_root,
-            #                     'sph/funcs/external_force.cl')).read()
+            src = solver.cl_read(
+                path.join(pysph_root, "sph/funcs/external_force.cl"),
+                precision=self.precision)
 
             self.prog = cl.Program(ctx, src).build(solver.get_cl_include())
             pa.setup_cl(ctx, q)
@@ -124,7 +121,7 @@ def test_cl_eval(self):
 
             k = base.CubicSplineKernel()
 
-            func.setup_cl(self.prog.NBodyForce, self.ctx)
+            func.setup_cl(self.prog, self.ctx)
 
             func.cl_eval(self.q, self.ctx, k)