Attachment 113306 Details for Bug 88444 – Patch to add even more debugging output

[patch] Patch to add even more debugging output

4-2-opencl-debug.diff (text/plain), 14.73 KB, created by How can I remove my account? on 2015-02-11 10:58:28 UTC

(hide)

Description:

Filename:

MIME Type:

Creator: How can I remove my account?

Created: 2015-02-11 10:58:28 UTC

Size: 14.73 KB

patch

obsolete

>diff --git a/sc/source/core/opencl/formulagroupcl.cxx b/sc/source/core/opencl/formulagroupcl.cxx
>index 724f675..6e4342a 100644
>--- a/sc/source/core/opencl/formulagroupcl.cxx
>+++ b/sc/source/core/opencl/formulagroupcl.cxx
>@@ -98,6 +98,7 @@ size_t VectorRef::Marshal(cl_kernel k, int argno, int, cl_program)
>                 pHostBuffer, &err);
>         if (CL_SUCCESS != err)
>             throw OpenCLError(err);
>+        SAL_INFO("sc.opencl", "Created buffer " << mpClmem << " size " << szHostBuffer << " from host buffer " << pHostBuffer);
>     }
>     else
>     {
>@@ -114,12 +115,14 @@ size_t VectorRef::Marshal(cl_kernel k, int argno, int, cl_program)
>                 szHostBuffer, 0, NULL, NULL, &err);
>         if (CL_SUCCESS != err)
>             throw OpenCLError(err);
>+        SAL_INFO("sc.opencl", "Created buffer " << mpClmem << " size " << szHostBuffer << ", mapped to " << pNanBuffer);
>         for (size_t i = 0; i < szHostBuffer/sizeof(double); i++)
>             pNanBuffer[i] = NAN;
>         err = clEnqueueUnmapMemObject(kEnv.mpkCmdQueue, mpClmem,
>                 pNanBuffer, 0, NULL, NULL);
>     }
> 
>+    SAL_INFO("sc.opencl", "Kernel " << k << " arg " << argno << ": cl_mem: " << mpClmem);
>     err = clSetKernelArg(k, argno, sizeof(cl_mem), (void*)&mpClmem);
>     if (CL_SUCCESS != err)
>         throw OpenCLError(err);
>@@ -179,6 +182,7 @@ public:
>         KernelEnv kEnv;
>         OpenclDevice::setKernelEnv(&kEnv);
>         // Pass the scalar result back to the rest of the formula kernel
>+        SAL_INFO("sc.opencl", "Kernel " << k << " arg " << argno << ": cl_uint: " << hashCode);
>         cl_int err = clSetKernelArg(k, argno, sizeof(cl_uint), (void*)&hashCode);
>         if (CL_SUCCESS != err)
>             throw OpenCLError(err);
>@@ -228,6 +232,7 @@ public:
>     {
>         double tmp = GetDouble();
>         // Pass the scalar result back to the rest of the formula kernel
>+        SAL_INFO("sc.opencl", "Kernel " << k << " arg " << argno << ": double: " << tmp);
>         cl_int err = clSetKernelArg(k, argno, sizeof(double), (void*)&tmp);
>         if (CL_SUCCESS != err)
>             throw OpenCLError(err);
>@@ -268,6 +273,7 @@ public:
>     {
>         double tmp = 0.0;
>         // Pass the scalar result back to the rest of the formula kernel
>+        SAL_INFO("sc.opencl", "Kernel " << k << " arg " << argno << ": double: " << tmp);
>         cl_int err = clSetKernelArg(k, argno, sizeof(double), (void*)&tmp);
>         if (CL_SUCCESS != err)
>             throw OpenCLError(err);
>@@ -310,6 +316,7 @@ public:
>     {
>         double tmp = 0.0;
>         // Pass the scalar result back to the rest of the formula kernel
>+        SAL_INFO("sc.opencl", "Kernel " << k << " arg " << argno << ": double: " << tmp);
>         cl_int err = clSetKernelArg(k, argno, sizeof(double), (void*)&tmp);
>         if (CL_SUCCESS != err)
>             throw OpenCLError(err);
>@@ -373,6 +380,7 @@ size_t DynamicKernelStringArgument::Marshal(cl_kernel k, int argno, int, cl_prog
>             szHostBuffer, 0, NULL, NULL, &err);
>     if (CL_SUCCESS != err)
>         throw OpenCLError(err);
>+    SAL_INFO("sc.opencl", "Created buffer " << mpClmem << " size " << szHostBuffer << ", mapped to " << pHashBuffer);
>     for (size_t i = 0; i < nStrings; i++)
>     {
>         if (vRef.mpStringArray[i])
>@@ -390,6 +398,7 @@ size_t DynamicKernelStringArgument::Marshal(cl_kernel k, int argno, int, cl_prog
>     if (CL_SUCCESS != err)
>         throw OpenCLError(err);
> 
>+    SAL_INFO("sc.opencl", "Kernel " << k << " arg " << argno << ": cl_mem: " << mpClmem);
>     err = clSetKernelArg(k, argno, sizeof(cl_mem), (void*)&mpClmem);
>     if (CL_SUCCESS != err)
>         throw OpenCLError(err);
>@@ -799,32 +808,39 @@ public:
>                 (cl_mem_flags) CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR,
>                 szHostBuffer,
>                 pHostBuffer, &err);
>+        SAL_INFO("sc.opencl", "Created buffer " << Base::mpClmem << " size " << nInput << "*" << sizeof(double) << "=" << szHostBuffer << " from host buffer " << pHostBuffer);
>         mpClmem2 = clCreateBuffer(kEnv.mpkContext, CL_MEM_WRITE_ONLY,
>                 sizeof(double)*w, NULL, NULL);
>         if (CL_SUCCESS != err)
>             throw OpenCLError(err);
>+        SAL_INFO("sc.opencl", "Created buffer " << mpClmem2 << " size " << sizeof(double) << "*" << w << "=" << (sizeof(double)*w));
>         // reproduce the reduction function name
>         std::string kernelName = Base::GetName() + "_reduction";
> 
>         cl_kernel redKernel = clCreateKernel(mpProgram, kernelName.c_str(), &err);
>         if (err != CL_SUCCESS)
>             throw OpenCLError(err);
>+        SAL_INFO("sc.opencl", "Created \"reduction\" kernel " << kernelName.c_str() << ": " << redKernel);
>         // set kernel arg of reduction kernel
>         // TODO(Wei Wei): use unique name for kernel
>         cl_mem buf = Base::GetCLBuffer();
>+        SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 0 << ": cl_mem: " << buf);
>         err = clSetKernelArg(redKernel, 0, sizeof(cl_mem),
>                 (void *)&buf);
>         if (CL_SUCCESS != err)
>             throw OpenCLError(err);
> 
>+        SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 1 << ": cl_mem: " << mpClmem2);
>         err = clSetKernelArg(redKernel, 1, sizeof(cl_mem), (void *)&mpClmem2);
>         if (CL_SUCCESS != err)
>             throw OpenCLError(err);
> 
>+        SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 2 << ": cl_int: " << nInput);
>         err = clSetKernelArg(redKernel, 2, sizeof(cl_int), (void*)&nInput);
>         if (CL_SUCCESS != err)
>             throw OpenCLError(err);
> 
>+        SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 3 << ": cl_ing: " << nCurWindowSize);
>         err = clSetKernelArg(redKernel, 3, sizeof(cl_int), (void*)&nCurWindowSize);
>         if (CL_SUCCESS != err)
>             throw OpenCLError(err);
>@@ -832,6 +848,7 @@ public:
>         // set work group size and execute
>         size_t global_work_size[] = {256, (size_t)w };
>         size_t local_work_size[] = {256, 1};
>+        SAL_INFO("sc.opencl", "Enqueueing kernel " << redKernel << " work sizes: {256," << w << "}, {256,1}");
>         err = clEnqueueNDRangeKernel(kEnv.mpkCmdQueue, redKernel, 2, NULL,
>                 global_work_size, local_work_size, 0, NULL, NULL);
>         if (CL_SUCCESS != err)
>@@ -839,8 +856,10 @@ public:
>         err = clFinish(kEnv.mpkCmdQueue);
>         if (CL_SUCCESS != err)
>             throw OpenCLError(err);
>+        SAL_INFO("sc.opencl", "Kernel " << redKernel << " finished");
> 
>         // set kernel arg
>+        SAL_INFO("sc.opencl", "Kernel " << k << " arg " << argno << ": cl_mem: " << mpClmem2);
>         err = clSetKernelArg(k, argno, sizeof(cl_mem), (void*)&(mpClmem2));
>         if (CL_SUCCESS != err)
>             throw OpenCLError(err);
>@@ -1482,14 +1501,21 @@ public:
>                         sizeof(double)*nVectorWidth, NULL, &err);
>                 if (CL_SUCCESS != err)
>                     throw OpenCLError(err);
>+                SAL_INFO("sc.opencl", "Created buffer " << mpClmem2 << " size " << sizeof(double) << "*" << nVectorWidth << "=" << (sizeof(double)*nVectorWidth));
> 
>                 std::string kernelName = mvSubArguments[0]->GetName() + "_SumIfs_reduction";
>                 cl_kernel redKernel = clCreateKernel(pProgram, kernelName.c_str(), &err);
>                 if (err != CL_SUCCESS)
>                     throw OpenCLError(err);
>+                SAL_INFO("sc.opencl", OUString("Created \"reduction\" kernel ") << kernelName.c_str() << ": " << redKernel);
> 
>                     // set kernel arg of reduction kernel
>                 for (size_t j=0; j< vclmem.size(); j++){
>+                    if (vclmem[j].mCLMem)
>+                        SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << j << ": cl_mem: " << vclmem[j].mCLMem);
>+                    else
>+                        SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << j << ": double: " << vclmem[j].mConst);
>+
>                     err = clSetKernelArg(redKernel, j,
>                             vclmem[j].mCLMem?sizeof(cl_mem):sizeof(double),
>                             vclmem[j].mCLMem?(void *)&vclmem[j].mCLMem:
>@@ -1497,20 +1523,24 @@ public:
>                     if (CL_SUCCESS != err)
>                         throw OpenCLError(err);
>                 }
>+                SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << vclmem.size() << ": cl_mem: " << mpClmem2);
>                 err = clSetKernelArg(redKernel, vclmem.size(), sizeof(cl_mem), (void *)&mpClmem2);
>                 if (CL_SUCCESS != err)
>                     throw OpenCLError(err);
> 
>+                SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << (vclmem.size()+1) << ": cl_int: " << nInput);
>                 err = clSetKernelArg(redKernel, vclmem.size()+1, sizeof(cl_int), (void*)&nInput);
>                 if (CL_SUCCESS != err)
>                     throw OpenCLError(err);
> 
>+                SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << (vclmem.size()+2) << ": cl_int: " << nCurWindowSize);
>                 err = clSetKernelArg(redKernel, vclmem.size()+2, sizeof(cl_int), (void*)&nCurWindowSize);
>                 if (CL_SUCCESS != err)
>                     throw OpenCLError(err);
>                 // set work group size and execute
>                 size_t global_work_size[] = {256, (size_t)nVectorWidth };
>                 size_t local_work_size[] = {256, 1};
>+                SAL_INFO("sc.opencl", "Enqueueing kernel " << redKernel << " work sizes: {256," << nVectorWidth << "}, {256,1}");
>                 err = clEnqueueNDRangeKernel(kEnv.mpkCmdQueue, redKernel, 2, NULL,
>                         global_work_size, local_work_size, 0, NULL, NULL);
>                 if (CL_SUCCESS != err)
>@@ -1518,8 +1548,10 @@ public:
>                 err = clFinish(kEnv.mpkCmdQueue);
>                 if (CL_SUCCESS != err)
>                     throw OpenCLError(err);
>+                SAL_INFO("sc.opencl", "Kernel " << redKernel << " finished");
>                 clReleaseKernel(redKernel);
>                  // Pass mpClmem2 to the "real" kernel
>+                SAL_INFO("sc.opencl", "Kernel " << k << " arg " << argno << ": cl_mem: " << mpClmem2);
>                 err = clSetKernelArg(k, argno, sizeof(cl_mem), (void *)&mpClmem2);
>                 if (CL_SUCCESS != err)
>                     throw OpenCLError(err);
>@@ -2624,16 +2656,20 @@ public:
>                 nr*sizeof(double), NULL, &err);
>         if (CL_SUCCESS != err)
>             throw OpenCLError(err);
>+        SAL_INFO("sc.opencl", "Created buffer " << mpResClmem << " size " << nr << "*" << sizeof(double) << "=" << (nr*sizeof(double)));
>+        SAL_INFO("sc.opencl", "Kernel " << mpKernel << " arg " << 0 << ": cl_mem: " << mpResClmem);
>         err = clSetKernelArg(mpKernel, 0, sizeof(cl_mem), (void*)&mpResClmem);
>         if (CL_SUCCESS != err)
>             throw OpenCLError(err);
>         // The rest of buffers
>         mSyms.Marshal(mpKernel, nr, mpProgram);
>         size_t global_work_size[] = {nr};
>+        SAL_INFO("sc.opencl", "Enqueueing kernel " << mpKernel << " work sizes: {" << nr << "}, NULL");
>         err = clEnqueueNDRangeKernel(kEnv.mpkCmdQueue, mpKernel, 1, NULL,
>             global_work_size, NULL, 0, NULL, NULL);
>         if (CL_SUCCESS != err)
>             throw OpenCLError(err);
>+        SAL_INFO("sc.opencl", "Kernel " << mpKernel << " finished");
>     }
>     ~DynamicKernel();
>     cl_mem GetResultBuffer(void) const { return mpResClmem; }
>@@ -2724,6 +2760,7 @@ void DynamicKernel::CreateKernel(void)
>     mpKernel = clCreateKernel(mpProgram, kname.c_str(), &err);
>     if (err != CL_SUCCESS)
>         throw OpenCLError(err);
>+    SAL_INFO("sc.opencl", "Created kernel " << kname.c_str() << ": " << mpKernel);
> }
> // Symbol lookup. If there is no such symbol created, allocate one
> // kernel with argument with unique name and return so.
>diff --git a/sc/source/core/opencl/op_math.cxx b/sc/source/core/opencl/op_math.cxx
>index e76a5d0..569d715 100644
>--- a/sc/source/core/opencl/op_math.cxx
>+++ b/sc/source/core/opencl/op_math.cxx
>@@ -466,6 +466,7 @@ void OpSumIfs::GenSlidingWindowFunction(std::stringstream &ss,
>         ss << ", __global double *result,int arrayLength,int windowSize";
> 
>         ss << ")\n{\n";
>+        ss << "printf(\"" << __LINE__ << ": " << vSubArguments[0]->GetName() << "_SumIfs_reduction: arrayLength=%d windowSize=%d get_local_id(0)=%d\\n\", arrayLength, windowSize, get_local_id(0));\n";
>         ss << "    double tmp =0;\n";
>         ss << "    int i ;\n";
> 
>@@ -488,6 +489,7 @@ void OpSumIfs::GenSlidingWindowFunction(std::stringstream &ss,
>         ss << "        int loopOffset = l*512;\n";
> 
>         ss << "        int p1 = loopOffset + lidx + offset, p2 = p1 + 256;\n";
>+        ss << "printf(\"" << __LINE__ << ": loopOffset=%d lidx=%d offset=%d p1=%d p2=%d\\n\", loopOffset, lidx, offset, p1, p2);\n";
>         ss << "        if (p2 < min(offset + windowSize, arrayLength)) {\n";
>         ss << "            tmp0 = 0.0;\n";
>         int mm=0;
>@@ -498,6 +500,7 @@ void OpSumIfs::GenSlidingWindowFunction(std::stringstream &ss,
>             CheckSubArgumentIsNan2(ss,vSubArguments,j,p1);
>             CheckSubArgumentIsNan2(ss,vSubArguments,j+1,p1);
>             ss << "";
>+            ss << "printf(\"" << __LINE__ << ": tmp" << j << "=%d tmp" << j+1 << "=%d\\n\", tmp" << j << ", tmp" << j+1 << ");\n";
>             ss <<"    if(isequal(";
>             ss <<"tmp";
>             ss <<j;
>@@ -522,6 +525,7 @@ void OpSumIfs::GenSlidingWindowFunction(std::stringstream &ss,
>         {
>             CheckSubArgumentIsNan2(ss,vSubArguments,j,p2);
>             CheckSubArgumentIsNan2(ss,vSubArguments,j+1,p2);
>+            ss << "printf(\"" << __LINE__ << ": tmp" << j << "=%d tmp" << j+1 << "=%d\\n\", tmp" << j << ", tmp" << j+1 << ");\n";
>             ss <<"    if(isequal(";
>             ss <<"tmp";
>             ss <<j;
>@@ -549,6 +553,7 @@ void OpSumIfs::GenSlidingWindowFunction(std::stringstream &ss,
>             CheckSubArgumentIsNan2(ss,vSubArguments,j,p1);
>             CheckSubArgumentIsNan2(ss,vSubArguments,j+1,p1);
> 
>+            ss << "printf(\"" << __LINE__ << ": tmp" << j << "=%d tmp" << j+1 << "=%d\\n\", tmp" << j << ", tmp" << j+1 << ");\n";
>             ss <<"    if(isequal(";
>             ss <<"tmp";
>             ss <<j;
>diff --git a/sc/source/core/opencl/opbase.cxx b/sc/source/core/opencl/opbase.cxx
>index fd6cdfb..b0cd8e4 100644
>--- a/sc/source/core/opencl/opbase.cxx
>+++ b/sc/source/core/opencl/opbase.cxx
>@@ -187,6 +187,14 @@ void CheckVariables::CheckSubArgumentIsNan2( std::stringstream & ss,
>     }
> 
> #ifdef ISNAN
>+    if(/* i == 4 && */ vSubArguments[i]->GetFormulaToken()->GetType() == formula::svSingleVectorRef)
>+    {
>+        ss << "printf(\"get_group_id(1)=%d ";
>+        vSubArguments[i]->GenDeclRef(ss);
>+        ss << "=%p\\n\",get_group_id(1),";
>+        vSubArguments[i]->GenDeclRef(ss);
>+        ss << ");\n";
>+    }
>     ss<< "    tmp";
>     ss<< i;
>     ss<< "= fsum(";

Actions: View

Attachments on bug 88444: 112270 | 112280 | 112281 | 112341 | 112453 | 112479 | 113075 | 113076 | 113306