summaryrefslogtreecommitdiff
path: root/sc/source/core/opencl
diff options
context:
space:
mode:
authorhaochen <haochen@multicorewareinc.com>2014-05-16 17:52:55 +0800
committerMarkus Mohrhard <markus.mohrhard@googlemail.com>2014-05-19 09:13:23 +0200
commit2d73c5cdefd4fe811ae33938be32772c84fee16c (patch)
treeb52cbb45452d0de2a912495e39183c4edd89a411 /sc/source/core/opencl
parentfbbf5c089d9e102f819df752dce0e1b3223ad269 (diff)
GPU Calc:Fixed style&error problem
in op_financial&opbase.cxx&opencl_device.cxx Change-Id: Ic4b8cee77a0a4faec44bcf3650f4f6fa691a9e71
Diffstat (limited to 'sc/source/core/opencl')
-rw-r--r--sc/source/core/opencl/op_financial.cxx74
-rw-r--r--sc/source/core/opencl/opbase.cxx2
-rw-r--r--sc/source/core/opencl/opencl_device.cxx18
3 files changed, 38 insertions, 56 deletions
diff --git a/sc/source/core/opencl/op_financial.cxx b/sc/source/core/opencl/op_financial.cxx
index ef669e74dd0d..6170bb77df7d 100644
--- a/sc/source/core/opencl/op_financial.cxx
+++ b/sc/source/core/opencl/op_financial.cxx
@@ -1253,14 +1253,10 @@ void XNPV::GenSlidingWindowFunction(
ss << "double result = 0.0;\n\t";
ss << "int gid0 = get_global_id(0);\n\t";
ss << "int i=0;\n\t";
- if (!pCurDVR->IsStartFixed() && pCurDVR->IsEndFixed()) {
- ss<< "i=gid0;\n\t";
- }
ss << "double date;\n\t";
ss << "double value;\n\t";
ss << "double rate;\n\t";
ss << "double dateNull;\n\t";
-#ifdef ISNAN
FormulaToken *tmpCur0 = vSubArguments[0]->GetFormulaToken();
const formula::SingleVectorRefToken*tmpCurDVR0= static_cast<const
formula::SingleVectorRefToken *>(tmpCur0);
@@ -1281,25 +1277,24 @@ void XNPV::GenSlidingWindowFunction(
ss<< "int buffer_date_len = ";
ss<< tmpCurDVR2->GetArrayLength();
ss << ";\n\t";
-#endif
-#ifdef ISNAN
- ss<<"if((i+gid0)>=buffer_date_len || isNan(";
+ ss<<"if((gid0)>=buffer_date_len || isNan(";
ss << vSubArguments[2]->GenSlidingWindowDeclRef();
ss<<"))\n\t\t";
- ss<<"dateNull = 0;\n\telse \n\t\t";
-#endif
+ ss<<"return NAN;\n\telse \n";
ss<<"dateNull = ";
ss << vSubArguments[2]->GenSlidingWindowDeclRef();
ss<<";\n\t";
-#ifdef ISNAN
- ss<<"if((i+gid0)>=buffer_rate_len || isNan(";
+ ss<<"if((gid0)>=buffer_rate_len || isNan(";
ss << vSubArguments[0]->GenSlidingWindowDeclRef();
ss<<"))\n\t\t";
- ss<<"rate = 0;\n\telse \n\t\t";
-#endif
+ ss<<"return NAN;\n\telse \n";
ss<<"rate = ";
ss << vSubArguments[0]->GenSlidingWindowDeclRef();
ss<<";\n\t";
+ ss<<"if(1 == buffer_date_len )\n";
+ ss<<"return ";
+ ss << vSubArguments[1]->GenSlidingWindowDeclRef();
+ ss<<";\n\t";
ss << "for (int i = ";
if (!pCurDVR->IsStartFixed() && pCurDVR->IsEndFixed())
{
@@ -1314,17 +1309,27 @@ void XNPV::GenSlidingWindowFunction(
ss << "0; i < "<< nCurWindowSize <<"; i++)\n\t\t";
}
ss << "{\n\t";
- ss << "result += ";
- ss << vSubArguments[1]->GenSlidingWindowDeclRef();
- ss << "/(pow((";
- ss<<vSubArguments[0]->GenSlidingWindowDeclRef();
- ss <<"+1),(";
- ss << vSubArguments[2]->GenSlidingWindowDeclRef();
- ss << "-dateNull)/365));\n\t";
- ss <<"}\n\t";
+ if (!pCurDVR->IsStartFixed() && !pCurDVR->IsEndFixed())
+ {
+ ss << "if((i+gid0)>=buffer_value_len || (i+gid0)>=buffer_date_len)\n\t\t";
+ ss << "return result;\n\telse \n\t\t";
+ }
+ else
+ {
+ ss << "if(i>=buffer_value_len || i>=buffer_date_len)\n\t\t";
+ ss << "return result;\n\telse \n\t\t";
+ }
- ss << "return result;\n";
- ss << "}";
+ ss << "value = ";
+ ss << vSubArguments[1]->GenSlidingWindowDeclRef(true);
+ ss << ";\n";
+ ss << " date = ";
+ ss << vSubArguments[2]->GenSlidingWindowDeclRef(true);
+ ss << ";\n";
+ ss << "result += value/(pow((rate+1),(date-dateNull)/365));\n";
+ ss << "}\n";
+ ss << "return result;\n";
+ ss << "}";
}
void PriceMat::BinInlineFun(std::set<std::string>& decls,
@@ -2891,6 +2896,7 @@ void OpPrice::GenSlidingWindowFunction(std::stringstream &ss,
ss <<";\n";
#endif
}
+ ss << " if(tmp4*tmp5 == 0) return NAN;\n";
ss << " tmp = getPrice_(tmp0,tmp1,tmp2,tmp3,tmp4,tmp5,tmp6);\n";
ss << " return tmp;\n";
ss << "}";
@@ -2932,7 +2938,6 @@ void OpOddlprice::GenSlidingWindowFunction(std::stringstream &ss,
ss <<" double tmp5=0;\n";
ss <<" double tmp6=0;\n";
ss <<" double tmp7=0;\n";
- size_t nItems = 0;
ss <<" \n";
for (size_t i = 0; i < vSubArguments.size(); i++)
{
@@ -2973,7 +2978,6 @@ void OpOddlprice::GenSlidingWindowFunction(std::stringstream &ss,
ss << "0; i < "<< nCurWindowSize << "; i++)\n";
#endif
}
- nItems += nCurWindowSize;
}
else if (pCur->GetType() == formula::svSingleVectorRef)
{
@@ -2982,7 +2986,6 @@ void OpOddlprice::GenSlidingWindowFunction(std::stringstream &ss,
static_cast< const formula::SingleVectorRefToken* >(pCur);
ss << " if (gid0 < " << pSVR->GetArrayLength() << "){\n";
#else
- nItems += 1;
#endif
}
else if (pCur->GetType() == formula::svDouble)
@@ -2990,13 +2993,11 @@ void OpOddlprice::GenSlidingWindowFunction(std::stringstream &ss,
#ifdef ISNAN
ss << "{\n";
#endif
- nItems += 1;
}
else
{
#ifdef ISNAN
#endif
- nItems += 1;
}
#ifdef ISNAN
if(ocPush==vSubArguments[i]->GetFormulaToken()->GetOpCode())
@@ -3064,7 +3065,6 @@ void OpOddlyield::GenSlidingWindowFunction(std::stringstream &ss,
ss <<" double tmp5=0;\n";
ss <<" double tmp6=0;\n";
ss <<" double tmp7=0;\n";
- size_t nItems = 0;
ss <<" \n";
for (size_t i = 0; i < vSubArguments.size(); i++)
{
@@ -3105,7 +3105,6 @@ void OpOddlyield::GenSlidingWindowFunction(std::stringstream &ss,
ss << "0; i < "<< nCurWindowSize << "; i++)\n";
#endif
}
- nItems += nCurWindowSize;
}
else if (pCur->GetType() == formula::svSingleVectorRef)
{
@@ -3114,7 +3113,6 @@ void OpOddlyield::GenSlidingWindowFunction(std::stringstream &ss,
static_cast< const formula::SingleVectorRefToken* >(pCur);
ss << " if (gid0 < " << pSVR->GetArrayLength() << "){\n";
#else
- nItems += 1;
#endif
}
else if (pCur->GetType() == formula::svDouble)
@@ -3122,13 +3120,11 @@ void OpOddlyield::GenSlidingWindowFunction(std::stringstream &ss,
#ifdef ISNAN
ss << "{\n";
#endif
- nItems += 1;
}
else
{
#ifdef ISNAN
#endif
- nItems += 1;
}
#ifdef ISNAN
if(ocPush==vSubArguments[i]->GetFormulaToken()->GetOpCode())
@@ -3271,7 +3267,7 @@ void OpNper::GenSlidingWindowFunction(std::stringstream &ss,
ss <<" double tmp2=0;\n";
ss <<" double tmp3=0;\n";
ss <<" double tmp4=0;\n";
- size_t nItems = 0;
+
for (size_t i = 0; i < vSubArguments.size(); i++)
{
FormulaToken *pCur = vSubArguments[i]->GetFormulaToken();
@@ -3283,7 +3279,6 @@ void OpNper::GenSlidingWindowFunction(std::stringstream &ss,
static_cast< const formula::SingleVectorRefToken* >(pCur);
ss << " if (gid0 < " << pSVR->GetArrayLength() << "){\n";
#else
- nItems += 1;
#endif
}
else if (pCur->GetType() == formula::svDouble)
@@ -3291,13 +3286,11 @@ void OpNper::GenSlidingWindowFunction(std::stringstream &ss,
#ifdef ISNAN
ss << "{\n";
#endif
- nItems += 1;
}
else
{
#ifdef ISNAN
#endif
- nItems += 1;
}
#ifdef ISNAN
if(ocPush==vSubArguments[i]->GetFormulaToken()->GetOpCode())
@@ -5010,13 +5003,10 @@ vSubArguments)
const formula::SingleVectorRefToken*tmpCurDVR2= static_cast<const
formula::SingleVectorRefToken *>(tmpCur2);
- const formula::SingleVectorRefToken*tmpCurDVR3;
- const formula::SingleVectorRefToken*tmpCurDVR4;
-
if(vSubArguments.size()>3)
{
FormulaToken *tmpCur3 = vSubArguments[3]->GetFormulaToken();
- tmpCurDVR3= static_cast<const formula::SingleVectorRefToken *>(
+ const formula::SingleVectorRefToken* tmpCurDVR3= static_cast<const formula::SingleVectorRefToken *>(
tmpCur3);
ss<< " int buffer_zw_len = ";
ss<< tmpCurDVR3->GetArrayLength();
@@ -5026,7 +5016,7 @@ tmpCur3);
if(vSubArguments.size()>4)
{
FormulaToken *tmpCur4 = vSubArguments[4]->GetFormulaToken();
- tmpCurDVR4= static_cast<const formula::SingleVectorRefToken *>(
+ const formula::SingleVectorRefToken* tmpCurDVR4= static_cast<const formula::SingleVectorRefToken *>(
tmpCur4);
ss<< " int buffer_flag_len = ";
ss<< tmpCurDVR4->GetArrayLength();
diff --git a/sc/source/core/opencl/opbase.cxx b/sc/source/core/opencl/opbase.cxx
index 63c07b9b04b5..6bc5214f809a 100644
--- a/sc/source/core/opencl/opbase.cxx
+++ b/sc/source/core/opencl/opbase.cxx
@@ -43,8 +43,6 @@ VectorRef::~VectorRef()
{
if (mpClmem) {
cl_int ret = clReleaseMemObject(mpClmem);
- if (ret != CL_SUCCESS)
- throw OpenCLError(ret, __FILE__, __LINE__);
}
}
diff --git a/sc/source/core/opencl/opencl_device.cxx b/sc/source/core/opencl/opencl_device.cxx
index 1a5edc545bec..9070cd8ae519 100644
--- a/sc/source/core/opencl/opencl_device.cxx
+++ b/sc/source/core/opencl/opencl_device.cxx
@@ -226,11 +226,6 @@ ds_status evaluateScoreForDevice(ds_device* device, void* evalData)
/* Evaluating an OpenCL device */
LOG_PRINTF("[DS] Device: \"" << device->oclDeviceName << "\" (OpenCL) evaluation...");
cl_int clStatus;
- cl_context clContext;
- cl_command_queue clQueue;
- cl_program clProgram;
- cl_kernel clKernel;
-
/* Check for 64-bit float extensions */
size_t aDevExtInfoSize = 0;
clStatus = clGetDeviceInfo(device->oclDeviceID, CL_DEVICE_EXTENSIONS, 0, NULL, &aDevExtInfoSize);
@@ -239,7 +234,6 @@ ds_status evaluateScoreForDevice(ds_device* device, void* evalData)
char* aExtInfo = new char[aDevExtInfoSize];
clStatus = clGetDeviceInfo(device->oclDeviceID, CL_DEVICE_EXTENSIONS, sizeof(char) * aDevExtInfoSize, aExtInfo, NULL);
DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clGetDeviceInfo");
-
bool bKhrFp64Flag = false;
bool bAmdFp64Flag = false;
const char* buildOption = NULL;
@@ -276,13 +270,13 @@ ds_status evaluateScoreForDevice(ds_device* device, void* evalData)
/* 64-bit float support present */
/* Create context and command queue */
- clContext = clCreateContext(NULL, 1, &device->oclDeviceID, NULL, NULL, &clStatus);
+ cl_context clContext = clCreateContext(NULL, 1, &device->oclDeviceID, NULL, NULL, &clStatus);
DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clCreateContext");
- clQueue = clCreateCommandQueue(clContext, device->oclDeviceID, 0, &clStatus);
+ cl_command_queue clQueue = clCreateCommandQueue(clContext, device->oclDeviceID, 0, &clStatus);
DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clCreateCommandQueue");
/* Build program */
- clProgram = clCreateProgramWithSource(clContext, 1, &source, sourceSize, &clStatus);
+ cl_program clProgram = clCreateProgramWithSource(clContext, 1, &source, sourceSize, &clStatus);
DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clCreateProgramWithSource");
clStatus = clBuildProgram(clProgram, 1, &device->oclDeviceID, buildOption, NULL, NULL);
DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clBuildProgram");
@@ -293,7 +287,7 @@ ds_status evaluateScoreForDevice(ds_device* device, void* evalData)
char* buildLog;
clStatus = clGetProgramBuildInfo(clProgram, device->oclDeviceID, CL_PROGRAM_BUILD_LOG, 0, NULL, &length);
buildLog = (char*)malloc(length);
- clStatus = clGetProgramBuildInfo(clProgram, device->oclDeviceID, CL_PROGRAM_BUILD_LOG, length, buildLog, &length);
+ clGetProgramBuildInfo(clProgram, device->oclDeviceID, CL_PROGRAM_BUILD_LOG, length, buildLog, &length);
LOG_PRINTF("[OCL] Build Errors" << std::endl << buildLog);
free(buildLog);
@@ -309,7 +303,7 @@ ds_status evaluateScoreForDevice(ds_device* device, void* evalData)
/* Run kernel */
LibreOfficeDeviceEvaluationIO* testData = (LibreOfficeDeviceEvaluationIO*)evalData;
- clKernel = clCreateKernel(clProgram, "DynamicKernel", &clStatus);
+ cl_kernel clKernel = clCreateKernel(clProgram, "DynamicKernel", &clStatus);
DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clCreateKernel");
cl_mem clResult = clCreateBuffer(clContext, CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR, sizeof(cl_double) * testData->outputSize, &testData->output[0], &clStatus);
DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clCreateBuffer::clResult");
@@ -560,7 +554,7 @@ ds_device getDeviceSelection(const char* sProfilePath, bool bForceSelection)
bIsDeviceSelected = true;
/* Release profile */
- status = releaseDSProfile(profile, releaseScore);
+ releaseDSProfile(profile, releaseScore);
}
return selectedDevice;
}