Adding floating point header to cl_read to avoid OpenCL compiler warn…

…ings.

source/pysph/base/kernels.h

@@ -2,7 +2,7 @@
 #ifndef CL_KERNEL_H
 #define CL_KERNEL_H
 
-#define PI 2.0f*acos(0.0f)
+#define PI 2.0F*acos(0.0f)
 
 enum KernelType {
     CUBIC_SPLINE = 1,
@@ -20,16 +20,16 @@ enum KernelType {
 
 REAL cubic_spline_fac(unsigned int dim, REAL h)
 {
-  REAL fac = 0.0;
+  REAL fac = 0.0F;
 
   if ( dim == 1 )
-    fac = ( 2.0/3.0 ) / (h);
+    fac = ( 2.0F/3.0F ) / (h);
 
   if (dim == 2 )
-    fac = ( 10.0 ) / ( 7.0*PI ) / ( h*h );
+    fac = ( 10.0F ) / ( 7.0F*PI ) / ( h*h );
 
   if ( dim == 3 )
-    fac = 1.0 /( PI*h*h*h );
+    fac = 1.0F /( PI*h*h*h );
 
   return fac;
 }
@@ -40,20 +40,20 @@ REAL cubic_spline_function(REAL4 pa, REAL4 pb, unsigned int dim)
 
   REAL4 rab = pa - pb;
 
-  REAL h = 0.5 * ( pa.s3 + pb.s3 );
+  REAL h = 0.5F * ( pa.s3 + pb.s3 );
   REAL q = length(rab)/h;
 
-  REAL val = 0.0;
+  REAL val = 0.0F;
   REAL fac = cubic_spline_fac(dim, h);
 
-  if ( q >= 0.0 )
-    val = 1.0 - 1.5 * (q*q) * (1.0 - 0.5 * q);
+  if ( q >= 0.0F )
+    val = 1.0F - 1.5F * (q*q) * (1.0F - 0.5F * q);
 
-  if ( q >= 1.0 )
-    val = 0.25 * (2.0-q) * (2.0-q) * (2.0-q);
+  if ( q >= 1.0F )
+    val = 0.25F * (2.0F-q) * (2.0F-q) * (2.0F-q);
 
-  if ( q >= 2.0 )
-      val = 0.0;
+  if ( q >= 2.0F )
+      val = 0.0F;
 
   return val * fac;
 }
@@ -63,23 +63,23 @@ void cubic_spline_gradient(REAL4 pa, REAL4 pb, REAL4* grad,
 {
   REAL4 rab = pa - pb;
 
-  REAL h = 0.5 * ( pa.s3 + pb.s3 );
+  REAL h = 0.5F * ( pa.s3 + pb.s3 );
   REAL q = length(rab)/h;
 
-  REAL val = 0.0;
+  REAL val = 0.0F;
   REAL fac = cubic_spline_fac(dim, h);
 
-  if ( q == 0.0 )
-    val = 0.0;
+  if ( q == 0.0F )
+    val = 0.0F;
 
-  if ( q > 0.0 )
-    val = 3.0*(0.75*q - 1.0)/(h*h);
+  if ( q > 0.0F )
+    val = 3.0F*(0.75F*q - 1.0F)/(h*h);
 
-  if ( q >= 1.0 )
-    val = -0.75 * (2.0-q) * (2.0-q)/(h * length(rab));
+  if ( q >= 1.0F )
+    val = -0.75F * (2.0F-q) * (2.0F-q)/(h * length(rab));
 
-  if ( q >= 2.0 )
-    val = 0.0;
+  if ( q >= 2.0F )
+    val = 0.0F;
 
   val *= fac;
 

source/pysph/solver/cl_common.cl

@@ -8,9 +8,9 @@ __kernel void set_tmp_to_zero(__global REAL* tmpx,
   unsigned int work_dim = get_work_dim();
   unsigned int gid = get_gid(work_dim);
 
-  tmpx[gid] = 0.0;
-  tmpy[gid] = 0.0;
-  tmpz[gid] = 0.0;
+  tmpx[gid] = 0.0F;
+  tmpy[gid] = 0.0F;
+  tmpz[gid] = 0.0F;
 
 }
 

source/pysph/solver/cl_utils.py

@@ -75,16 +75,17 @@ def cl_read(filename, precision='double'):
 
     if precision == 'single':
         typ = 'float'
-        hdr = ""
+        hdr = "#define F f \n"
     else:
         typ = 'double'
-        hdr = "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n" 
+        hdr = "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n"
+        hdr += '#define F \n'
 
     for x in ('', '2', '3', '4', '8'):
         hdr += '#define REAL%s %%(typ)s%s\n'%(x, x)
 
     hdr = hdr%(dict(typ=typ))
-        
+
     return hdr + src
 
 def get_real(val, precision):

source/pysph/sph/funcs/density_funcs.cl

@@ -1,7 +1,6 @@
 #include "cl_common.h"
 #include "cl_common.cl"
 #include "kernels.h"
-
 
 __kernel void SPHRho(int const kernel_type, int const dim, int const nbrs,
 		     __global REAL* d_x, __global REAL* d_y, 
@@ -17,7 +16,7 @@ __kernel void SPHRho(int const kernel_type, int const dim, int const nbrs,
   unsigned int gid = get_gid(work_dim);
 
   REAL4 pa = (REAL4)( d_x[gid], d_y[gid], d_z[gid], d_h[gid] );
-  REAL wmb = 0.0;
+  REAL wmb = 0.0F ;
   REAL w;
 
   for (unsigned int i = 0; i < nbrs; ++i)

source/pysph/sph/funcs/external_force.cl

@@ -42,7 +42,7 @@ __kernel void NBodyForce(int const nbrs, int const self,
       REAL4 pb = (REAL4)( s_x[i], s_y[i], s_z[i], s_h[i] );
       rba = pb - pa;
 
-      invr = 1.0/( length(rba) + eps );
+      invr = 1.0F/( length(rba) + eps );
       invr *= ( invr * invr );
 
       force_mag = s_m[i] * invr;