Latest cleanup and improvements of opencl backend.

Conflicts:
	sc/source/core/opencl/openclwrapper.cxx

Change-Id: I3fdc90570e90a156ccecb511fc04b473752018bd

1 files changed, 93 insertions, 88 deletions

diff --git a/sc/source/core/opencl/oclkernels.hxx b/sc/source/core/opencl/oclkernels.hxx
index 3269f3a3eba2..6c9012641b54 100644..100755
--- a/sc/source/core/opencl/oclkernels.hxx
+++ b/sc/source/core/opencl/oclkernels.hxx
@@ -6,153 +6,158 @@
  * License, v. 2.0. If a copy of the MPL was not distributed with this
  * file, You can obtain one at http://mozilla.org/MPL/2.0/.
  */
+
 #ifndef _OCL_KERNEL_H_
 #define _OCL_KERNEL_H_
 
 #ifndef USE_EXTERNAL_KERNEL
 #define KERNEL( ... )# __VA_ARGS__
 
-
 /////////////////////////////////////////////
 const char *kernel_src = KERNEL(
 __kernel void hello(__global uint *buffer)
 
 {
-size_t idx = get_global_id(0);
-
-buffer[idx]=idx;
-
+    size_t idx = get_global_id(0);
+    buffer[idx]=idx;
 }
 
 __kernel void oclformula(__global float *data,
-					   const uint type)
+                       const uint type)
 {
-	const unsigned int i = get_global_id(0);
-
-	switch (type)
-	{
-		case 0:          //MAX
-		{
-			//printf("%i ?%f>%f\n",i,data[2*i],data[2*i+1]);
-			if(data[2*i]>data[2*i+1])
-				data[i] = data[2*i];
-			else
-				data[i] = data[2*i+1];
-			break;
-		}
-		case 1:          //MIN
-		{
-			//printf("%d ?%d<%d\n",i,data[2*i],data[2*i+1]);
-			if(data[2*i]<data[2*i+1])
-				data[i] = data[2*i];
-			else
-				data[i] = data[2*i+1];
-			break;
-		}
-		case 2:          //SUM
-		case 3:          //AVG
-		{
-			//printf("%d %d+%d\n",i,data[2*i],data[2*i+1]);
-			data[i] = data[2*i] + data[2*i+1];
-			break;
-		}
-		default:
-			break;
-
-	}
+    const unsigned int i = get_global_id(0);
+
+    switch (type)
+    {
+        case 0:          //MAX
+        {
+            //printf("%i ?%f>%f\n",i,data[2*i],data[2*i+1]);
+            if(data[2*i]>data[2*i+1])
+                data[i] = data[2*i];
+            else
+                data[i] = data[2*i+1];
+            break;
+        }
+        case 1:          //MIN
+        {
+            //printf("%d ?%d<%d\n",i,data[2*i],data[2*i+1]);
+            if(data[2*i]<data[2*i+1])
+                data[i] = data[2*i];
+            else
+                data[i] = data[2*i+1];
+            break;
+        }
+        case 2:          //SUM
+        case 3:          //AVG
+        {
+            //printf("%d %d+%d\n",i,data[2*i],data[2*i+1]);
+            data[i] = data[2*i] + data[2*i+1];
+            break;
+        }
+        default:
+            break;
+
+    }
 }
 
 
 __kernel void oclSignedAdd(__global float *ltData,__global float *rtData,__global float *otData)
 {
-	const unsigned int id = get_global_id(0);
-	otData[id] = ltData[id] + rtData[id];
+    const unsigned int id = get_global_id(0);
+    otData[id] = ltData[id] + rtData[id];
 }
 
 
 __kernel void oclSignedSub(__global float *ltData,__global float *rtData,__global float *otData)
 {
-	const unsigned int id = get_global_id(0);
-	otData[id] = ltData[id] - rtData[id];
+    const unsigned int id = get_global_id(0);
+    otData[id] = ltData[id] - rtData[id];
 
 }
 
 __kernel void oclSignedMul(__global float *ltData,__global float *rtData,__global float *otData)
 {
-	int id = get_global_id(0);
-	otData[id] =ltData[id] * rtData[id];
+    int id = get_global_id(0);
+    otData[id] =ltData[id] * rtData[id];
 }
 
 __kernel void oclSignedDiv(__global float *ltData,__global float *rtData,__global float *otData)
 {
-	const unsigned int id = get_global_id(0);
-	otData[id] = ltData[id] / rtData[id];
+    const unsigned int id = get_global_id(0);
+    otData[id] = ltData[id] / rtData[id];
 }
 
 __kernel void oclFormulaMin(__global float *input,__global int *start,__global int *end,__global float *output)
 {
-	const unsigned int id = get_global_id(0);
-	int i=0;
-	unsigned int startFlag = start[id];
-	unsigned int endFlag = end[id];
-	float min = input[startFlag];
-	for(i=startFlag;i<=endFlag;i++)
-	{
-		if(input[i]<min)
-			min = input[i];
-	}
-	output[id] = min;
+    const unsigned int id = get_global_id(0);
+    int i=0;
+    unsigned int startFlag = start[id];
+    unsigned int endFlag = end[id];
+    float min = input[startFlag];
+    for(i=startFlag;i<=endFlag;i++)
+    {
+        if(input[i]<min)
+            min = input[i];
+    }
+    output[id] = min;
 
 }
 
 __kernel void oclFormulaMax(__global float *input,__global int *start,__global int *end,__global float *output)
 {
-	const unsigned int id = get_global_id(0);
-	int i=0;
-	unsigned int startFlag = start[id];
-	unsigned int endFlag = end[id];
-	float max = input[startFlag];
-	for(i=startFlag;i<=endFlag;i++)
-	{
-		if(input[i]>max)
-			max = input[i];
-	}
-	output[id] = max;
+    const unsigned int id = get_global_id(0);
+    int i=0;
+    unsigned int startFlag = start[id];
+    unsigned int endFlag = end[id];
+    float max = input[startFlag];
+    for(i=startFlag;i<=endFlag;i++)
+    {
+        if(input[i]>max)
+            max = input[i];
+    }
+    output[id] = max;
 
 }
-
-__kernel void oclFormulaSum(__global float *data,
-					   const uint type)
+//Sum
+__kernel void oclFormulaSum(__global float *input,__global int *start,__global int *end,__global float *output)
 {
-
+    const unsigned int nId = get_global_id(0);
+    float fSum = 0.0f;
+    for(int i = start[nId]; i<=end[nId]; i++)
+        fSum += input[i];
+    output[nId] = fSum ;
 }
-
-__kernel void oclFormulaCount(__global float *data,
-					   const uint type)
+//Count
+__kernel void oclFormulaCount(__global int *start,__global int *end,__global float *output)
 {
-
+    const unsigned int nId = get_global_id(0);
+    output[nId] = end[nId] - start[nId] + 1;
 }
 
 __kernel void oclFormulaAverage(__global float *input,__global int *start,__global int *end,__global float *output)
 {
-	const unsigned int id = get_global_id(0);
-	int i=0;
-	float sum=0;
-	for(i = start[id];i<=end[id];i++)
-		sum += input[i];
-	output[id] = sum / (end[id]-start[id]+1);
+    const unsigned int id = get_global_id(0);
+    int i=0;
+    float sum=0;
+    for(i = start[id];i<=end[id];i++)
+        sum += input[i];
+    output[id] = sum / (end[id]-start[id]+1);
 
 }
 
-
-__kernel void oclFormulaSumproduct(__global float *data,
-					   const uint type)
+//Sumproduct
+__kernel void oclFormulaSumproduct(__global float *firstCol,__global float *secondCol,__global int *start,__global int *end,__global float *output)
 {
-
+    const int nId = get_global_id(0);
+    int nCount     = start[nId] - end[nId] + 1;
+    int nStartA  = start[nId*2];
+    int nStartB  = start[nId*2+1];
+    for(int i = 0; i<nCount; i++)
+        output[nId] += firstCol[nStartA+i]*secondCol[nStartB+i];
 }
 
 __kernel void oclFormulaMinverse(__global float *data,
-					   const uint type)
+                       const uint type)
 {
 
 }