diff options
author | Luboš Luňák <l.lunak@collabora.com> | 2022-09-20 18:08:42 +0200 |
---|---|---|
committer | Luboš Luňák <l.lunak@collabora.com> | 2022-09-21 10:24:26 +0200 |
commit | bdb576ea889d944b2aa9d3dfaa8fbd7c07415c82 (patch) | |
tree | 9d6c3275f96367242e52af9946c8c241e710e484 /sc/source/core/opencl/formulagroupcl.cxx | |
parent | b764e5e0532a9c017f871b60026e507f283eb133 (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.cxx | 66 |
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; |