From 2549b21b5226c30d900f5994d45b7b08c15935d6 Mon Sep 17 00:00:00 2001 From: Kohei Yoshida Date: Fri, 3 Oct 2014 12:01:51 -0400 Subject: Run 'beautify' on these files as well. And some trivial manual fixing afterward... Change-Id: Ib60411f689b7971a2dd1acd1541864f38b34b570 --- sc/source/core/opencl/opbase.cxx | 318 ++++++++++++++++++++------------------- sc/source/core/opencl/opbase.hxx | 248 +++++++++++++++--------------- 2 files changed, 292 insertions(+), 274 deletions(-) (limited to 'sc/source') diff --git a/sc/source/core/opencl/opbase.cxx b/sc/source/core/opencl/opbase.cxx index b0751f53eea5..17996e9d4c08 100644 --- a/sc/source/core/opencl/opbase.cxx +++ b/sc/source/core/opencl/opbase.cxx @@ -13,22 +13,22 @@ using namespace formula; namespace sc { namespace opencl { -DynamicKernelArgument::DynamicKernelArgument(const std::string &s, - FormulaTreeNodeRef ft): - mSymName(s), mFormulaTree(ft) {} +DynamicKernelArgument::DynamicKernelArgument( const std::string& s, + FormulaTreeNodeRef ft ) : + mSymName(s), mFormulaTree(ft) { } /// Generate use/references to the argument -void DynamicKernelArgument::GenDeclRef(std::stringstream &ss) const +void DynamicKernelArgument::GenDeclRef( std::stringstream& ss ) const { ss << mSymName; } -FormulaToken* DynamicKernelArgument::GetFormulaToken(void) const +FormulaToken* DynamicKernelArgument::GetFormulaToken( void ) const { return mFormulaTree->GetFormulaToken(); } -VectorRef::VectorRef(const std::string &s, FormulaTreeNodeRef ft, int idx): +VectorRef::VectorRef( const std::string& s, FormulaTreeNodeRef ft, int idx ) : DynamicKernelArgument(s, ft), mpClmem(NULL), mnIndex(idx) { if (mnIndex) @@ -41,47 +41,48 @@ VectorRef::VectorRef(const std::string &s, FormulaTreeNodeRef ft, int idx): VectorRef::~VectorRef() { - if (mpClmem) { + if (mpClmem) + { clReleaseMemObject(mpClmem); } } /// Generate declaration -void VectorRef::GenDecl(std::stringstream &ss) const +void VectorRef::GenDecl( std::stringstream& ss ) const { - ss << "__global double *"<(DynamicKernelArgument::GetFormulaToken()); - if (pSVR&&!nested) + if (pSVR && !nested) ss << "(gid0 < " << pSVR->GetArrayLength() << "?"; ss << mSymName << "[gid0]"; - if (pSVR&&!nested) + if (pSVR && !nested) ss << ":NAN)"; return ss.str(); } -size_t VectorRef::GetWindowSize(void) const +size_t VectorRef::GetWindowSize( void ) const { - FormulaToken *pCur = mFormulaTree->GetFormulaToken(); + FormulaToken* pCur = mFormulaTree->GetFormulaToken(); assert(pCur); if (const formula::DoubleVectorRefToken* pCurDVR = - dynamic_cast(pCur)) + dynamic_cast(pCur)) { return pCurDVR->GetRefRowSize(); } - else if (dynamic_cast(pCur)) + else if (dynamic_cast(pCur)) { // Prepare intermediate results (on CPU for now) return 1; @@ -93,11 +94,11 @@ size_t VectorRef::GetWindowSize(void) const } void Normal::GenSlidingWindowFunction( - std::stringstream &ss, const std::string &sSymName, SubArguments &vSubArguments) + std::stringstream& ss, const std::string& sSymName, SubArguments& vSubArguments ) { ArgVector argVector; ss << "\ndouble " << sSymName; - ss << "_"<< BinFuncName() <<"("; + ss << "_" << BinFuncName() << "("; for (unsigned i = 0; i < vSubArguments.size(); i++) { if (i) @@ -106,7 +107,7 @@ void Normal::GenSlidingWindowFunction( argVector.push_back(vSubArguments[i]->GenSlidingWindowDeclRef()); } ss << ") {\n\t"; - ss << "double tmp = " << GetBottom() <<";\n\t"; + ss << "double tmp = " << GetBottom() << ";\n\t"; ss << "int gid0 = get_global_id(0);\n\t"; ss << "tmp = "; ss << Gen(argVector); @@ -116,193 +117,208 @@ void Normal::GenSlidingWindowFunction( } void CheckVariables::GenTmpVariables( - std::stringstream & ss, SubArguments & vSubArguments) + std::stringstream& ss, SubArguments& vSubArguments ) { - for(unsigned i=0;iGetFormulaToken()->GetType() == - formula::svSingleVectorRef) - { - const formula::SingleVectorRefToken*pTmpDVR1= static_cast(vSubArguments[i]->GetFormulaToken()); - ss<< " if(singleIndex>="; - ss<< pTmpDVR1->GetArrayLength(); - ss<<" ||"; - ss<< "isNan("; - ss<< vSubArguments[i]->GenSlidingWindowDeclRef(true); - ss<<"))\n"; - ss<< " tmp"; - ss<< i; - ss <<"=0;\n else \n"; - ss <<" tmp"; - ss <GetFormulaToken()->GetType() == - formula::svDoubleVectorRef) - { - const formula::DoubleVectorRefToken*pTmpDVR2= static_cast(vSubArguments[i]->GetFormulaToken()); - ss<< " if(doubleIndex>="; - ss<< pTmpDVR2->GetArrayLength(); - ss<<" ||"; - ss<< "isNan("; - ss<< vSubArguments[i]->GenSlidingWindowDeclRef(false); - ss<<"))\n"; - ss<< " tmp"; - ss<< i; - ss <<"=0;\n else \n"; - ss <<" tmp"; - ss <GetFormulaToken()->GetType() == formula::svDouble || - vSubArguments[i]->GetFormulaToken()->GetOpCode() != ocPush) - { - ss<< " if("; - ss<< "isNan("; - ss<< vSubArguments[i]->GenSlidingWindowDeclRef(); - ss<<"))\n"; - ss<< " tmp"; - ss<< i; - ss <<"=0;\n else \n"; - ss <<" tmp"; - ss <GetFormulaToken()->GetType() == + formula::svSingleVectorRef) + { + const formula::SingleVectorRefToken* pTmpDVR1 = + static_cast(vSubArguments[i]->GetFormulaToken()); + ss << " if(singleIndex>="; + ss << pTmpDVR1->GetArrayLength(); + ss << " ||"; + ss << "isNan("; + ss << vSubArguments[i]->GenSlidingWindowDeclRef(true); + ss << "))\n"; + ss << " tmp"; + ss << i; + ss << "=0;\n else \n"; + ss << " tmp"; + ss << i; + ss << "="; + ss << vSubArguments[i]->GenSlidingWindowDeclRef(true); + ss << ";\n"; + } + if (vSubArguments[i]->GetFormulaToken()->GetType() == + formula::svDoubleVectorRef) + { + const formula::DoubleVectorRefToken* pTmpDVR2 = + static_cast(vSubArguments[i]->GetFormulaToken()); + ss << " if(doubleIndex>="; + ss << pTmpDVR2->GetArrayLength(); + ss << " ||"; + ss << "isNan("; + ss << vSubArguments[i]->GenSlidingWindowDeclRef(false); + ss << "))\n"; + ss << " tmp"; + ss << i; + ss << "=0;\n else \n"; + ss << " tmp"; + ss << i; + ss << "="; + ss << vSubArguments[i]->GenSlidingWindowDeclRef(false); + ss << ";\n"; + } + if (vSubArguments[i]->GetFormulaToken()->GetType() == formula::svDouble || + vSubArguments[i]->GetFormulaToken()->GetOpCode() != ocPush) + { + ss << " if("; + ss << "isNan("; + ss << vSubArguments[i]->GenSlidingWindowDeclRef(); + ss << "))\n"; + ss << " tmp"; + ss << i; + ss << "=0;\n else \n"; + ss << " tmp"; + ss << i; + ss << "="; + ss << vSubArguments[i]->GenSlidingWindowDeclRef(); + ss << ";\n"; - } + } } -void CheckVariables::CheckSubArgumentIsNan2( std::stringstream & ss, - SubArguments &vSubArguments, int argumentNum, std::string p) +void CheckVariables::CheckSubArgumentIsNan2( std::stringstream& ss, + SubArguments& vSubArguments, int argumentNum, std::string p ) { int i = argumentNum; - if(vSubArguments[i]->GetFormulaToken()->GetType() == formula::svDouble) + if (vSubArguments[i]->GetFormulaToken()->GetType() == formula::svDouble) { - ss <<" tmp"; - ss <GenDeclRef(ss); - if(vSubArguments[i]->GetFormulaToken()->GetType() == - formula::svDoubleVectorRef) - ss<<"["<< p.c_str()<< "]"; - else if(vSubArguments[i]->GetFormulaToken()->GetType() == - formula::svSingleVectorRef) - ss<<"[get_group_id(1)]"; - ss<<", 0);\n"; + if (vSubArguments[i]->GetFormulaToken()->GetType() == + formula::svDoubleVectorRef) + ss << "[" << p.c_str() << "]"; + else if (vSubArguments[i]->GetFormulaToken()->GetType() == + formula::svSingleVectorRef) + ss << "[get_group_id(1)]"; + ss << ", 0);\n"; #else - ss <<" tmp"; - ss <GetFormulaToken()->GetType() == + ss << "[" << p.c_str() << "]"; + else if (vSubArguments[i]->GetFormulaToken()->GetType() == formula::svSingleVectorRef) - ss<<"[get_group_id(1)]"; + ss << "[get_group_id(1)]"; - ss<<";\n"; + ss << ";\n"; #endif } void CheckVariables::CheckAllSubArgumentIsNan( - std::stringstream & ss, SubArguments & vSubArguments) + std::stringstream& ss, SubArguments& vSubArguments ) { - ss<<" int k = gid0;\n"; - for(unsigned i=0;iIsStartFixed() && pCurDVR->IsEndFixed()) { - ss << " loop = ("<IsStartFixed() && !pCurDVR->IsEndFixed()) { - ss << " loop = ("<IsStartFixed() && pCurDVR->IsEndFixed()) + { + ss << " loop = (" << nCurWindowSize << " - gid0)/"; + ss << unrollSize << ";\n"; + } + else if (pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed()) + { + ss << " loop = (" << nCurWindowSize << " + gid0)/"; + ss << unrollSize << ";\n"; - } else { - ss << " loop = "<IsStartFixed() && pCurDVR->IsEndFixed()) + { + ss << "gid0 + j * " << unrollSize << ";\n"; + } + else + { + ss << "j * " << unrollSize << ";\n"; } - if(!pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed()) + if (!pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed()) { - ss << " int doubleIndex = i+gid0;\n"; - }else + ss << " int doubleIndex = i+gid0;\n"; + } + else { - ss << " int doubleIndex = i;\n"; + ss << " int doubleIndex = i;\n"; } - for(int j =0;j < unrollSize;j++) + for (int j = 0; j < unrollSize; j++) { ss << unrollstr.str(); ss << "i++;\n"; ss << "doubleIndex++;\n"; } - ss << " }\n"; - ss << " for (int i = "; - if (!pCurDVR->IsStartFixed() && pCurDVR->IsEndFixed()) { - ss << "gid0 + loop *"<IsStartFixed() && !pCurDVR->IsEndFixed()) { - ss << "0 + loop *"<IsStartFixed() && !pCurDVR->IsEndFixed()) - { + ss << " }\n"; + ss << " for (int i = "; + if (!pCurDVR->IsStartFixed() && pCurDVR->IsEndFixed()) + { + ss << "gid0 + loop *" << unrollSize << "; i < "; + ss << nCurWindowSize << "; i++)\n"; + } + else if (pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed()) + { + ss << "0 + loop *" << unrollSize << "; i < gid0+"; + ss << nCurWindowSize << "; i++)\n"; + } + else + { + ss << "0 + loop *" << unrollSize << "; i < "; + ss << nCurWindowSize << "; i++)\n"; + } + ss << " {\n"; + if (!pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed()) + { ss << " int doubleIndex = i+gid0;\n"; - }else - { + } + else + { ss << " int doubleIndex = i;\n"; - } - ss << unrollstr.str(); - ss << " }\n"; + } + ss << unrollstr.str(); + ss << " }\n"; } }} diff --git a/sc/source/core/opencl/opbase.hxx b/sc/source/core/opencl/opbase.hxx index 487fc04b09ee..795b7bb6623b 100644 --- a/sc/source/core/opencl/opbase.hxx +++ b/sc/source/core/opencl/opbase.hxx @@ -31,10 +31,11 @@ class FormulaTreeNode; class UnhandledToken { public: - UnhandledToken(formula::FormulaToken *t, - const char *const m, std::string fn="", int ln=0): - mToken(t), mMessage(m), mFile(fn), mLineNumber(ln) {} - formula::FormulaToken *mToken; + UnhandledToken( formula::FormulaToken* t, + const char* m, const std::string& fn = "", int ln = 0 ) : + mToken(t), mMessage(m), mFile(fn), mLineNumber(ln) {} + + formula::FormulaToken* mToken; std::string mMessage; std::string mFile; int mLineNumber; @@ -44,67 +45,67 @@ public: class OpenCLError { private: - const char *strerror(cl_int i) + const char* strerror( cl_int i ) { #define CASE(val) case val: return #val switch (i) { - CASE(CL_SUCCESS); - CASE(CL_DEVICE_NOT_FOUND); - CASE(CL_DEVICE_NOT_AVAILABLE); - CASE(CL_COMPILER_NOT_AVAILABLE); - CASE(CL_MEM_OBJECT_ALLOCATION_FAILURE); - CASE(CL_OUT_OF_RESOURCES); - CASE(CL_OUT_OF_HOST_MEMORY); - CASE(CL_PROFILING_INFO_NOT_AVAILABLE); - CASE(CL_MEM_COPY_OVERLAP); - CASE(CL_IMAGE_FORMAT_MISMATCH); - CASE(CL_IMAGE_FORMAT_NOT_SUPPORTED); - CASE(CL_BUILD_PROGRAM_FAILURE); - CASE(CL_MAP_FAILURE); - CASE(CL_INVALID_VALUE); - CASE(CL_INVALID_DEVICE_TYPE); - CASE(CL_INVALID_PLATFORM); - CASE(CL_INVALID_DEVICE); - CASE(CL_INVALID_CONTEXT); - CASE(CL_INVALID_QUEUE_PROPERTIES); - CASE(CL_INVALID_COMMAND_QUEUE); - CASE(CL_INVALID_HOST_PTR); - CASE(CL_INVALID_MEM_OBJECT); - CASE(CL_INVALID_IMAGE_FORMAT_DESCRIPTOR); - CASE(CL_INVALID_IMAGE_SIZE); - CASE(CL_INVALID_SAMPLER); - CASE(CL_INVALID_BINARY); - CASE(CL_INVALID_BUILD_OPTIONS); - CASE(CL_INVALID_PROGRAM); - CASE(CL_INVALID_PROGRAM_EXECUTABLE); - CASE(CL_INVALID_KERNEL_NAME); - CASE(CL_INVALID_KERNEL_DEFINITION); - CASE(CL_INVALID_KERNEL); - CASE(CL_INVALID_ARG_INDEX); - CASE(CL_INVALID_ARG_VALUE); - CASE(CL_INVALID_ARG_SIZE); - CASE(CL_INVALID_KERNEL_ARGS); - CASE(CL_INVALID_WORK_DIMENSION); - CASE(CL_INVALID_WORK_GROUP_SIZE); - CASE(CL_INVALID_WORK_ITEM_SIZE); - CASE(CL_INVALID_GLOBAL_OFFSET); - CASE(CL_INVALID_EVENT_WAIT_LIST); - CASE(CL_INVALID_EVENT); - CASE(CL_INVALID_OPERATION); - CASE(CL_INVALID_GL_OBJECT); - CASE(CL_INVALID_BUFFER_SIZE); - CASE(CL_INVALID_MIP_LEVEL); - CASE(CL_INVALID_GLOBAL_WORK_SIZE); - default: - return "Unknown OpenCL error code"; + CASE(CL_SUCCESS); + CASE(CL_DEVICE_NOT_FOUND); + CASE(CL_DEVICE_NOT_AVAILABLE); + CASE(CL_COMPILER_NOT_AVAILABLE); + CASE(CL_MEM_OBJECT_ALLOCATION_FAILURE); + CASE(CL_OUT_OF_RESOURCES); + CASE(CL_OUT_OF_HOST_MEMORY); + CASE(CL_PROFILING_INFO_NOT_AVAILABLE); + CASE(CL_MEM_COPY_OVERLAP); + CASE(CL_IMAGE_FORMAT_MISMATCH); + CASE(CL_IMAGE_FORMAT_NOT_SUPPORTED); + CASE(CL_BUILD_PROGRAM_FAILURE); + CASE(CL_MAP_FAILURE); + CASE(CL_INVALID_VALUE); + CASE(CL_INVALID_DEVICE_TYPE); + CASE(CL_INVALID_PLATFORM); + CASE(CL_INVALID_DEVICE); + CASE(CL_INVALID_CONTEXT); + CASE(CL_INVALID_QUEUE_PROPERTIES); + CASE(CL_INVALID_COMMAND_QUEUE); + CASE(CL_INVALID_HOST_PTR); + CASE(CL_INVALID_MEM_OBJECT); + CASE(CL_INVALID_IMAGE_FORMAT_DESCRIPTOR); + CASE(CL_INVALID_IMAGE_SIZE); + CASE(CL_INVALID_SAMPLER); + CASE(CL_INVALID_BINARY); + CASE(CL_INVALID_BUILD_OPTIONS); + CASE(CL_INVALID_PROGRAM); + CASE(CL_INVALID_PROGRAM_EXECUTABLE); + CASE(CL_INVALID_KERNEL_NAME); + CASE(CL_INVALID_KERNEL_DEFINITION); + CASE(CL_INVALID_KERNEL); + CASE(CL_INVALID_ARG_INDEX); + CASE(CL_INVALID_ARG_VALUE); + CASE(CL_INVALID_ARG_SIZE); + CASE(CL_INVALID_KERNEL_ARGS); + CASE(CL_INVALID_WORK_DIMENSION); + CASE(CL_INVALID_WORK_GROUP_SIZE); + CASE(CL_INVALID_WORK_ITEM_SIZE); + CASE(CL_INVALID_GLOBAL_OFFSET); + CASE(CL_INVALID_EVENT_WAIT_LIST); + CASE(CL_INVALID_EVENT); + CASE(CL_INVALID_OPERATION); + CASE(CL_INVALID_GL_OBJECT); + CASE(CL_INVALID_BUFFER_SIZE); + CASE(CL_INVALID_MIP_LEVEL); + CASE(CL_INVALID_GLOBAL_WORK_SIZE); + default: + return "Unknown OpenCL error code"; } #undef CASE } public: - OpenCLError(cl_int err, std::string fn, int ln): mError(err), - mFile(fn), mLineNumber(ln) + OpenCLError( cl_int err, const std::string& fn, int ln ) : mError(err), + mFile(fn), mLineNumber(ln) { SAL_INFO("sc.opencl", "OpenCLError:" << mError << ": " << strerror(mError)); } @@ -117,8 +118,8 @@ public: class Unhandled { public: - Unhandled(std::string fn="", int ln=0): - mFile(fn), mLineNumber(ln) {} + Unhandled( const std::string& fn = "", int ln = 0 ) : + mFile(fn), mLineNumber(ln) { } std::string mFile; int mLineNumber; }; @@ -128,15 +129,16 @@ typedef boost::shared_ptr FormulaTreeNodeRef; class FormulaTreeNode { public: - FormulaTreeNode(const formula::FormulaToken* ft): mpCurrentFormula(ft) + FormulaTreeNode( const formula::FormulaToken* ft ) : mpCurrentFormula(ft) { Children.reserve(8); } std::vector Children; - formula::FormulaToken *GetFormulaToken(void) const + formula::FormulaToken* GetFormulaToken( void ) const { return const_cast(mpCurrentFormula.get()); } + private: formula::FormulaConstTokenRef mpCurrentFormula; }; @@ -145,49 +147,49 @@ private: class DynamicKernelArgument : boost::noncopyable { public: - DynamicKernelArgument(const std::string &s, FormulaTreeNodeRef ft); + DynamicKernelArgument( const std::string& s, FormulaTreeNodeRef ft ); - const std::string &GetNameAsString(void) const { return mSymName; } + const std::string& GetNameAsString( void ) const { return mSymName; } /// Generate declaration - virtual void GenDecl(std::stringstream &ss) const = 0; + virtual void GenDecl( std::stringstream& ss ) const = 0; /// When declared as input to a sliding window function - virtual void GenSlidingWindowDecl(std::stringstream &ss) const = 0; + virtual void GenSlidingWindowDecl( std::stringstream& ss ) const = 0; /// When referenced in a sliding window function - virtual std::string GenSlidingWindowDeclRef(bool=false) const = 0; + virtual std::string GenSlidingWindowDeclRef( bool = false ) const = 0; /// When Mix, it will be called - virtual std::string GenDoubleSlidingWindowDeclRef(bool=false) const + virtual std::string GenDoubleSlidingWindowDeclRef( bool = false ) const { return std::string(""); } /// When Mix, it will be called - virtual std::string GenStringSlidingWindowDeclRef(bool=false) const + virtual std::string GenStringSlidingWindowDeclRef( bool = false ) const { return std::string(""); } virtual bool IsMixedArgument() const { return false; } /// Generate use/references to the argument - virtual void GenDeclRef(std::stringstream &ss) const; - virtual void GenNumDeclRef(std::stringstream &ss) const{ss << ",";} + virtual void GenDeclRef( std::stringstream& ss ) const; + virtual void GenNumDeclRef( std::stringstream& ss ) const { ss << ",";} - virtual void GenStringDeclRef(std::stringstream &ss) const{ss << ",";} + virtual void GenStringDeclRef( std::stringstream& ss ) const { ss << ",";} /// Create buffer and pass the buffer to a given kernel - virtual size_t Marshal(cl_kernel, int, int, cl_program) = 0; + virtual size_t Marshal( cl_kernel, int, int, cl_program ) = 0; - virtual ~DynamicKernelArgument() {} + virtual ~DynamicKernelArgument() { } - virtual void GenSlidingWindowFunction(std::stringstream &) {} - const std::string &GetSymName(void) const { return mSymName; } - formula::FormulaToken *GetFormulaToken(void) const; - virtual size_t GetWindowSize(void) const = 0; - virtual std::string DumpOpName(void) const { return std::string(""); } - virtual void DumpInlineFun(std::set& , - std::set& ) const {} - const std::string& GetName(void) const { return mSymName; } - virtual bool NeedParallelReduction(void) const { return false; } + virtual void GenSlidingWindowFunction( std::stringstream& ) { } + const std::string& GetSymName( void ) const { return mSymName; } + formula::FormulaToken* GetFormulaToken( void ) const; + virtual size_t GetWindowSize( void ) const = 0; + virtual std::string DumpOpName( void ) const { return std::string(""); } + virtual void DumpInlineFun( std::set&, + std::set& ) const { } + const std::string& GetName( void ) const { return mSymName; } + virtual bool NeedParallelReduction( void ) const { return false; } protected: std::string mSymName; @@ -202,31 +204,31 @@ protected: class VectorRef : public DynamicKernelArgument { public: - VectorRef(const std::string &s, FormulaTreeNodeRef ft, int index = 0); + VectorRef( const std::string& s, FormulaTreeNodeRef ft, int index = 0 ); - const std::string &GetNameAsString(void) const { return mSymName; } + const std::string& GetNameAsString( void ) const { return mSymName; } /// Generate declaration - virtual void GenDecl(std::stringstream &ss) const SAL_OVERRIDE; + virtual void GenDecl( std::stringstream& ss ) const SAL_OVERRIDE; /// When declared as input to a sliding window function - virtual void GenSlidingWindowDecl(std::stringstream &ss) const SAL_OVERRIDE; + virtual void GenSlidingWindowDecl( std::stringstream& ss ) const SAL_OVERRIDE; /// When referenced in a sliding window function - virtual std::string GenSlidingWindowDeclRef(bool=false) const SAL_OVERRIDE; + virtual std::string GenSlidingWindowDeclRef( bool = false ) const SAL_OVERRIDE; /// Create buffer and pass the buffer to a given kernel - virtual size_t Marshal(cl_kernel, int, int, cl_program) SAL_OVERRIDE; + virtual size_t Marshal( cl_kernel, int, int, cl_program ) SAL_OVERRIDE; virtual ~VectorRef(); - virtual void GenSlidingWindowFunction(std::stringstream &) SAL_OVERRIDE {} - const std::string &GetSymName(void) const { return mSymName; } - virtual size_t GetWindowSize(void) const SAL_OVERRIDE; - virtual std::string DumpOpName(void) const SAL_OVERRIDE { return std::string(""); } - virtual void DumpInlineFun(std::set& , - std::set& ) const SAL_OVERRIDE {} - const std::string& GetName(void) const { return mSymName; } - virtual cl_mem GetCLBuffer(void) const { return mpClmem; } - virtual bool NeedParallelReduction(void) const SAL_OVERRIDE { return false; } + virtual void GenSlidingWindowFunction( std::stringstream& ) SAL_OVERRIDE { } + const std::string& GetSymName( void ) const { return mSymName; } + virtual size_t GetWindowSize( void ) const SAL_OVERRIDE; + virtual std::string DumpOpName( void ) const SAL_OVERRIDE { return std::string(""); } + virtual void DumpInlineFun( std::set&, + std::set& ) const SAL_OVERRIDE { } + const std::string& GetName( void ) const { return mSymName; } + virtual cl_mem GetCLBuffer( void ) const { return mpClmem; } + virtual bool NeedParallelReduction( void ) const SAL_OVERRIDE { return false; } protected: // Used by marshaling @@ -234,26 +236,26 @@ protected: // index in multiple double vector refs that have multiple ranges const int mnIndex; }; -/// Abstract class for code generation +/// Abstract class for code generation class OpBase { public: typedef std::vector ArgVector; typedef std::vector::iterator ArgVectorIter; - virtual std::string GetBottom(void) {return "";}; - virtual std::string Gen2(const std::string &/*lhs*/, - const std::string &/*rhs*/) const {return "";} - virtual std::string Gen(ArgVector& /*argVector*/){return "";}; - virtual std::string BinFuncName(void)const {return "";}; - virtual void BinInlineFun(std::set& , - std::set& ) {} + virtual std::string GetBottom( void ) { return "";}; + virtual std::string Gen2( const std::string&/*lhs*/, + const std::string&/*rhs*/ ) const { return "";} + virtual std::string Gen( ArgVector& /*argVector*/ ) { return "";}; + virtual std::string BinFuncName( void ) const { return "";}; + virtual void BinInlineFun( std::set&, + std::set& ) { } virtual bool takeString() const = 0; virtual bool takeNumeric() const = 0; //Continue process 'Zero' or Not(like OpMul, not continue process when meet // 'Zero' - virtual bool ZeroReturnZero() {return false;} - virtual ~OpBase() {} + virtual bool ZeroReturnZero() { return false;} + virtual ~OpBase() { } }; class SlidingFunctionBase : public OpBase @@ -261,34 +263,34 @@ class SlidingFunctionBase : public OpBase public: typedef boost::shared_ptr SubArgument; typedef std::vector SubArguments; - virtual void GenSlidingWindowFunction(std::stringstream &, - const std::string&, SubArguments &) = 0; - virtual ~SlidingFunctionBase() {}; + virtual void GenSlidingWindowFunction( std::stringstream&, + const std::string&, SubArguments& ) = 0; + virtual ~SlidingFunctionBase() { }; }; -class Normal: public SlidingFunctionBase +class Normal : public SlidingFunctionBase { public: - virtual void GenSlidingWindowFunction(std::stringstream &ss, - const std::string &sSymName, SubArguments &vSubArguments) SAL_OVERRIDE; + virtual void GenSlidingWindowFunction( std::stringstream& ss, + const std::string& sSymName, SubArguments& vSubArguments ) SAL_OVERRIDE; virtual bool takeString() const SAL_OVERRIDE { return false; } virtual bool takeNumeric() const SAL_OVERRIDE { return true; } }; -class CheckVariables:public Normal +class CheckVariables : public Normal { public: - void GenTmpVariables(std::stringstream &ss, SubArguments &vSubArguments); - void CheckSubArgumentIsNan(std::stringstream &ss, - SubArguments &vSubArguments, int argumentNum); - void CheckAllSubArgumentIsNan(std::stringstream &ss, - SubArguments &vSubArguments); + void GenTmpVariables( std::stringstream& ss, SubArguments& vSubArguments ); + void CheckSubArgumentIsNan( std::stringstream& ss, + SubArguments& vSubArguments, int argumentNum ); + void CheckAllSubArgumentIsNan( std::stringstream& ss, + SubArguments& vSubArguments ); // only check isNan - void CheckSubArgumentIsNan2(std::stringstream &ss, - SubArguments &vSubArguments, int argumentNum, std::string p); - void UnrollDoubleVector(std::stringstream &ss, - std::stringstream &unrollstr, const formula::DoubleVectorRefToken* pCurDVR, - int nCurWindowSize); + void CheckSubArgumentIsNan2( std::stringstream& ss, + SubArguments& vSubArguments, int argumentNum, std::string p ); + void UnrollDoubleVector( std::stringstream& ss, + std::stringstream& unrollstr, const formula::DoubleVectorRefToken* pCurDVR, + int nCurWindowSize ); }; }} -- cgit