ODFstudio/odfstudio
2
0
Fork 0
Code Issues Packages Activity

fix and simplify openCL *A functions (COUNTA,MINA, etc.)

Browse Source 
These work just like their non-A counterparts, they just also accept
strings and treat them as zeros. So instead of having a duplicated
function for everything, use the basic functions, detect that string
arguments should be treated this way and then convert them in the input
data. This as a side-effect also makes
the ScCalcConfig::StringConversion::ZERO part work (no idea if that's
even been used).

Change-Id: I3385b5363c15b0ae45ff191df00ac357e80bd3c6
Reviewed-on: https://gerrit.libreoffice.org/c/core/+/140226
Tested-by: Jenkins
Reviewed-by: Luboš Luňák <l.lunak@collabora.com>
This commit is contained in:
Luboš Luňák
2022-09-20 17:54:03 +02:00
parent 68ae079807
commit 424595a599
9 changed files with 450 additions and 2370 deletions
Download Patch File Download Diff File Expand all files Collapse all files
sc/source/core/opencl/formulagroupcl.cxx
+104 -94
View File
@@ -108,6 +108,7 @@ OUString LimitedString( const OUString& str )
return OUString::Concat("\"") + str.subView( 0, 20 ) + "\"...";
}
const int MAX_PEEK_ELEMENTS = 5;
// Returns formatted contents of the data (possibly shortened), to be used in debug output.
std::string DebugPeekData(const FormulaToken* ref, int doubleRefIndex = 0)
{
@@ -117,16 +118,19 @@ std::string DebugPeekData(const FormulaToken* ref, int doubleRefIndex = 0)
static_cast<const formula::SingleVectorRefToken*>(ref);
outputstream buf;
buf << "SingleRef {";
for( size_t i = 0; i < std::min< size_t >( 4, pSVR->GetArrayLength()); ++i )
for( size_t i = 0; i < std::min< size_t >( MAX_PEEK_ELEMENTS, pSVR->GetArrayLength()); ++i )
{
if( i != 0 )
buf << ",";
if( pSVR->GetArray().mpNumericArray != nullptr )
buf << pSVR->GetArray().mpNumericArray[ i ];
else if( pSVR->GetArray().mpStringArray != nullptr )
if( pSVR->GetArray().mpStringArray != nullptr
&& pSVR->GetArray().mpStringArray[ i ] != nullptr )
{
buf << LimitedString( OUString( pSVR->GetArray().mpStringArray[ i ] ));
}
else if( pSVR->GetArray().mpNumericArray != nullptr )
buf << pSVR->GetArray().mpNumericArray[ i ];
}
if( pSVR->GetArrayLength() > 4 )
if( pSVR->GetArrayLength() > MAX_PEEK_ELEMENTS )
buf << ",...";
buf << "}";
return buf.str();
@@ -137,16 +141,19 @@ std::string DebugPeekData(const FormulaToken* ref, int doubleRefIndex = 0)
static_cast<const formula::DoubleVectorRefToken*>(ref);
outputstream buf;
buf << "DoubleRef {";
for( size_t i = 0; i < std::min< size_t >( 4, pDVR->GetArrayLength()); ++i )
for( size_t i = 0; i < std::min< size_t >( MAX_PEEK_ELEMENTS, pDVR->GetArrayLength()); ++i )
{
if( i != 0 )
buf << ",";
if( pDVR->GetArrays()[doubleRefIndex].mpNumericArray != nullptr )
buf << pDVR->GetArrays()[doubleRefIndex].mpNumericArray[ i ];
else if( pDVR->GetArrays()[doubleRefIndex].mpStringArray != nullptr )
if( pDVR->GetArrays()[doubleRefIndex].mpStringArray != nullptr
&& pDVR->GetArrays()[doubleRefIndex].mpStringArray[ i ] != nullptr )
{
buf << LimitedString( OUString( pDVR->GetArrays()[doubleRefIndex].mpStringArray[ i ] ));
}
else if( pDVR->GetArrays()[doubleRefIndex].mpNumericArray != nullptr )
buf << pDVR->GetArrays()[doubleRefIndex].mpNumericArray[ i ];
}
if( pDVR->GetArrayLength() > 4 )
if( pDVR->GetArrayLength() > MAX_PEEK_ELEMENTS )
buf << ",...";
buf << "}";
return buf.str();
@@ -172,13 +179,13 @@ std::string DebugPeekDoubles(const double* data, int size)
{
outputstream buf;
buf << "{";
for( int i = 0; i < std::min( 4, size ); ++i )
for( int i = 0; i < std::min( MAX_PEEK_ELEMENTS, size ); ++i )
{
if( i != 0 )
buf << ",";
buf << data[ i ];
}
if( size > 4 )
if( size > MAX_PEEK_ELEMENTS )
buf << ",...";
buf << "}";
return buf.str();
@@ -200,7 +207,21 @@ size_t VectorRef::Marshal( cl_kernel k, int argno, int, cl_program )
SAL_INFO("sc.opencl", "SingleVectorRef len=" << pSVR->GetArrayLength() << " mpNumericArray=" << pSVR->GetArray().mpNumericArray << " (mpStringArray=" << pSVR->GetArray().mpStringArray << ")");
pHostBuffer = const_cast<double*>(pSVR->GetArray().mpNumericArray);
if( forceStringsToZero && pSVR->GetArray().mpStringArray != nullptr )
{
dataBuffer.resize( pSVR->GetArrayLength());
for( size_t i = 0; i < pSVR->GetArrayLength(); ++i )
if( pSVR->GetArray().mpStringArray[ i ] != nullptr )
dataBuffer[ i ] = 0;
else
dataBuffer[ i ] = pSVR->GetArray().mpNumericArray[ i ];
pHostBuffer = dataBuffer.data();
SAL_INFO("sc.opencl", "Forced strings to zero : " << DebugPeekDoubles( pHostBuffer, pSVR->GetArrayLength()));
}
else
{
pHostBuffer = const_cast<double*>(pSVR->GetArray().mpNumericArray);
}
szHostBuffer = pSVR->GetArrayLength() * sizeof(double);
}
else if (ref->GetType() == formula::svDoubleVectorRef)
@@ -210,8 +231,21 @@ size_t VectorRef::Marshal( cl_kernel k, int argno, int, cl_program )
SAL_INFO("sc.opencl", "DoubleVectorRef index=" << mnIndex << " len=" << pDVR->GetArrayLength() << " mpNumericArray=" << pDVR->GetArrays()[mnIndex].mpNumericArray << " (mpStringArray=" << pDVR->GetArrays()[mnIndex].mpStringArray << ")");
pHostBuffer = const_cast<double*>(
pDVR->GetArrays()[mnIndex].mpNumericArray);
if( forceStringsToZero && pDVR->GetArrays()[mnIndex].mpStringArray != nullptr )
{
dataBuffer.resize( pDVR->GetArrayLength());
for( size_t i = 0; i < pDVR->GetArrayLength(); ++i )
if( pDVR->GetArrays()[mnIndex].mpStringArray[ i ] != nullptr )
dataBuffer[ i ] = 0;
else
dataBuffer[ i ] = pDVR->GetArrays()[mnIndex].mpNumericArray[ i ];
pHostBuffer = dataBuffer.data();
SAL_INFO("sc.opencl", "Forced strings to zero : " << DebugPeekDoubles( pHostBuffer, pDVR->GetArrayLength()));
}
else
{
pHostBuffer = const_cast<double*>(pDVR->GetArrays()[mnIndex].mpNumericArray);
}
szHostBuffer = pDVR->GetArrayLength() * sizeof(double);
}
else
@@ -333,57 +367,6 @@ public:
}
};
/// Arguments that are actually compile-time constants
class DynamicKernelConstantArgument : public DynamicKernelArgument
{
public:
DynamicKernelConstantArgument( const ScCalcConfig& config, const std::string& s,
const FormulaTreeNodeRef& ft ) :
DynamicKernelArgument(config, s, ft) { }
/// Generate declaration
virtual void GenDecl( outputstream& ss ) const override
{
ss << "double " << mSymName;
}
virtual void GenDeclRef( outputstream& ss ) const override
{
ss << mSymName;
}
virtual void GenSlidingWindowDecl( outputstream& ss ) const override
{
GenDecl(ss);
}
virtual std::string GenSlidingWindowDeclRef( bool = false ) const override
{
if (GetFormulaToken()->GetType() != formula::svDouble)
throw Unhandled(__FILE__, __LINE__);
return mSymName;
}
virtual size_t GetWindowSize() const override
{
return 1;
}
double GetDouble() const
{
FormulaToken* Tok = GetFormulaToken();
if (Tok->GetType() != formula::svDouble)
throw Unhandled(__FILE__, __LINE__);
return Tok->GetDouble();
}
/// Create buffer and pass the buffer to a given kernel
virtual size_t Marshal( cl_kernel k, int argno, int, cl_program ) override
{
OpenCLZone zone;
double tmp = GetDouble();
// Pass the scalar result back to the rest of the formula kernel
SAL_INFO("sc.opencl", "Kernel " << k << " arg " << argno << ": double: " << preciseFloat( tmp ));
cl_int err = clSetKernelArg(k, argno, sizeof(double), static_cast<void*>(&tmp));
if (CL_SUCCESS != err)
throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__);
return 1;
}
};
class DynamicKernelPiArgument : public DynamicKernelArgument
{
public:
@@ -1109,6 +1092,7 @@ size_t DynamicKernelSlidingArgument<Base>::GenReductionLoopHeader( outputstream&
}
template class DynamicKernelSlidingArgument<VectorRef>;
template class DynamicKernelSlidingArgument<VectorRefStringsToZero>;
template class DynamicKernelSlidingArgument<DynamicKernelStringArgument>;
namespace {
@@ -2070,26 +2054,38 @@ DynamicKernelSoPArguments::DynamicKernelSoPArguments(const ScCalcConfig& config,
}
else if (pDVR->GetArrays()[j].mpNumericArray &&
pCodeGen->takeNumeric() &&
(AllStringsAreNull(pDVR->GetArrays()[j].mpStringArray, pDVR->GetArrayLength()) || mCalcConfig.meStringConversion == ScCalcConfig::StringConversion::ZERO))
(AllStringsAreNull(pDVR->GetArrays()[j].mpStringArray, pDVR->GetArrayLength())
|| mCalcConfig.meStringConversion == ScCalcConfig::StringConversion::ZERO
|| pCodeGen->forceStringsToZero()))
{
// Function takes numbers, and either there
// are no strings, or there are strings but
// they are to be treated as zero
SAL_INFO("sc.opencl", "Numbers (no strings or strings treated as zero)");
mvSubArguments.push_back(
VectorRefFactory<VectorRef>(mCalcConfig,
ts, ft->Children[i], mpCodeGen, j));
if(!AllStringsAreNull(pDVR->GetArrays()[j].mpStringArray, pDVR->GetArrayLength()))
{
mvSubArguments.push_back(
VectorRefFactory<VectorRefStringsToZero>(mCalcConfig,
ts, ft->Children[i], mpCodeGen, j));
}
else
{
mvSubArguments.push_back(
VectorRefFactory<VectorRef>(mCalcConfig,
ts, ft->Children[i], mpCodeGen, j));
}
}
else if (pDVR->GetArrays()[j].mpNumericArray == nullptr &&
pCodeGen->takeNumeric() &&
pDVR->GetArrays()[j].mpStringArray &&
mCalcConfig.meStringConversion == ScCalcConfig::StringConversion::ZERO)
( mCalcConfig.meStringConversion == ScCalcConfig::StringConversion::ZERO
|| pCodeGen->forceStringsToZero()))
{
// Function takes numbers, and there are only
// strings, but they are to be treated as zero
SAL_INFO("sc.opencl", "Only strings even if want numbers but should be treated as zero");
mvSubArguments.push_back(
VectorRefFactory<VectorRef>(mCalcConfig,
VectorRefFactory<VectorRefStringsToZero>(mCalcConfig,
ts, ft->Children[i], mpCodeGen, j));
}
else if (pDVR->GetArrays()[j].mpStringArray &&
@@ -2145,26 +2141,34 @@ DynamicKernelSoPArguments::DynamicKernelSoPArguments(const ScCalcConfig& config,
}
else if (pSVR->GetArray().mpNumericArray &&
pCodeGen->takeNumeric() &&
(AllStringsAreNull(pSVR->GetArray().mpStringArray, pSVR->GetArrayLength()) || mCalcConfig.meStringConversion == ScCalcConfig::StringConversion::ZERO))
(AllStringsAreNull(pSVR->GetArray().mpStringArray, pSVR->GetArrayLength())
|| mCalcConfig.meStringConversion == ScCalcConfig::StringConversion::ZERO
|| pCodeGen->forceStringsToZero()))
{
// Function takes numbers, and either there
// are no strings, or there are strings but
// they are to be treated as zero
SAL_INFO("sc.opencl", "Numbers (no strings or strings treated as zero)");
mvSubArguments.push_back(
std::make_shared<VectorRef>(mCalcConfig, ts,
ft->Children[i]));
if( !AllStringsAreNull(pSVR->GetArray().mpStringArray, pSVR->GetArrayLength()))
mvSubArguments.push_back(
std::make_shared<VectorRefStringsToZero>(mCalcConfig, ts,
ft->Children[i]));
else
mvSubArguments.push_back(
std::make_shared<VectorRef>(mCalcConfig, ts,
ft->Children[i]));
}
else if (pSVR->GetArray().mpNumericArray == nullptr &&
pCodeGen->takeNumeric() &&
pSVR->GetArray().mpStringArray &&
mCalcConfig.meStringConversion == ScCalcConfig::StringConversion::ZERO)
(mCalcConfig.meStringConversion == ScCalcConfig::StringConversion::ZERO
|| pCodeGen->forceStringsToZero()))
{
// Function takes numbers, and there are only
// strings, but they are to be treated as zero
SAL_INFO("sc.opencl", "Only strings even if want numbers but should be treated as zero");
mvSubArguments.push_back(
std::make_shared<VectorRef>(mCalcConfig, ts,
std::make_shared<VectorRefStringsToZero>(mCalcConfig, ts,
ft->Children[i]));
}
else if (pSVR->GetArray().mpStringArray &&
@@ -2208,6 +2212,16 @@ DynamicKernelSoPArguments::DynamicKernelSoPArguments(const ScCalcConfig& config,
std::make_shared<ConstStringArgument>(mCalcConfig, ts,
ft->Children[i]));
}
else if (pChild->GetType() == formula::svString
&& !pCodeGen->takeString()
&& pCodeGen->takeNumeric()
&& pCodeGen->forceStringsToZero())
{
SAL_INFO("sc.opencl", "Constant string case, treated as zero");
mvSubArguments.push_back(
DynamicKernelArgumentRef(new DynamicKernelStringToZeroArgument(mCalcConfig, ts,
ft->Children[i])));
}
else
{
SAL_INFO("sc.opencl", "Unhandled operand, rejecting for OpenCL");
@@ -2230,15 +2244,27 @@ DynamicKernelSoPArguments::DynamicKernelSoPArguments(const ScCalcConfig& config,
case ocAverage:
mvSubArguments.push_back(SoPHelper(mCalcConfig, ts, ft->Children[i], std::make_shared<OpAverage>(nResultSize), nResultSize));
break;
case ocAverageA:
mvSubArguments.push_back(SoPHelper(mCalcConfig, ts, ft->Children[i], std::make_shared<OpAverageA>(nResultSize), nResultSize));
break;
case ocMin:
mvSubArguments.push_back(SoPHelper(mCalcConfig, ts, ft->Children[i], std::make_shared<OpMin>(nResultSize), nResultSize));
break;
case ocMinA:
mvSubArguments.push_back(SoPHelper(mCalcConfig, ts, ft->Children[i], std::make_shared<OpMinA>(nResultSize), nResultSize));
break;
case ocMax:
mvSubArguments.push_back(SoPHelper(mCalcConfig, ts, ft->Children[i], std::make_shared<OpMax>(nResultSize), nResultSize));
break;
case ocMaxA:
mvSubArguments.push_back(SoPHelper(mCalcConfig, ts, ft->Children[i], std::make_shared<OpMaxA>(nResultSize), nResultSize));
break;
case ocCount:
mvSubArguments.push_back(SoPHelper(mCalcConfig, ts, ft->Children[i], std::make_shared<OpCount>(nResultSize), nResultSize));
break;
case ocCount2:
mvSubArguments.push_back(SoPHelper(mCalcConfig, ts, ft->Children[i], std::make_shared<OpCountA>(nResultSize), nResultSize));
break;
case ocSumProduct:
mvSubArguments.push_back(SoPHelper(mCalcConfig, ts, ft->Children[i], std::make_shared<OpSumProduct>(), nResultSize));
break;
@@ -2849,22 +2875,6 @@ DynamicKernelSoPArguments::DynamicKernelSoPArguments(const ScCalcConfig& config,
mvSubArguments.push_back(SoPHelper(mCalcConfig, ts,
ft->Children[i], std::make_shared<OpFact>(), nResultSize));
break;
case ocMinA:
mvSubArguments.push_back(SoPHelper(mCalcConfig, ts,
ft->Children[i], std::make_shared<OpMinA>(), nResultSize));
break;
case ocCount2:
mvSubArguments.push_back(SoPHelper(mCalcConfig, ts,
ft->Children[i], std::make_shared<OpCountA>(), nResultSize));
break;
case ocMaxA:
mvSubArguments.push_back(SoPHelper(mCalcConfig, ts,
ft->Children[i], std::make_shared<OpMaxA>(), nResultSize));
break;
case ocAverageA:
mvSubArguments.push_back(SoPHelper(mCalcConfig, ts,
ft->Children[i], std::make_shared<OpAverageA>(), nResultSize));
break;
case ocVarA:
mvSubArguments.push_back(SoPHelper(mCalcConfig, ts,
ft->Children[i], std::make_shared<OpVarA>(), nResultSize));
sc/source/core/opencl/op_math.cxx
-114
View File
@@ -1650,102 +1650,6 @@ void SumOfProduct::GenSlidingWindowFunction( outputstream& ss,
ss << "}";
}
void Reduction::GenSlidingWindowFunction( outputstream& ss,
const std::string& sSymName, SubArguments& vSubArguments )
{
GenerateFunctionDeclaration( sSymName, vSubArguments, ss );
ss << "{\n";
ss << "double tmp = " << GetBottom() << ";\n";
ss << "int gid0 = get_global_id(0);\n";
if (isAverage() || isMinOrMax())
ss << "int nCount = 0;\n";
ss << "double tmpBottom;\n";
unsigned i = vSubArguments.size();
while (i--)
{
if (NumericRange* NR = dynamic_cast<NumericRange*>(vSubArguments[i].get()))
{
bool needBody;
NR->GenReductionLoopHeader(ss, needBody);
if (!needBody)
continue;
}
else if (ParallelNumericRange* PNR = dynamic_cast<ParallelNumericRange*>(vSubArguments[i].get()))
{
//did not handle yet
bool bNeedBody = false;
PNR->GenReductionLoopHeader(ss, mnResultSize, bNeedBody);
if (!bNeedBody)
continue;
}
else if (StringRange* SR = dynamic_cast<StringRange*>(vSubArguments[i].get()))
{
//did not handle yet
bool needBody;
SR->GenReductionLoopHeader(ss, needBody);
if (!needBody)
continue;
}
else
{
FormulaToken* pCur = vSubArguments[i]->GetFormulaToken();
assert(pCur);
assert(pCur->GetType() != formula::svDoubleVectorRef);
if (pCur->GetType() == formula::svSingleVectorRef ||
pCur->GetType() == formula::svDouble)
{
ss << "{\n";
}
}
if (ocPush == vSubArguments[i]->GetFormulaToken()->GetOpCode())
{
bool bNanHandled = HandleNaNArgument(ss, i, vSubArguments);
ss << "tmpBottom = " << GetBottom() << ";\n";
if (!bNanHandled)
{
ss << "if (isnan(";
ss << vSubArguments[i]->GenSlidingWindowDeclRef();
ss << "))\n";
if (ZeroReturnZero())
ss << " return 0;\n";
else
{
ss << " tmp = ";
ss << Gen2("tmpBottom", "tmp") << ";\n";
}
ss << "else\n";
}
ss << "{";
ss << " tmp = ";
ss << Gen2(vSubArguments[i]->GenSlidingWindowDeclRef(), "tmp");
ss << ";\n";
ss << " }\n";
ss << "}\n";
}
else
{
ss << "tmp = ";
ss << Gen2(vSubArguments[i]->GenSlidingWindowDeclRef(), "tmp");
ss << ";\n";
}
}
if (isAverage())
ss <<
"if (nCount==0)\n"
" return CreateDoubleError(DivisionByZero);\n";
else if (isMinOrMax())
ss <<
"if (nCount==0)\n"
" return 0;\n";
ss << "return tmp";
if (isAverage())
ss << "/(double)nCount";
ss << ";\n}";
}
void OpSum::BinInlineFun(std::set<std::string>& decls,std::set<std::string>& funs)
{
decls.insert(is_representable_integerDecl);
@@ -1756,12 +1660,6 @@ void OpSum::BinInlineFun(std::set<std::string>& decls,std::set<std::string>& fun
funs.insert(fsum_approx);
}
void OpAverage::BinInlineFun(std::set<std::string>& decls,std::set<std::string>& funs)
{
decls.insert(fsum_countDecl);
funs.insert(fsum_count);
}
void OpSub::BinInlineFun(std::set<std::string>& decls,std::set<std::string>& funs)
{
decls.insert(is_representable_integerDecl);
@@ -1772,18 +1670,6 @@ void OpSub::BinInlineFun(std::set<std::string>& decls,std::set<std::string>& fun
funs.insert(fsub_approx);
}
void OpMin::BinInlineFun(std::set<std::string>& decls,std::set<std::string>& funs)
{
decls.insert(fmin_countDecl);
funs.insert(fmin_count);
}
void OpMax::BinInlineFun(std::set<std::string>& decls,std::set<std::string>& funs)
{
decls.insert(fmax_countDecl);
funs.insert(fmax_count);
}
}
/* vim:set shiftwidth=4 softtabstop=4 expandtab: */
sc/source/core/opencl/op_math.hxx
-90
View File
@@ -570,29 +570,6 @@ public:
virtual std::string BinFuncName() const override { return "fsop"; }
};
class Reduction : public SlidingFunctionBase
{
int const mnResultSize;
public:
explicit Reduction(int nResultSize) : mnResultSize(nResultSize) {}
typedef DynamicKernelSlidingArgument<VectorRef> NumericRange;
typedef DynamicKernelSlidingArgument<DynamicKernelStringArgument> StringRange;
typedef ParallelReductionVectorRef<VectorRef> ParallelNumericRange;
virtual bool HandleNaNArgument( outputstream&, unsigned, SubArguments& ) const
{
return false;
}
virtual void GenSlidingWindowFunction( outputstream& ss,
const std::string& sSymName, SubArguments& vSubArguments ) override;
virtual bool isAverage() const { return false; }
virtual bool isMinOrMax() const { return false; }
virtual bool takeString() const override { return false; }
virtual bool takeNumeric() const override { return true; }
};
/// operator traits
class OpNop : public Reduction
{
@@ -607,22 +584,6 @@ public:
virtual std::string BinFuncName() const override { return "nop"; }
};
class OpCount : public Reduction
{
public:
explicit OpCount(int nResultSize) : Reduction(nResultSize) {}
virtual std::string GetBottom() override { return "0"; }
virtual std::string Gen2( const std::string& lhs, const std::string& rhs ) const override
{
outputstream ss;
ss << "(isnan(" << lhs << ")?" << rhs << ":" << rhs << "+1.0)";
return ss.str();
}
virtual std::string BinFuncName() const override { return "fcount"; }
virtual bool canHandleMultiVector() const override { return true; }
};
class OpSum : public Reduction
{
public:
@@ -641,24 +602,6 @@ public:
virtual bool canHandleMultiVector() const override { return true; }
};
class OpAverage : public Reduction
{
public:
explicit OpAverage(int nResultSize) : Reduction(nResultSize) {}
virtual std::string GetBottom() override { return "0"; }
virtual std::string Gen2( const std::string& lhs, const std::string& rhs ) const override
{
outputstream ss;
ss << "fsum_count(" << lhs << "," << rhs << ", &nCount)";
return ss.str();
}
virtual void BinInlineFun(std::set<std::string>& decls,std::set<std::string>& funs) override;
virtual std::string BinFuncName() const override { return "average"; }
virtual bool isAverage() const override { return true; }
virtual bool canHandleMultiVector() const override { return true; }
};
class OpSub : public Reduction
{
public:
@@ -723,39 +666,6 @@ public:
};
class OpMin : public Reduction
{
public:
explicit OpMin(int nResultSize) : Reduction(nResultSize) {}
virtual std::string GetBottom() override { return "NAN"; }
virtual std::string Gen2( const std::string& lhs, const std::string& rhs ) const override
{
return "fmin_count(" + lhs + "," + rhs + ", &nCount)";
}
virtual void BinInlineFun(std::set<std::string>& decls,std::set<std::string>& funs) override;
virtual std::string BinFuncName() const override { return "min"; }
virtual bool isMinOrMax() const override { return true; }
virtual bool canHandleMultiVector() const override { return true; }
};
class OpMax : public Reduction
{
public:
explicit OpMax(int nResultSize) : Reduction(nResultSize) {}
virtual std::string GetBottom() override { return "NAN"; }
virtual std::string Gen2( const std::string& lhs, const std::string& rhs ) const override
{
return "fmax_count(" + lhs + "," + rhs + ", &nCount)";
}
virtual void BinInlineFun(std::set<std::string>& decls,std::set<std::string>& funs) override;
virtual std::string BinFuncName() const override { return "max"; }
virtual bool isMinOrMax() const override { return true; }
virtual bool canHandleMultiVector() const override { return true; }
};
}
/* vim:set shiftwidth=4 softtabstop=4 expandtab: */
sc/source/core/opencl/op_math_helpers.hxx
-26
View File
@@ -86,32 +86,6 @@ const char atan2Content[] =
" return a;\n"
"}\n";
const char fsum_countDecl[] = "double fsum_count(double a, double b, __private int *p);\n";
const char fsum_count[] =
"double fsum_count(double a, double b, __private int *p) {\n"
" bool t = isnan(a);\n"
" (*p) += t?0:1;\n"
" return t?b:a+b;\n"
"}\n";
const char fmin_countDecl[] = "double fmin_count(double a, double b, __private int *p);\n";
const char fmin_count[] =
"double fmin_count(double a, double b, __private int *p) {\n"
" double result = fmin(a, b);\n"
" bool t = isnan(result);\n"
" (*p) += t?0:1;\n"
" return result;\n"
"}\n";
const char fmax_countDecl[] = "double fmax_count(double a, double b, __private int *p);\n";
const char fmax_count[] =
"double fmax_count(double a, double b, __private int *p) {\n"
" double result = fmax(a, b);\n"
" bool t = isnan(result);\n"
" (*p) += t?0:1;\n"
" return result;\n"
"}\n";
const char is_representable_integerDecl[] = "int is_representable_integer(double a);\n";
const char is_representable_integer[] =
"int is_representable_integer(double a) {\n"
sc/source/core/opencl/op_statistical.cxx
+9 -1997
View File
File diff suppressed because it is too large Load Diff
sc/source/core/opencl/op_statistical.hxx
+86 -46
View File
@@ -503,83 +503,123 @@ public:
virtual std::string BinFuncName() const override { return "OpRsq"; }
};
class OpMinA: public Normal
class OpMin : public Reduction
{
public:
virtual void GenSlidingWindowFunction(outputstream &ss,
const std::string &sSymName, SubArguments &vSubArguments) override;
virtual std::string BinFuncName() const override { return "OpMinA"; }
virtual bool takeString() const override { return true; }
virtual bool takeNumeric() const override { return true; }
explicit OpMin(int nResultSize) : Reduction(nResultSize) {}
virtual std::string GetBottom() override { return "NAN"; }
virtual std::string Gen2( const std::string& lhs, const std::string& rhs ) const override
{
return "fmin_count(" + lhs + "," + rhs + ", &nCount)";
}
virtual void BinInlineFun(std::set<std::string>& decls,std::set<std::string>& funs) override;
virtual std::string BinFuncName() const override { return "min"; }
virtual bool isMinOrMax() const override { return true; }
virtual bool canHandleMultiVector() const override { return true; }
};
class OpCountA: public Normal
class OpMax : public Reduction
{
public:
virtual void GenSlidingWindowFunction(outputstream &ss,
const std::string &sSymName, SubArguments &vSubArguments) override;
explicit OpMax(int nResultSize) : Reduction(nResultSize) {}
virtual std::string GetBottom() override { return "NAN"; }
virtual std::string Gen2( const std::string& lhs, const std::string& rhs ) const override
{
return "fmax_count(" + lhs + "," + rhs + ", &nCount)";
}
virtual void BinInlineFun(std::set<std::string>& decls,std::set<std::string>& funs) override;
virtual std::string BinFuncName() const override { return "max"; }
virtual bool isMinOrMax() const override { return true; }
virtual bool canHandleMultiVector() const override { return true; }
};
class OpAverage : public Reduction
{
public:
explicit OpAverage(int nResultSize) : Reduction(nResultSize) {}
virtual std::string GetBottom() override { return "0"; }
virtual std::string Gen2( const std::string& lhs, const std::string& rhs ) const override
{
outputstream ss;
ss << "fsum_count(" << lhs << "," << rhs << ", &nCount)";
return ss.str();
}
virtual void BinInlineFun(std::set<std::string>& decls,std::set<std::string>& funs) override;
virtual std::string BinFuncName() const override { return "average"; }
virtual bool isAverage() const override { return true; }
virtual bool canHandleMultiVector() const override { return true; }
};
class OpCount : public Reduction
{
public:
explicit OpCount(int nResultSize) : Reduction(nResultSize) {}
virtual std::string GetBottom() override { return "0"; }
virtual std::string Gen2( const std::string& lhs, const std::string& rhs ) const override
{
outputstream ss;
ss << "(isnan(" << lhs << ")?" << rhs << ":" << rhs << "+1.0)";
return ss.str();
}
virtual std::string BinFuncName() const override { return "fcount"; }
virtual bool canHandleMultiVector() const override { return true; }
};
class OpCountA: public OpCount
{
public:
explicit OpCountA(int nResultSize) : OpCount(nResultSize) {}
virtual std::string BinFuncName() const override { return "OpCountA"; }
virtual bool takeString() const override { return true; }
virtual bool takeNumeric() const override { return true; }
virtual bool canHandleMultiVector() const override { return true; }
virtual bool forceStringsToZero() const override { return true; }
};
class OpMaxA: public Normal
class OpMaxA: public OpMax
{
public:
virtual void GenSlidingWindowFunction(outputstream &ss,
const std::string &sSymName, SubArguments &vSubArguments) override;
explicit OpMaxA(int nResultSize) : OpMax(nResultSize) {}
virtual std::string BinFuncName() const override { return "OpMaxA"; }
virtual bool takeString() const override { return true; }
virtual bool takeNumeric() const override { return true; }
virtual bool canHandleMultiVector() const override { return true; }
virtual bool forceStringsToZero() const override { return true; }
};
class OpVarA: public Normal
class OpMinA : public OpMin
{
public:
explicit OpMinA(int nResultSize) : OpMin(nResultSize) {}
virtual std::string BinFuncName() const override { return "OpMinA"; }
virtual bool forceStringsToZero() const override { return true; }
};
class OpVarA: public OpVar
{
public:
virtual void GenSlidingWindowFunction(outputstream &ss,
const std::string &sSymName, SubArguments &vSubArguments) override;
virtual std::string BinFuncName() const override { return "OpVarA"; }
virtual bool takeString() const override { return true; }
virtual bool takeNumeric() const override { return true; }
virtual bool canHandleMultiVector() const override { return true; }
virtual bool forceStringsToZero() const override { return true; }
};
class OpVarPA: public Normal
class OpVarPA: public OpVarP
{
public:
virtual void GenSlidingWindowFunction(outputstream &ss,
const std::string &sSymName, SubArguments &vSubArguments) override;
virtual std::string BinFuncName() const override { return "OpVarPA"; }
virtual bool takeString() const override { return true; }
virtual bool takeNumeric() const override { return true; }
virtual bool canHandleMultiVector() const override { return true; }
virtual bool forceStringsToZero() const override { return true; }
};
class OpStDevPA: public Normal
class OpStDevPA: public OpStDevP
{
public:
virtual void GenSlidingWindowFunction(outputstream &ss,
const std::string &sSymName, SubArguments &vSubArguments) override;
virtual std::string BinFuncName() const override { return "OpStDevPA"; }
virtual bool takeString() const override { return true; }
virtual bool takeNumeric() const override { return true; }
virtual bool forceStringsToZero() const override { return true; }
};
class OpAverageA: public Normal
class OpAverageA: public OpAverage
{
public:
virtual void GenSlidingWindowFunction(outputstream &ss,
const std::string &sSymName, SubArguments &vSubArguments) override;
explicit OpAverageA(int nResultSize) : OpAverage(nResultSize) {}
virtual std::string BinFuncName() const override { return "OpAverageA"; }
virtual bool takeString() const override { return true; }
virtual bool takeNumeric() const override { return true; }
virtual bool canHandleMultiVector() const override { return true; }
virtual bool forceStringsToZero() const override { return true; }
};
class OpStDevA: public Normal
class OpStDevA: public OpStDev
{
public:
virtual void GenSlidingWindowFunction(outputstream &ss,
const std::string &sSymName, SubArguments &vSubArguments) override;
virtual std::string BinFuncName() const override { return "OpStDevA"; }
virtual bool takeString() const override { return true; }
virtual bool takeNumeric() const override { return true; }
virtual bool forceStringsToZero() const override { return true; }
};
}
sc/source/core/opencl/op_statistical_helpers.hxx
+23
View File
@@ -23,6 +23,29 @@ const char GetValue[] =
"{\n"
" return fp - 2 * GetTDist(x, fDF);\n"
"}\n";
const char fmin_countDecl[] = "double fmin_count(double a, double b, __private int *p);\n";
const char fmin_count[] =
"double fmin_count(double a, double b, __private int *p) {\n"
" double result = fmin(a, b);\n"
" bool t = isnan(result);\n"
" (*p) += t?0:1;\n"
" return result;\n"
"}\n";
const char fmax_countDecl[] = "double fmax_count(double a, double b, __private int *p);\n";
const char fmax_count[] =
"double fmax_count(double a, double b, __private int *p) {\n"
" double result = fmax(a, b);\n"
" bool t = isnan(result);\n"
" (*p) += t?0:1;\n"
" return result;\n"
"}\n";
const char fsum_countDecl[] = "double fsum_count(double a, double b, __private int *p);\n";
const char fsum_count[] =
"double fsum_count(double a, double b, __private int *p) {\n"
" bool t = isnan(a);\n"
" (*p) += t?0:1;\n"
" return t?b:a+b;\n"
"}\n";
const char GetGammaSeriesDecl[] =
"double GetGammaSeries( double fA, double fX );\n";
const char GetGammaSeries[] =
sc/source/core/opencl/opbase.cxx
+122 -2
View File
@@ -84,7 +84,7 @@ bool DynamicKernelArgument::NeedParallelReduction() const
}
VectorRef::VectorRef( const ScCalcConfig& config, const std::string& s, const FormulaTreeNodeRef& ft, int idx ) :
DynamicKernelArgument(config, s, ft), mpClmem(nullptr), mnIndex(idx)
DynamicKernelArgument(config, s, ft), mpClmem(nullptr), mnIndex(idx), forceStringsToZero( false )
{
if (mnIndex)
{
@@ -174,6 +174,13 @@ bool VectorRef::NeedParallelReduction() const
return false;
}
VectorRefStringsToZero::VectorRefStringsToZero( const ScCalcConfig& config, const std::string& s,
const FormulaTreeNodeRef& ft, int index )
: VectorRef( config, s, ft, index )
{
forceStringsToZero = true;
}
void SlidingFunctionBase::GenerateArg( const char* name, int arg, SubArguments& vSubArguments, outputstream& ss )
{
assert( arg < int( vSubArguments.size()));
@@ -196,6 +203,11 @@ void SlidingFunctionBase::GenerateArg( const char* name, int arg, SubArguments&
}
else if(token->GetType() == formula::svDouble)
ss << " " << name << " = " << token->GetDouble() << ";\n";
else if(token->GetType() == formula::svString)
{
assert( dynamic_cast<DynamicKernelStringToZeroArgument*>(vSubArguments[arg].get()));
ss << " " << name << " = 0.0;\n";
}
else
throw Unhandled( __FILE__, __LINE__ );
}
@@ -259,13 +271,23 @@ void SlidingFunctionBase::GenerateRangeArgs( int firstArg, int lastArg, SubArgum
ss << code;
ss << " }\n";
}
else
else if(token->GetType() == formula::svDouble)
{
ss << " {\n";
ss << " double arg = " << token->GetDouble() << ";\n";
ss << code;
ss << " }\n";
}
else if(token->GetType() == formula::svString)
{
assert( dynamic_cast<DynamicKernelStringToZeroArgument*>(vSubArguments[i].get()));
ss << " {\n";
ss << " double arg = 0.0;\n";
ss << code;
ss << " }\n";
}
else
throw Unhandled( __FILE__, __LINE__ );
}
else
{
@@ -611,6 +633,104 @@ void CheckVariables::UnrollDoubleVector( outputstream& ss,
ss << " }\n";
}
void Reduction::GenSlidingWindowFunction( outputstream& ss,
const std::string& sSymName, SubArguments& vSubArguments )
{
GenerateFunctionDeclaration( sSymName, vSubArguments, ss );
ss << "{\n";
ss << "double tmp = " << GetBottom() << ";\n";
ss << "int gid0 = get_global_id(0);\n";
if (isAverage() || isMinOrMax())
ss << "int nCount = 0;\n";
ss << "double tmpBottom;\n";
unsigned i = vSubArguments.size();
while (i--)
{
if (NumericRange* NR = dynamic_cast<NumericRange*>(vSubArguments[i].get()))
{
bool needBody;
NR->GenReductionLoopHeader(ss, needBody);
if (!needBody)
continue;
}
else if (NumericRangeStringsToZero* NRS = dynamic_cast<NumericRangeStringsToZero*>(vSubArguments[i].get()))
{
bool needBody;
NRS->GenReductionLoopHeader(ss, needBody);
if (!needBody)
continue;
}
else if (ParallelNumericRange* PNR = dynamic_cast<ParallelNumericRange*>(vSubArguments[i].get()))
{
//did not handle yet
bool bNeedBody = false;
PNR->GenReductionLoopHeader(ss, mnResultSize, bNeedBody);
if (!bNeedBody)
continue;
}
else if (StringRange* SR = dynamic_cast<StringRange*>(vSubArguments[i].get()))
{
//did not handle yet
bool needBody;
SR->GenReductionLoopHeader(ss, needBody);
if (!needBody)
continue;
}
else
{
FormulaToken* pCur = vSubArguments[i]->GetFormulaToken();
if( pCur == nullptr || pCur->GetType() == formula::svDoubleVectorRef )
{
throw Unhandled(__FILE__, __LINE__);
}
ss << "{\n";
}
if (ocPush == vSubArguments[i]->GetFormulaToken()->GetOpCode())
{
bool bNanHandled = HandleNaNArgument(ss, i, vSubArguments);
ss << " tmpBottom = " << GetBottom() << ";\n";
if (!bNanHandled)
{
ss << " if (isnan(";
ss << vSubArguments[i]->GenSlidingWindowDeclRef();
ss << "))\n";
if (ZeroReturnZero())
ss << " return 0;\n";
else
{
ss << " tmp = ";
ss << Gen2("tmpBottom", "tmp") << ";\n";
}
ss << " else\n";
}
ss << " tmp = ";
ss << Gen2(vSubArguments[i]->GenSlidingWindowDeclRef(), "tmp");
ss << ";\n";
}
else
{
ss << " tmp = ";
ss << Gen2(vSubArguments[i]->GenSlidingWindowDeclRef(), "tmp");
ss << ";\n";
}
ss << "}\n";
}
if (isAverage())
ss <<
"if (nCount==0)\n"
" return CreateDoubleError(DivisionByZero);\n";
else if (isMinOrMax())
ss <<
"if (nCount==0)\n"
" return 0;\n";
ss << "return tmp";
if (isAverage())
ss << "/(double)nCount";
ss << ";\n}";
}
}
/* vim:set shiftwidth=4 softtabstop=4 expandtab: */
sc/source/core/opencl/opbase.hxx
+106 -1
View File
@@ -13,6 +13,8 @@
#include <formula/token.hxx>
#include <formula/types.hxx>
#include <formula/vectortoken.hxx>
#include <opencl/OpenCLZone.hxx>
#include <sal/log.hxx>
#include <memory>
#include <set>
#include <vector>
@@ -94,7 +96,7 @@ public:
} while( false )
#define CHECK_PARAMETER_DOUBLEVECTORREF(arg) \
do { \
FormulaToken *token = vSubArguments[arg]->GetFormulaToken(); \
formula::FormulaToken *token = vSubArguments[arg]->GetFormulaToken(); \
if (token == nullptr || token->GetType() != formula::svDoubleVectorRef) \
throw Unhandled(__FILE__, __LINE__); \
} while( false )
@@ -206,6 +208,17 @@ protected:
cl_mem mpClmem;
// index in multiple double vector refs that have multiple ranges
const int mnIndex;
// Makes Marshall convert strings to 0 values.
bool forceStringsToZero;
// Used for storing when the data needs to be modified before sending to OpenCL.
std::vector< double > dataBuffer;
};
// Sets VectorRef::forceStringsToZero.
class VectorRefStringsToZero : public VectorRef
{
public:
VectorRefStringsToZero( const ScCalcConfig& config, const std::string& s, const FormulaTreeNodeRef& ft, int index = 0 );
};
/// A vector of strings
@@ -229,6 +242,72 @@ public:
virtual size_t Marshal( cl_kernel, int, int, cl_program ) override;
};
/// Arguments that are actually compile-time constants
class DynamicKernelConstantArgument : public DynamicKernelArgument
{
public:
DynamicKernelConstantArgument( const ScCalcConfig& config, const std::string& s,
const FormulaTreeNodeRef& ft ) :
DynamicKernelArgument(config, s, ft) { }
/// Generate declaration
virtual void GenDecl( outputstream& ss ) const override
{
ss << "double " << mSymName;
}
virtual void GenDeclRef( outputstream& ss ) const override
{
ss << mSymName;
}
virtual void GenSlidingWindowDecl( outputstream& ss ) const override
{
GenDecl(ss);
}
virtual std::string GenSlidingWindowDeclRef( bool = false ) const override
{
if (GetFormulaToken()->GetType() != formula::svDouble)
throw Unhandled(__FILE__, __LINE__);
return mSymName;
}
virtual size_t GetWindowSize() const override
{
return 1;
}
virtual double GetDouble() const
{
formula::FormulaToken* Tok = GetFormulaToken();
if (Tok->GetType() != formula::svDouble)
throw Unhandled(__FILE__, __LINE__);
return Tok->GetDouble();
}
/// Create buffer and pass the buffer to a given kernel
virtual size_t Marshal( cl_kernel k, int argno, int, cl_program ) override
{
OpenCLZone zone;
double tmp = GetDouble();
// Pass the scalar result back to the rest of the formula kernel
SAL_INFO("sc.opencl", "Kernel " << k << " arg " << argno << ": double: " << preciseFloat( tmp ));
cl_int err = clSetKernelArg(k, argno, sizeof(double), static_cast<void*>(&tmp));
if (CL_SUCCESS != err)
throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__);
return 1;
}
};
// Constant 0 argument when a string is forced to zero.
class DynamicKernelStringToZeroArgument : public DynamicKernelConstantArgument
{
public:
DynamicKernelStringToZeroArgument( const ScCalcConfig& config, const std::string& s,
const FormulaTreeNodeRef& ft ) :
DynamicKernelConstantArgument(config, s, ft) { }
virtual std::string GenSlidingWindowDeclRef( bool = false ) const override
{
return mSymName;
}
virtual double GetDouble() const override { return 0; }
};
/// Abstract class for code generation
class OpBase
{
@@ -250,6 +329,8 @@ public:
//Continue process 'Zero' or Not(like OpMul, not continue process when meet
// 'Zero'
virtual bool ZeroReturnZero() { return false;}
// For use with COUNTA() etc, input strings will be converted to 0 in data.
virtual bool forceStringsToZero() const { return false; }
virtual ~OpBase() { }
};
@@ -393,6 +474,30 @@ protected:
cl_mem mpClmem2;
};
class Reduction : public SlidingFunctionBase
{
int const mnResultSize;
public:
explicit Reduction(int nResultSize) : mnResultSize(nResultSize) {}
typedef DynamicKernelSlidingArgument<VectorRef> NumericRange;
typedef DynamicKernelSlidingArgument<VectorRefStringsToZero> NumericRangeStringsToZero;
typedef DynamicKernelSlidingArgument<DynamicKernelStringArgument> StringRange;
typedef ParallelReductionVectorRef<VectorRef> ParallelNumericRange;
virtual bool HandleNaNArgument( outputstream&, unsigned, SubArguments& ) const
{
return false;
}
virtual void GenSlidingWindowFunction( outputstream& ss,
const std::string& sSymName, SubArguments& vSubArguments ) override;
virtual bool isAverage() const { return false; }
virtual bool isMinOrMax() const { return false; }
virtual bool takeString() const override { return false; }
virtual bool takeNumeric() const override { return true; }
};
}
/* vim:set shiftwidth=4 softtabstop=4 expandtab: */