diff --git a/sc/source/core/opencl/formulagroupcl.cxx b/sc/source/core/opencl/formulagroupcl.cxx index af81125e9c99..a8dc885d1aa0 100644 --- a/sc/source/core/opencl/formulagroupcl.cxx +++ b/sc/source/core/opencl/formulagroupcl.cxx @@ -1642,66 +1642,6 @@ public: { i += rxSubArgument->Marshal(k, argno + i, nVectorWidth, pProgram); } - if (dynamic_cast(mpCodeGen.get())) - { - openclwrapper::KernelEnv kEnv; - openclwrapper::setKernelEnv(&kEnv); - cl_int err; - cl_mem pClmem2; - - std::vector vclmem; - for (const auto& rxSubArgument : mvSubArguments) - { - if (VectorRef* VR = dynamic_cast(rxSubArgument.get())) - vclmem.push_back(VR->GetCLBuffer()); - else - vclmem.push_back(nullptr); - } - pClmem2 = clCreateBuffer(kEnv.mpkContext, CL_MEM_READ_WRITE, - sizeof(double) * nVectorWidth, nullptr, &err); - if (CL_SUCCESS != err) - throw OpenCLError("clCreateBuffer", err, __FILE__, __LINE__); - SAL_INFO("sc.opencl", "Created buffer " << pClmem2 << " size " << sizeof(double) << "*" << nVectorWidth << "=" << (sizeof(double)*nVectorWidth)); - - std::string kernelName = "GeoMean_reduction"; - cl_kernel redKernel = clCreateKernel(pProgram, kernelName.c_str(), &err); - if (err != CL_SUCCESS) - throw OpenCLError("clCreateKernel", err, __FILE__, __LINE__); - SAL_INFO("sc.opencl", "Created kernel " << redKernel << " with name " << kernelName << " in program " << pProgram); - - // set kernel arg of reduction kernel - for (size_t j = 0; j < vclmem.size(); j++) - { - SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << j << ": " << (vclmem[j] ? "cl_mem" : "double") << ": " << vclmem[j]); - err = clSetKernelArg(redKernel, j, - vclmem[j] ? sizeof(cl_mem) : sizeof(double), - static_cast(&vclmem[j])); - if (CL_SUCCESS != err) - throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); - } - SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << vclmem.size() << ": cl_mem: " << pClmem2); - err = clSetKernelArg(redKernel, vclmem.size(), sizeof(cl_mem), static_cast(&pClmem2)); - if (CL_SUCCESS != err) - throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); - - // set work group size and execute - size_t global_work_size[] = { 256, static_cast(nVectorWidth) }; - size_t const local_work_size[] = { 256, 1 }; - SAL_INFO("sc.opencl", "Enqueuing kernel " << redKernel); - err = clEnqueueNDRangeKernel(kEnv.mpkCmdQueue, redKernel, 2, nullptr, - global_work_size, local_work_size, 0, nullptr, nullptr); - if (CL_SUCCESS != err) - throw OpenCLError("clEnqueueNDRangeKernel", err, __FILE__, __LINE__); - err = clFinish(kEnv.mpkCmdQueue); - if (CL_SUCCESS != err) - throw OpenCLError("clFinish", err, __FILE__, __LINE__); - - // Pass pClmem2 to the "real" kernel - SAL_INFO("sc.opencl", "Kernel " << k << " arg " << argno << ": cl_mem: " << pClmem2); - err = clSetKernelArg(k, argno, sizeof(cl_mem), static_cast(&pClmem2)); - if (CL_SUCCESS != err) - throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); - } if (OpSumIfs* OpSumCodeGen = dynamic_cast(mpCodeGen.get())) { openclwrapper::KernelEnv kEnv; @@ -2303,9 +2243,9 @@ DynamicKernelSoPArguments::DynamicKernelSoPArguments(const ScCalcConfig& config, case ocGauss: mvSubArguments.push_back(SoPHelper(mCalcConfig, ts, ft->Children[i], std::make_shared(), nResultSize)); break; - /*case ocGeoMean: - mvSubArguments.push_back(SoPHelper(mCalcConfig, ts, ft->Children[i], std::make_sharedChildren[i], std::make_shared(), nResultSize)); + break; case ocHarMean: mvSubArguments.push_back(SoPHelper(mCalcConfig, ts, ft->Children[i], std::make_shared(), nResultSize)); break; diff --git a/sc/source/core/opencl/op_statistical.cxx b/sc/source/core/opencl/op_statistical.cxx index 598f81df9d83..7095a23e7066 100644 --- a/sc/source/core/opencl/op_statistical.cxx +++ b/sc/source/core/opencl/op_statistical.cxx @@ -477,160 +477,26 @@ vSubArguments) } void OpGeoMean::GenSlidingWindowFunction( - outputstream &ss, const std::string &sSymName, SubArguments &vSubArguments) + outputstream &ss, const std::string &sSymName, SubArguments & +vSubArguments) { - ss << "__kernel void "; - ss << "GeoMean_reduction( "; - for (size_t i = 0; i < vSubArguments.size(); i++) - { - if (i) - ss << ","; - vSubArguments[i]->GenSlidingWindowDecl(ss); - } - ss << ", __global double *result)\n"; - ss << "{\n"; - ss << " double tmp =0;\n"; - ss << " int count = 0;\n"; - ss << " int i ;\n"; - GenTmpVariables(ss,vSubArguments); - ss << " double current_sum = 0.0;\n"; - ss << " int windowSize;\n"; - ss << " int arrayLength;\n"; - ss << " int current_count = 0;\n"; - ss << " int writePos = get_group_id(1);\n"; - ss << " int lidx = get_local_id(0);\n"; - ss << " __local double shm_buf[256];\n"; - ss << " __local int count_buf[256];\n"; - ss << " int loop;\n"; - ss << " int offset;\n"; - ss << " barrier(CLK_LOCAL_MEM_FENCE);\n"; - - for(const DynamicKernelArgumentRef & rArg : vSubArguments) - { - assert(rArg->GetFormulaToken()); - - if(rArg->GetFormulaToken()->GetType() == - formula::svDoubleVectorRef) - { - FormulaToken *tmpCur = rArg->GetFormulaToken(); - const formula::DoubleVectorRefToken*pCurDVR= static_cast(tmpCur); - size_t nCurWindowSize = pCurDVR->GetArrayLength() < - pCurDVR->GetRefRowSize() ? pCurDVR->GetArrayLength(): - pCurDVR->GetRefRowSize() ; - - if (pCurDVR->IsStartFixed() && pCurDVR->IsEndFixed()) - ss << " offset = 0;\n"; - else if (!pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed()) - ss << " offset = get_group_id(1);\n"; - else - throw Unhandled(__FILE__, __LINE__); - ss << " windowSize = "; - ss << nCurWindowSize; - ss << ";\n"; - ss << " arrayLength = "; - ss << pCurDVR->GetArrayLength(); - ss << ";\n"; - ss << " loop = arrayLength/512 + 1;\n"; - ss << " for (int l=0; l0; i/=2) {\n"; - ss << " if (lidx < i)\n"; - ss << " {\n"; - ss << " shm_buf[lidx] += shm_buf[lidx + i];\n"; - ss << " count_buf[lidx] += count_buf[lidx + i];\n"; - ss << " }\n"; - ss << " barrier(CLK_LOCAL_MEM_FENCE);\n"; - ss << " }\n"; - ss << " if (lidx == 0)\n"; - ss << " {\n"; - ss << " current_sum += shm_buf[0];\n"; - ss << " current_count += count_buf[0];\n"; - ss << " }\n"; - // ss << "if(writePos == 14 && lidx ==0)\n"; - //ss <<"printf(\"\\n********************sum is %f, count is%d\",current_sum,current_count);\n"; - ss << " barrier(CLK_LOCAL_MEM_FENCE);\n"; - ss << " }\n"; - }else - { - ss << " if (lidx == 0)\n"; - ss << " {\n"; - ss << " tmp0 ="; - if(rArg->GetFormulaToken()->GetType() == formula::svSingleVectorRef) - { - rArg->GenDeclRef(ss); - ss << "[writePos];\n"; - } - else - { - rArg->GenDeclRef(ss); - ss <<";\n"; - //ss <<"printf(\"\\n********************tmp0 is %f\",tmp0);\n"; - } - ss << " if(!isnan(tmp0))\n"; - ss << " {\n"; - ss << " current_sum += log(tmp0);\n"; - ss << " current_count++;\n"; - ss << " }\n"; - ss << " }\n"; - } - } - - ss << " if (lidx == 0)\n"; - ss << " result[writePos] = exp(current_sum/current_count);\n"; - ss << "}\n"; - + CHECK_PARAMETER_COUNT( 1, 30 ); GenerateFunctionDeclaration( sSymName, vSubArguments, ss ); ss << "{\n"; - ss <<" int gid0=get_global_id(0);\n"; - ss << " double tmp =0;\n"; - ss << " tmp ="; - vSubArguments[0]->GenDeclRef(ss); - ss << "[gid0];\n"; - ss << " return tmp;\n"; + ss << " int gid0 = get_global_id(0);\n"; + ss << " double nVal=0.0;\n"; + ss << " double tmp = 0;\n"; + ss << " int length;\n"; + ss << " int totallength=0;\n"; + GenerateRangeArgs( vSubArguments, ss, SkipEmpty, + " if( arg < 0 )\n" + " return CreateDoubleError(IllegalArgument);\n" + " if( arg == 0 )\n" + " return 0;\n" + " nVal += log(arg);\n" + " ++totallength;\n" + ); + ss << " return exp(nVal/totallength);\n"; ss << "}"; } @@ -652,8 +518,7 @@ vSubArguments) " nVal += (1.0 / arg);\n" " ++totallength;\n" ); - ss << " tmp = totallength/nVal;\n"; - ss << " return tmp;\n"; + ss << " return totallength/nVal;\n"; ss << "}"; } diff --git a/sc/source/core/opencl/op_statistical.hxx b/sc/source/core/opencl/op_statistical.hxx index 64ca0c1cdf7c..2741efe49656 100644 --- a/sc/source/core/opencl/op_statistical.hxx +++ b/sc/source/core/opencl/op_statistical.hxx @@ -125,13 +125,13 @@ public: virtual std::string BinFuncName() const override { return "Gauss"; } }; -class OpGeoMean: public CheckVariables +class OpGeoMean: public Normal { public: - OpGeoMean(): CheckVariables() {} virtual void GenSlidingWindowFunction(outputstream &ss, const std::string &sSymName, SubArguments &vSubArguments) override; virtual std::string BinFuncName() const override { return "GeoMean"; } + virtual bool canHandleMultiVector() const override { return true; } }; class OpHarMean: public Normal