Skip to content

C
core
Kaydet (Commit) 2d73c5cd authored tarafından haochen's avatar haochen Kaydeden (comit) Markus Mohrhard
Dosyalara gözat
Seçenekler

GPU Calc:Fixed style&error problem 

in op_financial&opbase.cxx&opencl_device.cxx

Change-Id: Ic4b8cee77a0a4faec44bcf3650f4f6fa691a9e71
üst fbbf5c08
Hide whitespace changes
Inline Side-by-side
Showing with 38 additions and 56 deletions
sc/source/core/opencl/op_financial.cxx
Dosyayı görüntüle @ 2d73c5cd
......@@ -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();
......
sc/source/core/opencl/opbase.cxx
Dosyayı görüntüle @ 2d73c5cd
......@@ -43,8 +43,6 @@ VectorRef::~VectorRef()
{
if (mpClmem) {
cl_int ret = clReleaseMemObject(mpClmem);
if (ret != CL_SUCCESS)
throw OpenCLError(ret, __FILE__, __LINE__);
}
}
......
sc/source/core/opencl/opencl_device.cxx
Dosyayı görüntüle @ 2d73c5cd
......@@ -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;
}
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment