summaryrefslogtreecommitdiff
path: root/sc/source/core/opencl/formulagroupcl.cxx
diff options
context:
space:
mode:
authorLuboš Luňák <l.lunak@collabora.com>2022-09-20 18:08:42 +0200
committerLuboš Luňák <l.lunak@collabora.com>2022-09-21 10:24:26 +0200
commitbdb576ea889d944b2aa9d3dfaa8fbd7c07415c82 (patch)
tree9d6c3275f96367242e52af9946c8c241e710e484 /sc/source/core/opencl/formulagroupcl.cxx
parentb764e5e0532a9c017f871b60026e507f283eb133 (diff)
fix/simplify opencl GEOMEAN()
I don't quite see why this one would need such a special handling, when most other functions (e.g. the very similar HARMEAN()) can do with just generic handling. Change-Id: I31f3772ffdf9540178a42f11ae4e376023ad2413 Reviewed-on: https://gerrit.libreoffice.org/c/core/+/140257 Tested-by: Jenkins Reviewed-by: Luboš Luňák <l.lunak@collabora.com>
Diffstat (limited to 'sc/source/core/opencl/formulagroupcl.cxx')
-rw-r--r--sc/source/core/opencl/formulagroupcl.cxx66
1 files changed, 3 insertions, 63 deletions
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<OpGeoMean*>(mpCodeGen.get()))
- {
- openclwrapper::KernelEnv kEnv;
- openclwrapper::setKernelEnv(&kEnv);
- cl_int err;
- cl_mem pClmem2;
-
- std::vector<cl_mem> vclmem;
- for (const auto& rxSubArgument : mvSubArguments)
- {
- if (VectorRef* VR = dynamic_cast<VectorRef*>(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<void*>(&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<void*>(&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<size_t>(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<void*>(&pClmem2));
- if (CL_SUCCESS != err)
- throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__);
- }
if (OpSumIfs* OpSumCodeGen = dynamic_cast<OpSumIfs*>(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<OpGauss>(), nResultSize));
break;
- /*case ocGeoMean:
- mvSubArguments.push_back(SoPHelper(mCalcConfig, ts, ft->Children[i], std::make_shared<OpGeoMean));
- break;*/
+ case ocGeoMean:
+ mvSubArguments.push_back(SoPHelper(mCalcConfig, ts, ft->Children[i], std::make_shared<OpGeoMean>(), nResultSize));
+ break;
case ocHarMean:
mvSubArguments.push_back(SoPHelper(mCalcConfig, ts, ft->Children[i], std::make_shared<OpHarMean>(), nResultSize));
break;