ODFstudio/odfstudio
2
0
Fork 0
Code Issues Packages Activity

fix/simplify opencl GEOMEAN()

Browse Source 
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>
This commit is contained in:
Luboš Luňák
2022-09-21 10:24:26 +02:00
parent b764e5e053
commit bdb576ea88
3 changed files with 23 additions and 218 deletions
Download Patch File Download Diff File Expand all files Collapse all files
sc/source/core/opencl/formulagroupcl.cxx
+3 -63
View File
@@ -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;
sc/source/core/opencl/op_statistical.cxx
+18 -153
View File
@@ -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<const
formula::DoubleVectorRefToken *>(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; l<loop; l++){\n";
ss << " tmp = 0.0;\n";
ss << " count = 0;\n";
ss << " int loopOffset = l*512;\n";
ss << " int p1 = loopOffset + lidx + offset, p2 = p1 + 256;\n";
ss << " if (p2 < min(offset + windowSize, arrayLength)) {\n";
ss << " tmp0 = 0.0;\n";
std::string p1 = "p1";
std::string p2 = "p2";
ss << " tmp0 =";
rArg->GenDeclRef(ss);
ss << "["<<p1.c_str()<<"];\n";
ss << " if(!isnan(tmp0))\n";
ss << " {\n";
ss << " tmp += log(tmp0);\n";
ss << " count++;\n";
ss << " }\n";
ss << " tmp0 =";
rArg->GenDeclRef(ss);
ss << "["<<p2.c_str()<<"];\n";
ss << " if(!isnan(tmp0))\n";
ss << " {\n";
ss << " tmp += log(tmp0);\n";
ss << " count++;\n";
ss << " }\n";
ss << " }\n";
ss << " else if (p1 < min(arrayLength, offset + windowSize)) {\n";
ss << " tmp0 =";
rArg->GenDeclRef(ss);
ss << "["<<p1.c_str()<<"];\n";
ss << " if(!isnan(tmp0))\n";
ss << " {\n";
ss << " tmp += log(tmp0);\n";
ss << " count++;\n";
ss << " }\n";
ss << " }\n";
ss << " shm_buf[lidx] = tmp;\n";
ss << " count_buf[lidx] = count;\n";
ss << " barrier(CLK_LOCAL_MEM_FENCE);\n";
ss << " for (int i = 128; i >0; 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 << "}";
}
sc/source/core/opencl/op_statistical.hxx
+2 -2
View File
@@ -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