summaryrefslogtreecommitdiff
path: root/sc/source
diff options
context:
space:
mode:
authorKohei Yoshida <kohei.yoshida@collabora.com>2014-10-03 12:01:51 -0400
committerKohei Yoshida <kohei.yoshida@collabora.com>2014-10-03 15:31:25 -0400
commit2549b21b5226c30d900f5994d45b7b08c15935d6 (patch)
tree1ebbe39ae5dc307b6afcb3ea0d7e353bcc878f9f /sc/source
parent08d93506c21fc4899392074f853ae3175fb86aee (diff)
Run 'beautify' on these files as well.
And some trivial manual fixing afterward... Change-Id: Ib60411f689b7971a2dd1acd1541864f38b34b570
Diffstat (limited to 'sc/source')
-rw-r--r--sc/source/core/opencl/opbase.cxx318
-rw-r--r--sc/source/core/opencl/opbase.hxx248
2 files changed, 292 insertions, 274 deletions
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 *"<<mSymName;
+ ss << "__global double *" << mSymName;
}
/// When declared as input to a sliding window function
-void VectorRef::GenSlidingWindowDecl(std::stringstream &ss) const
+void VectorRef::GenSlidingWindowDecl( std::stringstream& ss ) const
{
VectorRef::GenDecl(ss);
}
/// When referenced in a sliding window function
-std::string VectorRef::GenSlidingWindowDeclRef(bool nested) const
+std::string VectorRef::GenSlidingWindowDeclRef( bool nested ) const
{
std::stringstream ss;
- formula::SingleVectorRefToken *pSVR =
+ formula::SingleVectorRefToken* pSVR =
dynamic_cast<formula::SingleVectorRefToken*>(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<const formula::DoubleVectorRefToken *>(pCur))
+ dynamic_cast<const formula::DoubleVectorRefToken*>(pCur))
{
return pCurDVR->GetRefRowSize();
}
- else if (dynamic_cast<const formula::SingleVectorRefToken *>(pCur))
+ else if (dynamic_cast<const formula::SingleVectorRefToken*>(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;i<vSubArguments.size();i++)
+ for (unsigned i = 0; i < vSubArguments.size(); i++)
{
- ss << " double tmp";
- ss << i;
- ss <<";\n";
+ ss << " double tmp";
+ ss << i;
+ ss << ";\n";
}
}
-void CheckVariables::CheckSubArgumentIsNan( std::stringstream & ss,
- SubArguments &vSubArguments, int argumentNum)
+void CheckVariables::CheckSubArgumentIsNan( std::stringstream& ss,
+ SubArguments& vSubArguments, int argumentNum )
{
int i = argumentNum;
- if(vSubArguments[i]->GetFormulaToken()->GetType() ==
- formula::svSingleVectorRef)
- {
- const formula::SingleVectorRefToken*pTmpDVR1= static_cast<const
- formula::SingleVectorRefToken *>(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<const
- formula::DoubleVectorRefToken *>(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";
+ if (vSubArguments[i]->GetFormulaToken()->GetType() ==
+ formula::svSingleVectorRef)
+ {
+ const formula::SingleVectorRefToken* pTmpDVR1 =
+ static_cast<const formula::SingleVectorRefToken*>(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<const formula::DoubleVectorRefToken*>(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 <<i;
+ ss << " tmp";
+ ss << i;
ss << "=";
vSubArguments[i]->GenDeclRef(ss);
- ss<<";\n";
+ ss << ";\n";
return;
}
#ifdef ISNAN
- ss<< " tmp";
- ss<< i;
- ss<< "= fsum(";
+ ss << " tmp";
+ ss << i;
+ ss << "= fsum(";
vSubArguments[i]->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 <<i;
+ ss << " tmp";
+ ss << i;
ss << "=";
vSubArguments[i]->GenDeclRef(ss);
- if(vSubArguments[i]->GetFormulaToken()->GetType() ==
+ if (vSubArguments[i]->GetFormulaToken()->GetType() ==
formula::svDoubleVectorRef)
- ss<<"["<< p.c_str()<< "]";
- else if(vSubArguments[i]->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;i<vSubArguments.size();i++)
+ ss << " int k = gid0;\n";
+ for (unsigned i = 0; i < vSubArguments.size(); i++)
{
- CheckSubArgumentIsNan(ss,vSubArguments,i);
+ CheckSubArgumentIsNan(ss, vSubArguments, i);
}
}
-void CheckVariables::UnrollDoubleVector( std::stringstream & ss,
-std::stringstream & unrollstr,const formula::DoubleVectorRefToken* pCurDVR,
-int nCurWindowSize)
+void CheckVariables::UnrollDoubleVector( std::stringstream& ss,
+ std::stringstream& unrollstr, const formula::DoubleVectorRefToken* pCurDVR,
+ int nCurWindowSize )
{
int unrollSize = 16;
- if (!pCurDVR->IsStartFixed() && pCurDVR->IsEndFixed()) {
- ss << " loop = ("<<nCurWindowSize<<" - gid0)/";
- ss << unrollSize<<";\n";
- } else if (pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed()) {
- ss << " loop = ("<<nCurWindowSize<<" + gid0)/";
- ss << unrollSize<<";\n";
+ if (!pCurDVR->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 = "<<nCurWindowSize<<"/"<< unrollSize<<";\n";
+ }
+ else
+ {
+ ss << " loop = " << nCurWindowSize << "/" << unrollSize << ";\n";
}
ss << " for ( int j = 0;j< loop; j++)\n";
ss << " {\n";
ss << " int i = ";
- if (!pCurDVR->IsStartFixed()&& pCurDVR->IsEndFixed()) {
- ss << "gid0 + j * "<< unrollSize <<";\n";
- }else {
- ss << "j * "<< unrollSize <<";\n";
+ if (!pCurDVR->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 *"<<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 << " }\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<FormulaTreeNode> FormulaTreeNodeRef;
class FormulaTreeNode
{
public:
- FormulaTreeNode(const formula::FormulaToken* ft): mpCurrentFormula(ft)
+ FormulaTreeNode( const formula::FormulaToken* ft ) : mpCurrentFormula(ft)
{
Children.reserve(8);
}
std::vector<FormulaTreeNodeRef> Children;
- formula::FormulaToken *GetFormulaToken(void) const
+ formula::FormulaToken* GetFormulaToken( void ) const
{
return const_cast<formula::FormulaToken*>(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::string>& ,
- std::set<std::string>& ) 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::string>&,
+ std::set<std::string>& ) 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::string>& ,
- std::set<std::string>& ) 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::string>&,
+ std::set<std::string>& ) 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<std::string> ArgVector;
typedef std::vector<std::string>::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::string>& ,
- std::set<std::string>& ) {}
+ 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::string>&,
+ std::set<std::string>& ) { }
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<DynamicKernelArgument> SubArgument;
typedef std::vector<SubArgument> 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 );
};
}}