diff options
author | Kohei Yoshida <kohei.yoshida@collabora.com> | 2013-10-29 16:40:57 -0400 |
---|---|---|
committer | Kohei Yoshida <kohei.yoshida@collabora.com> | 2013-10-29 17:00:31 -0400 |
commit | f571c985c0cef19b821ca34582de6abcbf940e0e (patch) | |
tree | b2e91186fa2302a0b641d89b672b7645d23762c3 /sc | |
parent | b0aec6296d480a4444203032cf05e0ec71259b6d (diff) |
Properly split opbase.hxx into header and source files.
Change-Id: I6ab5a792b5c31b44b8f302e45ec3511d3e2d6cdc
Diffstat (limited to 'sc')
-rw-r--r-- | sc/Library_scopencl.mk | 1 | ||||
-rw-r--r-- | sc/source/core/opencl/op_financial.cxx | 1 | ||||
-rw-r--r-- | sc/source/core/opencl/op_math.cxx | 1 | ||||
-rw-r--r-- | sc/source/core/opencl/op_statistical.cxx | 1 | ||||
-rw-r--r-- | sc/source/core/opencl/opbase.cxx | 107 | ||||
-rw-r--r-- | sc/source/core/opencl/opbase.hxx | 135 |
6 files changed, 147 insertions, 99 deletions
diff --git a/sc/Library_scopencl.mk b/sc/Library_scopencl.mk index 0efe3cf61532..2fb8cfc81f0b 100644 --- a/sc/Library_scopencl.mk +++ b/sc/Library_scopencl.mk @@ -38,6 +38,7 @@ $(eval $(call gb_Library_use_libraries,scopencl,\ $(eval $(call gb_Library_add_exception_objects,scopencl,\ sc/source/core/opencl/formulagroupcl \ sc/source/core/opencl/openclwrapper \ + sc/source/core/opencl/opbase \ sc/source/core/opencl/op_financial \ sc/source/core/opencl/op_database \ sc/source/core/opencl/op_math \ diff --git a/sc/source/core/opencl/op_financial.cxx b/sc/source/core/opencl/op_financial.cxx index a638530e98ef..04d505a43735 100644 --- a/sc/source/core/opencl/op_financial.cxx +++ b/sc/source/core/opencl/op_financial.cxx @@ -15,6 +15,7 @@ #include "tokenarray.hxx" #include "compiler.hxx" #include "interpre.hxx" +#include "formula/vectortoken.hxx" #include <list> #include <iostream> #include <sstream> diff --git a/sc/source/core/opencl/op_math.cxx b/sc/source/core/opencl/op_math.cxx index 03fc9d3edb28..83f780a9cb37 100644 --- a/sc/source/core/opencl/op_math.cxx +++ b/sc/source/core/opencl/op_math.cxx @@ -15,6 +15,7 @@ #include "tokenarray.hxx" #include "compiler.hxx" #include "interpre.hxx" +#include "formula/vectortoken.hxx" #include <list> #include <iostream> #include <sstream> diff --git a/sc/source/core/opencl/op_statistical.cxx b/sc/source/core/opencl/op_statistical.cxx index 221d2419e7bf..ec816c2df59b 100644 --- a/sc/source/core/opencl/op_statistical.cxx +++ b/sc/source/core/opencl/op_statistical.cxx @@ -15,6 +15,7 @@ #include "tokenarray.hxx" #include "compiler.hxx" #include "interpre.hxx" +#include "formula/vectortoken.hxx" #include <list> #include <iostream> #include <sstream> diff --git a/sc/source/core/opencl/opbase.cxx b/sc/source/core/opencl/opbase.cxx new file mode 100644 index 000000000000..eef9eba25547 --- /dev/null +++ b/sc/source/core/opencl/opbase.cxx @@ -0,0 +1,107 @@ +/* -*- Mode: C++; tab-width: 4; indent-tabs-mode: nil; c-basic-offset: 4 -*- */ +/* + * This file is part of the LibreOffice project. + * + * This Source Code Form is subject to the terms of the Mozilla Public + * License, v. 2.0. If a copy of the MPL was not distributed with this + * file, You can obtain one at http://mozilla.org/MPL/2.0/. + */ + +#include "opbase.hxx" + +#include "formula/vectortoken.hxx" + +using namespace formula; + +namespace sc { namespace opencl { + +DynamicKernelArgument::DynamicKernelArgument(const std::string &s, + std::shared_ptr<FormulaTreeNode> ft): + mSymName(s), mFormulaTree(ft), mpClmem(NULL) {} + +/// Generate declaration +void DynamicKernelArgument::GenDecl(std::stringstream &ss) const +{ + ss << "__global double *"<<mSymName; +} + +/// When declared as input to a sliding window function +void DynamicKernelArgument::GenSlidingWindowDecl(std::stringstream &ss) const +{ + GenDecl(ss); +} + +/// When referenced in a sliding window function +std::string DynamicKernelArgument::GenSlidingWindowDeclRef(bool) const +{ + std::stringstream ss; + ss << mSymName << "[gid0]"; + return ss.str(); +} + +/// Generate use/references to the argument +void DynamicKernelArgument::GenDeclRef(std::stringstream &ss) const +{ + ss << mSymName; +} + +DynamicKernelArgument::~DynamicKernelArgument() +{ + //std::cerr << "~DynamicKernelArgument: " << mSymName <<"\n"; + if (mpClmem) { + //std::cerr << "\tFreeing cl_mem of " << mSymName <<"\n"; + cl_int ret = clReleaseMemObject(mpClmem); + if (ret != CL_SUCCESS) + throw OpenCLError(ret); + } +} + +FormulaToken* DynamicKernelArgument::GetFormulaToken(void) const +{ + return mFormulaTree->GetFormulaToken(); +} + +size_t DynamicKernelArgument::GetWindowSize(void) const +{ + FormulaToken *pCur = mFormulaTree->GetFormulaToken(); + assert(pCur); + if (auto *pCurDVR = + dynamic_cast<const formula::DoubleVectorRefToken *>(pCur)) + { + return pCurDVR->GetRefRowSize(); + } else if (dynamic_cast<const formula::SingleVectorRefToken *>(pCur)) + { + // Prepare intermediate results (on CPU for now) + return 1; + } else { + throw Unhandled(); + } + return 0; +} + +void Normal::GenSlidingWindowFunction( + std::stringstream &ss, const std::string sSymName, SubArguments &vSubArguments) +{ + ArgVector argVector; + ss << "\ndouble " << sSymName; + ss << "_"<< BinFuncName() <<"("; + for (unsigned i = 0; i < vSubArguments.size(); i++) + { + if (i) + ss << ","; + vSubArguments[i]->GenSlidingWindowDecl(ss); + argVector.push_back(vSubArguments[i]->GenSlidingWindowDeclRef()); + } + ss << ") {\n\t"; + ss << "double tmp = " << GetBottom() <<";\n\t"; + ss << "int gid0 = get_global_id(0);\n\t"; + ss << "tmp = "; + ss << Gen(argVector); + ss << ";\n\t"; + ss << "return tmp;\n"; + ss << "}"; +} + +}} + +/* vim:set shiftwidth=4 softtabstop=4 expandtab: */ diff --git a/sc/source/core/opencl/opbase.hxx b/sc/source/core/opencl/opbase.hxx index acde3a974362..5f17d85e7a21 100644 --- a/sc/source/core/opencl/opbase.hxx +++ b/sc/source/core/opencl/opbase.hxx @@ -10,44 +10,34 @@ #ifndef SC_OPENCL_OPBASE_HXX #define SC_OPENCL_OPBASE_HXX -#include "formulagroup.hxx" -#include "document.hxx" -#include "formulacell.hxx" -#include "tokenarray.hxx" -#include "compiler.hxx" -#include "interpre.hxx" -#include "formula/vectortoken.hxx" - #include "clcc/clew.h" -#include <list> -#include <iostream> -#include <sstream> -#include <algorithm> +#include "formula/token.hxx" -#include <memory> #define ISNAN -using namespace formula; namespace sc { namespace opencl { /// Exceptions /// Failed in parsing -class UnhandledToken { +class UnhandledToken +{ public: - UnhandledToken(FormulaToken *t): mToken(t) {} - FormulaToken *mToken; + UnhandledToken(formula::FormulaToken *t): mToken(t) {} + formula::FormulaToken *mToken; }; /// Failed in marshaling -class OpenCLError { +class OpenCLError +{ public: OpenCLError(cl_int err): mError(err) {} cl_int mError; }; /// Inconsistent state -class Unhandled { +class Unhandled +{ public: Unhandled() {} }; @@ -55,17 +45,17 @@ public: class FormulaTreeNode { public: - FormulaTreeNode(FormulaToken *ft): mpCurrentFormula(ft) + FormulaTreeNode(formula::FormulaToken *ft): mpCurrentFormula(ft) { Children.reserve(8); } std::vector<std::shared_ptr<FormulaTreeNode>> Children; - FormulaToken *GetFormulaToken(void) const + formula::FormulaToken *GetFormulaToken(void) const { return mpCurrentFormula; } private: - FormulaToken *const mpCurrentFormula; + formula::FormulaToken *const mpCurrentFormula; }; /// Holds an input (read-only) argument reference to a SingleVectorRef. @@ -73,69 +63,33 @@ private: /// like SumOfProduct /// In most of the cases the argument is introduced /// by a Push operation in the given RPN. -class DynamicKernelArgument { +class DynamicKernelArgument +{ public: - DynamicKernelArgument(const std::string &s, - std::shared_ptr<FormulaTreeNode> ft): - mSymName(s), mFormulaTree(ft), mpClmem(NULL) {} + DynamicKernelArgument(const std::string &s, std::shared_ptr<FormulaTreeNode> ft); + const std::string &GetNameAsString(void) const { return mSymName; } /// Generate declaration - virtual void GenDecl(std::stringstream &ss) const - { - ss << "__global double *"<<mSymName; - } + virtual void GenDecl(std::stringstream &ss) const; + /// When declared as input to a sliding window function - virtual void GenSlidingWindowDecl(std::stringstream &ss) const - { - GenDecl(ss); - } + virtual void GenSlidingWindowDecl(std::stringstream &ss) const; + /// When referenced in a sliding window function - virtual std::string GenSlidingWindowDeclRef(bool=false) const - { - std::stringstream ss; - ss << mSymName << "[gid0]"; - return ss.str(); - } + virtual std::string GenSlidingWindowDeclRef(bool=false) const; + /// Generate use/references to the argument - virtual void GenDeclRef(std::stringstream &ss) const - { - ss << mSymName; - } + virtual void GenDeclRef(std::stringstream &ss) const; + /// Create buffer and pass the buffer to a given kernel virtual size_t Marshal(cl_kernel, int, int); - virtual ~DynamicKernelArgument() - { - //std::cerr << "~DynamicKernelArgument: " << mSymName <<"\n"; - if (mpClmem) { - //std::cerr << "\tFreeing cl_mem of " << mSymName <<"\n"; - cl_int ret = clReleaseMemObject(mpClmem); - if (ret != CL_SUCCESS) - throw OpenCLError(ret); - } - } - virtual void GenSlidingWindowFunction(std::stringstream &) {}; + + virtual ~DynamicKernelArgument(); + + virtual void GenSlidingWindowFunction(std::stringstream &) {} const std::string &GetSymName(void) const { return mSymName; } - FormulaToken *GetFormulaToken(void) const - { - return mFormulaTree->GetFormulaToken(); - } - virtual size_t GetWindowSize(void) const - { - FormulaToken *pCur = mFormulaTree->GetFormulaToken(); - assert(pCur); - if (auto *pCurDVR = - dynamic_cast<const formula::DoubleVectorRefToken *>(pCur)) - { - return pCurDVR->GetRefRowSize(); - } else if (dynamic_cast<const formula::SingleVectorRefToken *>(pCur)) - { - // Prepare intermediate results (on CPU for now) - return 1; - } else { - throw Unhandled(); - } - return 0; - } + formula::FormulaToken *GetFormulaToken(void) const; + virtual size_t GetWindowSize(void) const; virtual std::string DumpOpName(void) const { return std::string(""); } const std::string& GetName(void) const { return mSymName; } protected: @@ -147,7 +101,8 @@ protected: /// Abstract class for code generation -class SlidingFunctionBase { +class SlidingFunctionBase +{ public: typedef std::unique_ptr<DynamicKernelArgument> SubArgument; typedef std::vector<SubArgument> SubArguments; @@ -156,7 +111,8 @@ public: virtual ~SlidingFunctionBase() {}; }; -class OpBase { +class OpBase +{ public: typedef std::vector<std::string> ArgVector; typedef std::vector<std::string>::iterator ArgVectorIter; @@ -172,28 +128,9 @@ class Normal: public SlidingFunctionBase, public OpBase { public: virtual void GenSlidingWindowFunction(std::stringstream &ss, - const std::string sSymName, SubArguments &vSubArguments) - { - ArgVector argVector; - ss << "\ndouble " << sSymName; - ss << "_"<< BinFuncName() <<"("; - for (unsigned i = 0; i < vSubArguments.size(); i++) - { - if (i) - ss << ","; - vSubArguments[i]->GenSlidingWindowDecl(ss); - argVector.push_back(vSubArguments[i]->GenSlidingWindowDeclRef()); - } - ss << ") {\n\t"; - ss << "double tmp = " << GetBottom() <<";\n\t"; - ss << "int gid0 = get_global_id(0);\n\t"; - ss << "tmp = "; - ss << Gen(argVector); - ss << ";\n\t"; - ss << "return tmp;\n"; - ss << "}"; - } + const std::string sSymName, SubArguments &vSubArguments); }; + }} #endif |