fix handling of string arguments in opencl

As one of the code comments said the code used string hashes to represent
strings and this was a broken idea. But the basic idea of that is actually
valid, so just implement that properly and use it only for comparing
strings. See the code comment in opbase.cxx for technical details.

Change-Id: I113d6b4d5e1e78bbe2c05aafc0572605e2595ad8
Reviewed-on: https://gerrit.libreoffice.org/c/core/+/140395
Tested-by: Jenkins
Reviewed-by: Luboš Luňák <l.lunak@collabora.com>

6 files changed, 256 insertions, 119 deletions

diff --git a/sc/source/core/opencl/formulagroupcl.cxx b/sc/source/core/opencl/formulagroupcl.cxx
index 4f274749484b..967f52c4ce43 100644
--- a/sc/source/core/opencl/formulagroupcl.cxx
+++ b/sc/source/core/opencl/formulagroupcl.cxx
@@ -299,73 +299,8 @@ size_t VectorRef::Marshal( cl_kernel k, int argno, int, cl_program )
     return 1;
 }
 
-/// Arguments that are actually compile-time constant string
-/// Currently, only the hash is passed.
-/// TBD(IJSUNG): pass also length and the actual string if there is a
-/// hash function collision
-
-/// FIXME: This idea of passing of hashes of uppercased strings into OpenCL code is fairly potent
-/// crack. It is hopefully not used at all any more, but noticing that there are string arguments
-/// automatically disables use of OpenCL for a formula group. If at some point there are resources
-/// to drain the OpenCL swamp, this should go away.
-
 namespace {
 
-class ConstStringArgument : public DynamicKernelArgument
-{
-public:
-    ConstStringArgument( const ScCalcConfig& config, const std::string& s,
-        const FormulaTreeNodeRef& ft ) :
-        DynamicKernelArgument(config, s, ft) { }
-    /// Generate declaration
-    virtual void GenDecl( outputstream& ss ) const override
-    {
-        ss << "unsigned " << mSymName;
-    }
-    virtual void GenDeclRef( outputstream& ss ) const override
-    {
-        ss << GenSlidingWindowDeclRef();
-    }
-    virtual void GenSlidingWindowDecl( outputstream& ss ) const override
-    {
-        GenDecl(ss);
-    }
-    virtual std::string GenSlidingWindowDeclRef( bool = false ) const override
-    {
-        outputstream ss;
-        if (GetFormulaToken()->GetType() != formula::svString)
-            throw Unhandled(__FILE__, __LINE__);
-        FormulaToken* Tok = GetFormulaToken();
-        ss << Tok->GetString().getString().toAsciiUpperCase().hashCode() << "U";
-        return ss.str();
-    }
-    virtual size_t GetWindowSize() const override
-    {
-        return 1;
-    }
-    /// Pass the 32-bit hash of the string to the kernel
-    virtual size_t Marshal( cl_kernel k, int argno, int, cl_program ) override
-    {
-        OpenCLZone zone;
-        FormulaToken* ref = mFormulaTree->GetFormulaToken();
-        cl_uint hashCode = 0;
-        if (ref->GetType() != formula::svString)
-        {
-            throw Unhandled(__FILE__, __LINE__);
-        }
-
-        const OUString s = ref->GetString().getString().toAsciiUpperCase();
-        hashCode = s.hashCode();
-
-        // Pass the scalar result back to the rest of the formula kernel
-        SAL_INFO("sc.opencl", "Kernel " << k << " arg " << argno << ": cl_uint: " << hashCode << "(" << DebugPeekData(ref) << ")" );
-        cl_int err = clSetKernelArg(k, argno, sizeof(cl_uint), static_cast<void*>(&hashCode));
-        if (CL_SUCCESS != err)
-            throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__);
-        return 1;
-    }
-};
-
 class DynamicKernelPiArgument : public DynamicKernelArgument
 {
 public:
@@ -773,9 +708,65 @@ threefry2x32 (threefry2x32_ctr_t in, threefry2x32_key_t k)\n\
     }
 };
 
-}
+// Arguments that are actually compile-time constant string
+class ConstStringArgument : public DynamicKernelArgument
+{
+public:
+    ConstStringArgument( 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 << GenSlidingWindowDeclRef();
+    }
+    virtual void GenSlidingWindowDecl( outputstream& ss ) const override
+    {
+        GenDecl(ss);
+    }
+    virtual std::string GenSlidingWindowDeclRef( bool = false ) const override
+    {
+        outputstream ss;
+        if (GetFormulaToken()->GetType() != formula::svString)
+            throw Unhandled(__FILE__, __LINE__);
+        FormulaToken* Tok = GetFormulaToken();
+        ss << GetStringId(Tok->GetString().getData());
+        return ss.str();
+    }
+    virtual std::string GenIsString( bool = false ) const override
+    {
+        return "true";
+    }
+    virtual size_t GetWindowSize() const override
+    {
+        return 1;
+    }
+    virtual size_t Marshal( cl_kernel k, int argno, int, cl_program ) override
+    {
+        FormulaToken* ref = mFormulaTree->GetFormulaToken();
+        if (ref->GetType() != formula::svString)
+        {
+            throw Unhandled(__FILE__, __LINE__);
+        }
+        cl_double stringId = GetStringId(ref->GetString().getData());
+
+        // Pass the scalar result back to the rest of the formula kernel
+        SAL_INFO("sc.opencl", "Kernel " << k << " arg " << argno
+            << ": stringId: " << stringId << " (" << DebugPeekData(ref) << ")" );
+        cl_int err = clSetKernelArg(k, argno, sizeof(cl_double), static_cast<void*>(&stringId));
+        if (CL_SUCCESS != err)
+            throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__);
+        return 1;
+    }
+};
+
+} // namespace
 
-/// Marshal a string vector reference
+// Marshal a string vector reference
 size_t DynamicKernelStringArgument::Marshal( cl_kernel k, int argno, int, cl_program )
 {
     OpenCLZone zone;
@@ -800,12 +791,12 @@ size_t DynamicKernelStringArgument::Marshal( cl_kernel k, int argno, int, cl_pro
         nStrings = pDVR->GetArrayLength();
         vRef = pDVR->GetArrays()[mnIndex];
     }
-    size_t szHostBuffer = nStrings * sizeof(cl_int);
-    cl_uint* pHashBuffer = nullptr;
+    size_t szHostBuffer = nStrings * sizeof(cl_double);
+    cl_double* pStringIdsBuffer = nullptr;
 
     if (vRef.mpStringArray != nullptr)
     {
-        // Marshal strings. Right now we pass hashes of these string
+        // Marshal strings. See GetStringId().
         mpClmem = clCreateBuffer(kEnv.mpkContext,
             cl_mem_flags(CL_MEM_READ_ONLY) | CL_MEM_ALLOC_HOST_PTR,
             szHostBuffer, nullptr, &err);
@@ -813,7 +804,7 @@ size_t DynamicKernelStringArgument::Marshal( cl_kernel k, int argno, int, cl_pro
             throw OpenCLError("clCreateBuffer", err, __FILE__, __LINE__);
         SAL_INFO("sc.opencl", "Created buffer " << mpClmem << " size " << szHostBuffer);
 
-        pHashBuffer = static_cast<cl_uint*>(clEnqueueMapBuffer(
+        pStringIdsBuffer = static_cast<cl_double*>(clEnqueueMapBuffer(
             kEnv.mpkCmdQueue, mpClmem, CL_TRUE, CL_MAP_WRITE, 0,
             szHostBuffer, 0, nullptr, nullptr, &err));
         if (CL_SUCCESS != err)
@@ -822,20 +813,15 @@ size_t DynamicKernelStringArgument::Marshal( cl_kernel k, int argno, int, cl_pro
         for (size_t i = 0; i < nStrings; i++)
         {
             if (vRef.mpStringArray[i])
-            {
-                const OUString tmp(vRef.mpStringArray[i]);
-                pHashBuffer[i] = tmp.hashCode();
-            }
+                pStringIdsBuffer[i] = GetStringId(vRef.mpStringArray[i]);
             else
-            {
-                pHashBuffer[i] = 0;
-            }
+                rtl::math::setNan(&pStringIdsBuffer[i]);
         }
     }
     else
     {
         if (nStrings == 0)
-            szHostBuffer = sizeof(cl_int); // a dummy small value
+            szHostBuffer = sizeof(cl_double); // a dummy small value
                                            // Marshal as a buffer of NANs
         mpClmem = clCreateBuffer(kEnv.mpkContext,
             cl_mem_flags(CL_MEM_READ_ONLY) | CL_MEM_ALLOC_HOST_PTR,
@@ -844,27 +830,54 @@ size_t DynamicKernelStringArgument::Marshal( cl_kernel k, int argno, int, cl_pro
             throw OpenCLError("clCreateBuffer", err, __FILE__, __LINE__);
         SAL_INFO("sc.opencl", "Created buffer " << mpClmem << " size " << szHostBuffer);
 
-        pHashBuffer = static_cast<cl_uint*>(clEnqueueMapBuffer(
+        pStringIdsBuffer = static_cast<cl_double*>(clEnqueueMapBuffer(
             kEnv.mpkCmdQueue, mpClmem, CL_TRUE, CL_MAP_WRITE, 0,
             szHostBuffer, 0, nullptr, nullptr, &err));
         if (CL_SUCCESS != err)
             throw OpenCLError("clEnqueueMapBuffer", err, __FILE__, __LINE__);
 
-        for (size_t i = 0; i < szHostBuffer / sizeof(cl_int); i++)
-            pHashBuffer[i] = 0;
+        for (size_t i = 0; i < szHostBuffer / sizeof(cl_double); i++)
+            rtl::math::setNan(&pStringIdsBuffer[i]);
     }
     err = clEnqueueUnmapMemObject(kEnv.mpkCmdQueue, mpClmem,
-        pHashBuffer, 0, nullptr, nullptr);
+        pStringIdsBuffer, 0, nullptr, nullptr);
     if (CL_SUCCESS != err)
         throw OpenCLError("clEnqueueUnmapMemObject", err, __FILE__, __LINE__);
 
-    SAL_INFO("sc.opencl", "Kernel " << k << " arg " << argno << ": cl_mem: " << mpClmem << " (" << DebugPeekData(ref,mnIndex) << ")");
+    SAL_INFO("sc.opencl", "Kernel " << k << " arg " << argno << ": cl_mem: " << mpClmem
+        << " (stringIds: " << DebugPeekDoubles(pStringIdsBuffer, nStrings) << " "
+        << DebugPeekData(ref,mnIndex) << ")");
     err = clSetKernelArg(k, argno, sizeof(cl_mem), static_cast<void*>(&mpClmem));
     if (CL_SUCCESS != err)
         throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__);
     return 1;
 }
 
+std::string DynamicKernelStringArgument::GenIsString( bool nested ) const
+{
+    if( nested )
+        return "!isnan(" + mSymName + "[gid0])";
+    FormulaToken* ref = mFormulaTree->GetFormulaToken();
+    size_t nStrings = 0;
+    if (ref->GetType() == formula::svSingleVectorRef)
+    {
+        const formula::SingleVectorRefToken* pSVR =
+            static_cast<const formula::SingleVectorRefToken*>(ref);
+        nStrings = pSVR->GetArrayLength();
+    }
+    else if (ref->GetType() == formula::svDoubleVectorRef)
+    {
+        const formula::DoubleVectorRefToken* pDVR =
+            static_cast<const formula::DoubleVectorRefToken*>(ref);
+        nStrings = pDVR->GetArrayLength();
+    }
+    else
+        return "!isnan(" + mSymName + "[gid0])";
+    outputstream ss;
+    ss << "(gid0 < " << nStrings << "? !isnan(" << mSymName << "[gid0]):NAN)";
+    return ss.str();
+}
+
 namespace {
 
 /// A mixed string/numeric vector
@@ -903,18 +916,22 @@ public:
         ss << ")";
         return ss.str();
     }
-    virtual std::string GenDoubleSlidingWindowDeclRef( bool = false ) const override
+    virtual std::string GenDoubleSlidingWindowDeclRef( bool nested = false ) const override
     {
         outputstream ss;
-        ss << VectorRef::GenSlidingWindowDeclRef();
+        ss << VectorRef::GenSlidingWindowDeclRef( nested );
         return ss.str();
     }
-    virtual std::string GenStringSlidingWindowDeclRef( bool = false ) const override
+    virtual std::string GenStringSlidingWindowDeclRef( bool nested = false ) const override
     {
         outputstream ss;
-        ss << mStringArgument.GenSlidingWindowDeclRef();
+        ss << mStringArgument.GenSlidingWindowDeclRef( nested );
         return ss.str();
     }
+    virtual std::string GenIsString( bool nested = false ) const override
+    {
+        return mStringArgument.GenIsString( nested );
+    }
     virtual size_t Marshal( cl_kernel k, int argno, int vw, cl_program p ) override
     {
         int i = VectorRef::Marshal(k, argno, vw, p);
@@ -3513,19 +3530,19 @@ class CLInterpreterContext
     SCROW mnGroupLength;
 
 public:
-    explicit CLInterpreterContext(SCROW nGroupLength)
-        : mpKernel(nullptr)
+    explicit CLInterpreterContext(SCROW nGroupLength, std::shared_ptr<DynamicKernel> pKernel )
+        : mpKernelStore(std::move(pKernel))
+        , mpKernel(mpKernelStore.get())
         , mnGroupLength(nGroupLength) {}
 
-    bool isValid() const
+    ~CLInterpreterContext()
     {
-        return mpKernel != nullptr;
+        DynamicKernelArgument::ClearStringIds();
     }
 
-    void setManagedKernel( std::shared_ptr<DynamicKernel> pKernel )
+    bool isValid() const
     {
-        mpKernelStore = std::move(pKernel);
-        mpKernel = mpKernelStore.get();
+        return mpKernel != nullptr;
     }
 
     CLInterpreterResult launchKernel()
@@ -3571,11 +3588,7 @@ public:
 CLInterpreterContext createCLInterpreterContext( const ScCalcConfig& rConfig,
     const ScFormulaCellGroupRef& xGroup, const ScTokenArray& rCode )
 {
-    CLInterpreterContext aCxt(xGroup->mnLength);
-
-    aCxt.setManagedKernel(DynamicKernel::create(rConfig, rCode, xGroup->mnLength));
-
-    return aCxt;
+    return CLInterpreterContext(xGroup->mnLength, DynamicKernel::create(rConfig, rCode, xGroup->mnLength));
 }
 
 void genRPNTokens( ScDocument& rDoc, const ScAddress& rTopPos, ScTokenArray& rCode )
diff --git a/sc/source/core/opencl/op_math.cxx b/sc/source/core/opencl/op_math.cxx
index 1bf3e231b4a2..d0e033f8843b 100644
--- a/sc/source/core/opencl/op_math.cxx
+++ b/sc/source/core/opencl/op_math.cxx
@@ -481,11 +481,22 @@ void OpEqual::BinInlineFun(std::set<std::string>& decls,
     funs.insert(is_representable_integer);
     decls.insert(approx_equalDecl);
     funs.insert(approx_equal);
+    decls.insert(cell_equalDecl);
+    funs.insert(cell_equal);
 }
 
-void OpEqual::GenerateCode( outputstream& ss ) const
+void OpEqual::GenSlidingWindowFunction(outputstream &ss,
+        const std::string &sSymName, SubArguments &vSubArguments)
 {
-    ss << "    return approx_equal( arg0, arg1 );\n";
+    CHECK_PARAMETER_COUNT( 2, 2 );
+    GenerateFunctionDeclaration( sSymName, vSubArguments, ss );
+    ss << "{\n";
+    ss << "    double tmp = 0;\n";
+    ss << "    int gid0 = get_global_id(0);\n";
+    GenerateArg( 0, vSubArguments, ss, EmptyIsNan, GenerateArgType );
+    GenerateArg( 1, vSubArguments, ss, EmptyIsNan, GenerateArgType );
+    ss << "    return cell_equal( arg0, arg1, arg0_is_string, arg1_is_string );\n";
+    ss << "}";
 }
 
 void OpNotEqual::BinInlineFun(std::set<std::string>& decls,
@@ -495,11 +506,22 @@ void OpNotEqual::BinInlineFun(std::set<std::string>& decls,
     funs.insert(is_representable_integer);
     decls.insert(approx_equalDecl);
     funs.insert(approx_equal);
+    decls.insert(cell_equalDecl);
+    funs.insert(cell_equal);
 }
 
-void OpNotEqual::GenerateCode( outputstream& ss ) const
+void OpNotEqual::GenSlidingWindowFunction(outputstream &ss,
+        const std::string &sSymName, SubArguments &vSubArguments)
 {
-    ss << "    return !approx_equal( arg0, arg1 );\n";
+    CHECK_PARAMETER_COUNT( 2, 2 );
+    GenerateFunctionDeclaration( sSymName, vSubArguments, ss );
+    ss << "{\n";
+    ss << "    double tmp = 0;\n";
+    ss << "    int gid0 = get_global_id(0);\n";
+    GenerateArg( 0, vSubArguments, ss, EmptyIsNan, GenerateArgType );
+    GenerateArg( 1, vSubArguments, ss, EmptyIsNan, GenerateArgType );
+    ss << "    return !cell_equal( arg0, arg1, arg0_is_string, arg1_is_string );\n";
+    ss << "}";
 }
 
 void OpLessEqual::BinInlineFun(std::set<std::string>& decls,
diff --git a/sc/source/core/opencl/op_math.hxx b/sc/source/core/opencl/op_math.hxx
index 8ffe1ee779e5..92f9701eefd0 100644
--- a/sc/source/core/opencl/op_math.hxx
+++ b/sc/source/core/opencl/op_math.hxx
@@ -503,20 +503,24 @@ public:
     virtual void GenerateCode( outputstream& ss ) const override;
 };
 
-class OpEqual : public OpMathTwoArguments
+class OpEqual : public Normal
 {
 public:
-    virtual void GenerateCode( outputstream& ss ) const override;
+    virtual void GenSlidingWindowFunction(outputstream &ss,
+            const std::string &sSymName, SubArguments &vSubArguments) override;
     virtual std::string BinFuncName() const override { return "eq"; }
     virtual void BinInlineFun(std::set<std::string>& , std::set<std::string>& ) override;
+    virtual bool takeString() const override { return true; }
 };
 
-class OpNotEqual : public OpMathTwoArguments
+class OpNotEqual : public Normal
 {
 public:
-    virtual void GenerateCode( outputstream& ss ) const override;
+    virtual void GenSlidingWindowFunction(outputstream &ss,
+            const std::string &sSymName, SubArguments &vSubArguments) override;
     virtual std::string BinFuncName() const override { return "neq"; }
     virtual void BinInlineFun(std::set<std::string>& , std::set<std::string>& ) override;
+    virtual bool takeString() const override { return true; }
 };
 
 class OpLessEqual : public OpMathTwoArguments
diff --git a/sc/source/core/opencl/op_math_helpers.hxx b/sc/source/core/opencl/op_math_helpers.hxx
index 6a751dd9c35b..24aa73c56089 100644
--- a/sc/source/core/opencl/op_math_helpers.hxx
+++ b/sc/source/core/opencl/op_math_helpers.hxx
@@ -189,4 +189,20 @@ const char value_approx[] =
 "    return copysign(fValue, fOrigValue);\n"
 "}\n";
 
+// This compares two cell contents, including possible string equalities (based on string ids
+// coming from DynamicKernelArgument::GetStringId()).
+// The EmptyIsZero conversion must not be have been done for the arguments.
+const char cell_equalDecl[] = "bool cell_equal(double a, double b, bool a_is_string, bool b_is_string);\n";
+const char cell_equal[] =
+"bool cell_equal(double a, double b, bool a_is_string, bool b_is_string) {\n"
+"    if( !a_is_string && !b_is_string )\n"
+"        return approx_equal( isnan(a) ? 0 : a, isnan(b) ? 0 : b );\n"
+"    if( a_is_string && b_is_string )\n"
+"        return a == b;\n"
+// empty strings and empty cells compare equal
+"    if(( a_is_string && a == 0 && isnan(b)) || ( b_is_string && b == 0 && isnan(a)))\n"
+"        return true;\n"
+"    return false;\n"
+"}\n";
+
 /* vim:set shiftwidth=4 softtabstop=4 expandtab: */
diff --git a/sc/source/core/opencl/opbase.cxx b/sc/source/core/opencl/opbase.cxx
index 36605ef8e977..51ff89e043be 100644
--- a/sc/source/core/opencl/opbase.cxx
+++ b/sc/source/core/opencl/opbase.cxx
@@ -11,6 +11,7 @@
 #include <formula/vectortoken.hxx>
 #include <sal/log.hxx>
 #include <utility>
+#include <unordered_map>
 
 #include "opbase.hxx"
 
@@ -83,6 +84,49 @@ bool DynamicKernelArgument::NeedParallelReduction() const
     return false;
 }
 
+// Strings and OpenCL:
+// * Strings are non-trivial types and so passing them to OpenCL and handling them there
+// would be rather complex. However, in practice most string operations are checking
+// string equality, so only such string usage is supported (other cases will be
+// handled by Calc core when they get rejected for OpenCL).
+// * Strings from Calc core come from svl::SharedString, which already ensures that
+// equal strings have equal rtl_uString.
+// * Strings are passed to opencl as integer IDs, each uniquely identifying a different
+// string.
+// * OpenCL code generally handles all values as doubles, so merely converting rtl_uString*
+// to double could lead to loss of precision (double can store 52bits of precision).
+// This could lead to two strings possibly being considered equal by mistake (unlikely,
+// but not impossible). Therefore all rtl_uString* are mapped to internal integer IDs.
+// * Functions that can handle strings properly should override OpBase::takeString()
+// to return true. They should
+// * Empty string Id is 0. Empty cell Id is NAN.
+// * Since strings are marshalled as doubles too, it is important to check whether a value
+// is a real double or a string. Use e.g. GenerateArgType to generate also 'xxx_is_string'
+// variable, there is cell_equal() function to compare two cells.
+
+static std::unordered_map<const rtl_uString*, int>* stringIdsMap;
+
+int DynamicKernelArgument::GetStringId( const rtl_uString* string )
+{
+    assert( string != nullptr );
+    if( string->length == 0 )
+        return 0;
+    if( stringIdsMap == nullptr )
+        stringIdsMap = new std::unordered_map<const rtl_uString*, int>;
+    std::unordered_map<const rtl_uString*, int>::iterator it = stringIdsMap->find( string );
+    if( it != stringIdsMap->end())
+        return it->second;
+    int newId = stringIdsMap->size() + 1;
+    stringIdsMap->insert( std::pair( string, newId ));
+    return newId;
+}
+
+void DynamicKernelArgument::ClearStringIds()
+{
+    delete stringIdsMap;
+    stringIdsMap = nullptr;
+}
+
 VectorRef::VectorRef( const ScCalcConfig& config, const std::string& s, const FormulaTreeNodeRef& ft, int idx ) :
     DynamicKernelArgument(config, s, ft), mpClmem(nullptr), mnIndex(idx), forceStringsToZero( false )
 {
@@ -182,7 +226,7 @@ VectorRefStringsToZero::VectorRefStringsToZero( const ScCalcConfig& config, cons
 }
 
 void SlidingFunctionBase::GenerateArg( const char* name, int arg, SubArguments& vSubArguments,
-    outputstream& ss, EmptyArgType empty )
+    outputstream& ss, EmptyArgType empty, GenerateArgTypeType generateType )
 {
     assert( arg < int( vSubArguments.size()));
     FormulaToken *token = vSubArguments[arg]->GetFormulaToken();
@@ -195,9 +239,19 @@ void SlidingFunctionBase::GenerateArg( const char* name, int arg, SubArguments&
             const formula::SingleVectorRefToken* svr =
                 static_cast<const formula::SingleVectorRefToken *>(token);
             ss << "    double " << name << " = NAN;\n";
+            if( generateType == GenerateArgType )
+                ss << "    bool " << name << "_is_string = false;\n";
             ss << "    if (gid0 < " << svr->GetArrayLength() << ")\n";
+            if( generateType == GenerateArgType )
+                ss << "    {\n";
             ss << "        " << name << " = ";
-            ss << vSubArguments[arg]->GenSlidingWindowDeclRef() << ";\n";
+            ss << vSubArguments[arg]->GenSlidingWindowDeclRef( true ) << ";\n";
+            if( generateType == GenerateArgType )
+            {
+                ss << "        " << name << "_is_string = ";
+                ss << vSubArguments[arg]->GenIsString( true ) << ";\n";
+                ss << "    }\n";
+            }
             switch( empty )
             {
                 case EmptyIsZero:
@@ -212,11 +266,22 @@ void SlidingFunctionBase::GenerateArg( const char* name, int arg, SubArguments&
             }
         }
         else if(token->GetType() == formula::svDouble)
+        {
             ss << "    double " << name << " = " << token->GetDouble() << ";\n";
+            if( generateType == GenerateArgType )
+                ss << "    bool " << name << "_is_string = "
+                    << vSubArguments[arg]->GenIsString() << ";\n";
+        }
         else if(token->GetType() == formula::svString)
         {
-            assert( dynamic_cast<DynamicKernelStringToZeroArgument*>(vSubArguments[arg].get()));
+            if( forceStringsToZero())
+                assert( dynamic_cast<DynamicKernelStringToZeroArgument*>(vSubArguments[arg].get()));
+            else if( !takeString())
+                throw Unhandled( __FILE__, __LINE__ );
             ss << "    double " << name << " = 0.0;\n";
+            if( generateType == GenerateArgType )
+                ss << "    bool " << name << "_is_string = "
+                    << vSubArguments[arg]->GenIsString() << ";\n";
         }
         else
             throw Unhandled( __FILE__, __LINE__ );
@@ -225,15 +290,18 @@ void SlidingFunctionBase::GenerateArg( const char* name, int arg, SubArguments&
     {
         ss << "    double " << name << " = ";
         ss << vSubArguments[arg]->GenSlidingWindowDeclRef() << ";\n";
+        if( generateType == GenerateArgType )
+            ss << "    bool " << name << "_is_string = "
+                << vSubArguments[arg]->GenIsString() << ";\n";
     }
 }
 
 void SlidingFunctionBase::GenerateArg( int arg, SubArguments& vSubArguments, outputstream& ss,
-    EmptyArgType empty )
+    EmptyArgType empty, GenerateArgTypeType generateType )
 {
     char buf[ 30 ];
     sprintf( buf, "arg%d", arg );
-    GenerateArg( buf, arg, vSubArguments, ss, empty );
+    GenerateArg( buf, arg, vSubArguments, ss, empty, generateType );
 }
 
 void SlidingFunctionBase::GenerateArgWithDefault( const char* name, int arg, double def,
diff --git a/sc/source/core/opencl/opbase.hxx b/sc/source/core/opencl/opbase.hxx
index 8e1a167258cf..04ebab1ee7ed 100644
--- a/sc/source/core/opencl/opbase.hxx
+++ b/sc/source/core/opencl/opbase.hxx
@@ -153,6 +153,9 @@ public:
     /// When Mix, it will be called
     virtual std::string GenStringSlidingWindowDeclRef( bool = false ) const;
 
+    /// Will generate value saying whether the value is a string.
+    virtual std::string GenIsString( bool = false ) const { return "false"; }
+
     /// Generate use/references to the argument
     virtual void GenDeclRef( outputstream& ss ) const;
 
@@ -165,10 +168,13 @@ public:
     /// If there's actually no argument, i.e. it expands to no code.
     virtual bool IsEmpty() const { return false; }
 
+    static void ClearStringIds();
+
 protected:
     const ScCalcConfig& mCalcConfig;
     std::string mSymName;
     FormulaTreeNodeRef mFormulaTree;
+    static int GetStringId( const rtl_uString* string );
 };
 
 typedef std::shared_ptr<DynamicKernelArgument> DynamicKernelArgumentRef;
@@ -233,12 +239,13 @@ public:
     /// Generate declaration
     virtual void GenDecl( outputstream& ss ) const override
     {
-        ss << "__global unsigned int *" << mSymName;
+        ss << "__global double *" << mSymName;
     }
     virtual void GenSlidingWindowDecl( outputstream& ss ) const override
     {
         DynamicKernelStringArgument::GenDecl(ss);
     }
+    virtual std::string GenIsString( bool = false ) const override;
     virtual size_t Marshal( cl_kernel, int, int, cl_program ) override;
 };
 
@@ -348,17 +355,24 @@ protected:
         EmptyIsNan,  // empty cells become NAN, use isnan() to check in code
         SkipEmpty    // empty cells will be skipped
     };
+    // This enum controls whether the generated code will also include variable
+    // <name>_is_string that will be set depending on the value type.
+    enum GenerateArgTypeType
+    {
+        DoNotGenerateArgType,
+        GenerateArgType
+    };
     void GenerateFunctionDeclaration( const std::string& sSymName,
         SubArguments& vSubArguments, outputstream& ss );
     // Generate code for "double <name> = <value>;" from vSubArguments, svDoubleVectorRef is not supported.
-    static void GenerateArg( const char* name, int arg, SubArguments& vSubArguments, outputstream& ss,
-        EmptyArgType empty = EmptyIsZero );
+    void GenerateArg( const char* name, int arg, SubArguments& vSubArguments, outputstream& ss,
+        EmptyArgType empty = EmptyIsZero, GenerateArgTypeType generateType = DoNotGenerateArgType );
     // overload, variable will be named "arg<arg>"
-    static void GenerateArg( int arg, SubArguments& vSubArguments, outputstream& ss,
-        EmptyArgType empty = EmptyIsZero );
+    void GenerateArg( int arg, SubArguments& vSubArguments, outputstream& ss,
+        EmptyArgType empty = EmptyIsZero, GenerateArgTypeType generateType = DoNotGenerateArgType );
     // generate code for "double <name> = <value>;" from vSubArguments, if it exists,
     // otherwise set to <def>
-    static void GenerateArgWithDefault( const char* name, int arg, double def, SubArguments& vSubArguments,
+    void GenerateArgWithDefault( const char* name, int arg, double def, SubArguments& vSubArguments,
         outputstream& ss, EmptyArgType empty = EmptyIsZero );
     // Generate code that will handle all arguments firstArg-lastArg (zero-based, inclusive),
     // including range arguments (svDoubleVectorRef) and each value will be processed by 'code',