Changeset 4090 for branches/wezzy/linbox/algorithms/opencl-kernels/kernel_muladd_partial_8_dp.cl

Timestamp:

12/04/11 22:08:32 (19 months ago)

Author:

wezowicz

Message:

Fixed memory leak in opencl-domain-memory.inl.
Add OpenCL specialization to muladd() along with 8 more kernels.

File:

• branches/wezzy/linbox/algorithms/opencl-kernels/kernel_muladd_partial_8_dp.cl (modified) (2 diffs)

branches/wezzy/linbox/algorithms/opencl-kernels/kernel_muladd_partial_8_dp.cl

@@ -9 +9 @@
 #pragma OPENCL EXTENSION cl_khr_fp64 : enable
 
-__kernel void matrix_mul_kernel(__global double* C, __global double* A, __global double* B,
-                int width_A, int width_B, double mod){
+__kernel void matrix_mul_kernel(__global double* D, double alpha, __global double* A, __global double* B,
+                double beta, __global double* C, int width_A, int width_B, double mod){
         //Get Workgroup ID
         int bx = get_group_id(0);
@@ -64 +64 @@
         
         //Scale Dsub by alpha
-        Dsub = fmod((alpha * Dsub), mod);
+        Dsub = alpha * Dsub;
+        Dsub = fmod(Dsub, mod);
         
-        if(beta != 0.0){
-                //Add C scaled by beta to Dsub
-                double Csub = fmod((beta * C[d + ty * width_B + tx]), mod);
-                Dsub = fmod((Dsub + Csub), mod);
-        }
+        //Scalse Csub by beta
+        double Csub = C[d + ty * width_B + tx];
+        Csub = beta * Csub;
+        Csub = fmod(Csub, mod);
+        
+        //Add Dsub and Dsub
+        Dsub = Dsub + Csub;
+        Dsub = fmod(Dsub, mod);
         
         //Add the sum to the appropriate spot