From 6a2576150b5152244dc3f8c31b745fa634b31a47 Mon Sep 17 00:00:00 2001 From: Tor Lillqvist Date: Fri, 6 Feb 2015 00:15:04 +0200 Subject: Improve OpenCL error messages Add the name of the function that failed to the OpenCLError class. Log OpenCL failure in a couple of more places (where OpenCLError is not thrown). Print OpenCL errors symbolically instead of numerically where applicable. Change-Id: I60f910e9ea7b75af7ec506553d7a73ad99ba4366 --- sc/source/core/opencl/formulagroupcl.cxx | 136 +++++++++++++++++-------------- sc/source/core/opencl/opbase.cxx | 10 ++- sc/source/core/opencl/opbase.hxx | 3 +- 3 files changed, 82 insertions(+), 67 deletions(-) (limited to 'sc') diff --git a/sc/source/core/opencl/formulagroupcl.cxx b/sc/source/core/opencl/formulagroupcl.cxx index 18aed00b6376..381fb1d75720 100644 --- a/sc/source/core/opencl/formulagroupcl.cxx +++ b/sc/source/core/opencl/formulagroupcl.cxx @@ -169,7 +169,7 @@ size_t VectorRef::Marshal( cl_kernel k, int argno, int, cl_program ) szHostBuffer, pHostBuffer, &err); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clCreateBuffer", err, __FILE__, __LINE__); SAL_INFO("sc.opencl", "Created buffer " << mpClmem << " size " << szHostBuffer << " using host buffer " << pHostBuffer); } else @@ -181,24 +181,28 @@ size_t VectorRef::Marshal( cl_kernel k, int argno, int, cl_program ) (cl_mem_flags)CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, szHostBuffer, NULL, &err); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clCreateBuffer", err, __FILE__, __LINE__); SAL_INFO("sc.opencl", "Created buffer " << mpClmem << " size " << szHostBuffer); double* pNanBuffer = (double*)clEnqueueMapBuffer( kEnv.mpkCmdQueue, mpClmem, CL_TRUE, CL_MAP_WRITE, 0, szHostBuffer, 0, NULL, NULL, &err); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clEnqueueMapBuffer", err, __FILE__, __LINE__); + for (size_t i = 0; i < szHostBuffer / sizeof(double); i++) pNanBuffer[i] = NAN; err = clEnqueueUnmapMemObject(kEnv.mpkCmdQueue, mpClmem, pNanBuffer, 0, NULL, NULL); + // FIXME: Is it intentional to not throw an OpenCLError even if the clEnqueueUnmapMemObject() fails? + if (CL_SUCCESS != err) + SAL_WARN("sc.opencl", "clEnqueueUnmapMemObject failed: " << ::opencl::errorString(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, __FILE__, __LINE__); + throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); return 1; } @@ -257,7 +261,7 @@ public: 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, __FILE__, __LINE__); + throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); return 1; } }; @@ -307,7 +311,7 @@ public: 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, __FILE__, __LINE__); + throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); return 1; } }; @@ -347,7 +351,7 @@ public: 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, __FILE__, __LINE__); + throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); return 1; } }; @@ -707,7 +711,7 @@ threefry2x32 (threefry2x32_ctr_t in, threefry2x32_key_t k)\n\ SAL_INFO("sc.opencl", "Kernel " << k << " arg " << argno << ": cl_int: " << seed); cl_int err = clSetKernelArg(k, argno, sizeof(cl_int), (void*)&seed); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); return 1; } }; @@ -767,14 +771,14 @@ size_t DynamicKernelStringArgument::Marshal( cl_kernel k, int argno, int, cl_pro (cl_mem_flags)CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, szHostBuffer, NULL, &err); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clCreateBuffer", err, __FILE__, __LINE__); SAL_INFO("sc.opencl", "Created buffer " << mpClmem << " size " << szHostBuffer); pHashBuffer = (cl_uint*)clEnqueueMapBuffer( kEnv.mpkCmdQueue, mpClmem, CL_TRUE, CL_MAP_WRITE, 0, szHostBuffer, 0, NULL, NULL, &err); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clEnqueueMapBuffer", err, __FILE__, __LINE__); for (size_t i = 0; i < nStrings; i++) { @@ -798,14 +802,14 @@ size_t DynamicKernelStringArgument::Marshal( cl_kernel k, int argno, int, cl_pro (cl_mem_flags)CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, szHostBuffer, NULL, &err); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clCreateBuffer", err, __FILE__, __LINE__); SAL_INFO("sc.opencl", "Created buffer " << mpClmem << " size " << szHostBuffer); pHashBuffer = (cl_uint*)clEnqueueMapBuffer( kEnv.mpkCmdQueue, mpClmem, CL_TRUE, CL_MAP_WRITE, 0, szHostBuffer, 0, NULL, NULL, &err); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clEnqueueMapBuffer", err, __FILE__, __LINE__); for (size_t i = 0; i < szHostBuffer / sizeof(cl_int); i++) pHashBuffer[i] = 0; @@ -813,12 +817,12 @@ size_t DynamicKernelStringArgument::Marshal( cl_kernel k, int argno, int, cl_pro err = clEnqueueUnmapMemObject(kEnv.mpkCmdQueue, mpClmem, pHashBuffer, 0, NULL, NULL); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clEnqueueUnmapMemObject", err, __FILE__, __LINE__); 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, __FILE__, __LINE__); + throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); return 1; } @@ -1449,7 +1453,7 @@ public: CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, sizeof(double) * w, NULL, NULL); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clCreateBuffer", err, __FILE__, __LINE__); SAL_INFO("sc.opencl", "Created buffer " << mpClmem2 << " size " << sizeof(double) << "*" << w << "=" << (sizeof(double)*w)); // reproduce the reduction function name @@ -1460,7 +1464,7 @@ public: kernelName = Base::GetName() + "_sum_reduction"; cl_kernel redKernel = clCreateKernel(mpProgram, kernelName.c_str(), &err); if (err != CL_SUCCESS) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clCreateKernel", err, __FILE__, __LINE__); // set kernel arg of reduction kernel // TODO(Wei Wei): use unique name for kernel cl_mem buf = Base::GetCLBuffer(); @@ -1468,22 +1472,22 @@ public: err = clSetKernelArg(redKernel, 0, sizeof(cl_mem), (void*)&buf); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); 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, __FILE__, __LINE__); + throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); 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, __FILE__, __LINE__); + throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 3 << ": cl_int: " << nCurWindowSize); err = clSetKernelArg(redKernel, 3, sizeof(cl_int), (void*)&nCurWindowSize); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); // set work group size and execute size_t global_work_size[] = { 256, (size_t)w }; @@ -1492,10 +1496,10 @@ public: err = clEnqueueNDRangeKernel(kEnv.mpkCmdQueue, redKernel, 2, NULL, global_work_size, local_work_size, 0, NULL, NULL); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clEnqueueNDRangeKernel", err, __FILE__, __LINE__); err = clFinish(kEnv.mpkCmdQueue); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clFinish", err, __FILE__, __LINE__); if (dynamic_cast(mpCodeGen.get())) { /*average need more reduction kernel for count computing*/ @@ -1506,40 +1510,40 @@ public: sizeof(double) * w, 0, NULL, NULL, &err); if (err != CL_SUCCESS) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clEnqueueMapBuffer", err, __FILE__, __LINE__); for (int i = 0; i < w; i++) pAllBuffer[i] = resbuf[i]; err = clEnqueueUnmapMemObject(kEnv.mpkCmdQueue, mpClmem2, resbuf, 0, NULL, NULL); if (err != CL_SUCCESS) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clEnqueueUnmapMemObject", err, __FILE__, __LINE__); kernelName = Base::GetName() + "_count_reduction"; redKernel = clCreateKernel(mpProgram, kernelName.c_str(), &err); if (err != CL_SUCCESS) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clCreateKernel", err, __FILE__, __LINE__); // set kernel arg of reduction kernel 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, __FILE__, __LINE__); + throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); 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, __FILE__, __LINE__); + throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); 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, __FILE__, __LINE__); + throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 3 << ": cl_int: " << nCurWindowSize); err = clSetKernelArg(redKernel, 3, sizeof(cl_int), (void*)&nCurWindowSize); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); // set work group size and execute size_t global_work_size1[] = { 256, (size_t)w }; @@ -1548,20 +1552,23 @@ public: err = clEnqueueNDRangeKernel(kEnv.mpkCmdQueue, redKernel, 2, NULL, global_work_size1, local_work_size1, 0, NULL, NULL); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clEnqueueNDRangeKernel", err, __FILE__, __LINE__); err = clFinish(kEnv.mpkCmdQueue); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clFinish", err, __FILE__, __LINE__); resbuf = (double*)clEnqueueMapBuffer(kEnv.mpkCmdQueue, mpClmem2, CL_TRUE, CL_MAP_READ, 0, sizeof(double) * w, 0, NULL, NULL, &err); if (err != CL_SUCCESS) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clEnqueueMapBuffer", err, __FILE__, __LINE__); for (int i = 0; i < w; i++) pAllBuffer[i + w] = resbuf[i]; err = clEnqueueUnmapMemObject(kEnv.mpkCmdQueue, mpClmem2, resbuf, 0, NULL, NULL); + // FIXME: Is it intentional to not throw an OpenCLError even if the clEnqueueUnmapMemObject() fails? + if (CL_SUCCESS != err) + SAL_WARN("sc.opencl", "clEnqueueUnmapMemObject failed: " << ::opencl::errorString(err)); if (mpClmem2) { clReleaseMemObject(mpClmem2); @@ -1571,14 +1578,14 @@ public: (cl_mem_flags)CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, w * sizeof(double) * 2, pAllBuffer.get(), &err); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clCreateBuffer", err, __FILE__, __LINE__); SAL_INFO("sc.opencl", "Created buffer " << mpClmem2 << " size " << w << "*" << sizeof(double) << "=" << (w*sizeof(double)) << " copying host buffer " << pAllBuffer.get()); } // 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, __FILE__, __LINE__); + throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); return 1; } ~ParallelReductionVectorRef() @@ -2192,13 +2199,13 @@ public: pClmem2 = clCreateBuffer(kEnv.mpkContext, CL_MEM_READ_WRITE, sizeof(double) * nVectorWidth, NULL, &err); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clCreateBuffer", err, __FILE__, __LINE__); SAL_INFO("sc.opencl", "Created buffer " << pClmem2 << " size " << sizeof(double) << "*" << nVectorWidth << "=" << (sizeof(double)*nVectorWidth)); std::string kernelName = "GeoMean_reduction"; cl_kernel redKernel = clCreateKernel(pProgram, kernelName.c_str(), &err); if (err != CL_SUCCESS) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clCreateKernel", err, __FILE__, __LINE__); // set kernel arg of reduction kernel for (size_t j = 0; j < vclmem.size(); j++) { @@ -2207,12 +2214,12 @@ public: vclmem[j] ? sizeof(cl_mem) : sizeof(double), (void*)&vclmem[j]); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); } SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << vclmem.size() << ": cl_mem: " << pClmem2); err = clSetKernelArg(redKernel, vclmem.size(), sizeof(cl_mem), (void*)&pClmem2); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); // set work group size and execute size_t global_work_size[] = { 256, (size_t)nVectorWidth }; @@ -2221,16 +2228,16 @@ public: err = clEnqueueNDRangeKernel(kEnv.mpkCmdQueue, redKernel, 2, NULL, global_work_size, local_work_size, 0, NULL, NULL); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clEnqueueNDRangeKernel", err, __FILE__, __LINE__); err = clFinish(kEnv.mpkCmdQueue); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clFinish", err, __FILE__, __LINE__); // Pass pClmem2 to the "real" kernel SAL_INFO("sc.opencl", "Kernel " << k << " arg " << argno << ": cl_mem: " << pClmem2); err = clSetKernelArg(k, argno, sizeof(cl_mem), (void*)&pClmem2); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); } } if (OpSumIfs* OpSumCodeGen = dynamic_cast(mpCodeGen.get())) @@ -2263,13 +2270,13 @@ public: mpClmem2 = clCreateBuffer(kEnv.mpkContext, CL_MEM_READ_WRITE, sizeof(double) * nVectorWidth, NULL, &err); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clCreateBuffer", err, __FILE__, __LINE__); 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, __FILE__, __LINE__); + throw OpenCLError("clCreateKernel", err, __FILE__, __LINE__); // set kernel arg of reduction kernel for (size_t j = 0; j < vclmem.size(); j++) @@ -2283,22 +2290,22 @@ public: vclmem[j].mCLMem ? (void*)&vclmem[j].mCLMem : (void*)&vclmem[j].mConst); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); } 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, __FILE__, __LINE__); + throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); 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, __FILE__, __LINE__); + throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); 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, __FILE__, __LINE__); + throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); // set work group size and execute size_t global_work_size[] = { 256, (size_t)nVectorWidth }; size_t local_work_size[] = { 256, 1 }; @@ -2306,16 +2313,16 @@ public: err = clEnqueueNDRangeKernel(kEnv.mpkCmdQueue, redKernel, 2, NULL, global_work_size, local_work_size, 0, NULL, NULL); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clEnqueueNDRangeKernel", err, __FILE__, __LINE__); err = clFinish(kEnv.mpkCmdQueue); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clFinish", err, __FILE__, __LINE__); 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, __FILE__, __LINE__); + throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); } } return i; @@ -3712,7 +3719,7 @@ void DynamicKernel::CreateKernel() mpProgram = clCreateProgramWithSource(kEnv.mpkContext, 1, &src, NULL, &err); if (err != CL_SUCCESS) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clCreateProgramWithSource", err, __FILE__, __LINE__); err = clBuildProgram(mpProgram, 1, ::opencl::gpuEnv.mpArryDevsID, "", NULL, NULL); if (err != CL_SUCCESS) @@ -3729,7 +3736,7 @@ void DynamicKernel::CreateKernel() e != CL_SUCCESS, "sc.opencl", "after CL_BUILD_PROGRAM_FAILURE," " clGetProgramBuildInfo(CL_PROGRAM_BUILD_STATUS)" - " fails with " << e); + " fails with " << ::opencl::errorString(e)); if (e == CL_SUCCESS) { size_t n; @@ -3740,7 +3747,7 @@ void DynamicKernel::CreateKernel() e != CL_SUCCESS || n == 0, "sc.opencl", "after CL_BUILD_PROGRAM_FAILURE," " clGetProgramBuildInfo(CL_PROGRAM_BUILD_LOG)" - " fails with " << e << ", n=" << n); + " fails with " << ::opencl::errorString(e) << ", n=" << n); if (e == CL_SUCCESS && n != 0) { std::vector log(n); @@ -3751,7 +3758,7 @@ void DynamicKernel::CreateKernel() e != CL_SUCCESS || n == 0, "sc.opencl", "after CL_BUILD_PROGRAM_FAILURE," " clGetProgramBuildInfo(" - "CL_PROGRAM_BUILD_LOG) fails with " << e); + "CL_PROGRAM_BUILD_LOG) fails with " << ::opencl::errorString(e)); if (e == CL_SUCCESS) SAL_WARN( "sc.opencl", @@ -3761,7 +3768,7 @@ void DynamicKernel::CreateKernel() } } #endif - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clBuildProgram", err, __FILE__, __LINE__); } // Generate binary out of compiled kernel. ::opencl::generatBinFromKernelSource(mpProgram, @@ -3774,7 +3781,7 @@ void DynamicKernel::CreateKernel() } mpKernel = clCreateKernel(mpProgram, kname.c_str(), &err); if (err != CL_SUCCESS) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clCreateKernel", err, __FILE__, __LINE__); } void DynamicKernel::Launch( size_t nr ) @@ -3788,13 +3795,13 @@ void DynamicKernel::Launch( size_t nr ) (cl_mem_flags)CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, nr * sizeof(double), NULL, &err); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clCreateBuffer", err, __FILE__, __LINE__); 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, __FILE__, __LINE__); + throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); // The rest of buffers mSyms.Marshal(mpKernel, nr, mpProgram); size_t global_work_size[] = { nr }; @@ -3802,10 +3809,10 @@ void DynamicKernel::Launch( size_t nr ) err = clEnqueueNDRangeKernel(kEnv.mpkCmdQueue, mpKernel, 1, NULL, global_work_size, NULL, 0, NULL, NULL); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clEnqueueNDRangeKernel", err, __FILE__, __LINE__); err = clFlush(kEnv.mpkCmdQueue); if (CL_SUCCESS != err) - throw OpenCLError(err, __FILE__, __LINE__); + throw OpenCLError("clFlush", err, __FILE__, __LINE__); } // Symbol lookup. If there is no such symbol created, allocate one @@ -3954,7 +3961,7 @@ public: if (err != CL_SUCCESS) { - SAL_WARN("sc.opencl", "Dynamic formula compiler: OpenCL error: " << err); + SAL_WARN("sc.opencl", "clEnqueueMapBuffer failed:: " << ::opencl::errorString(err)); mpResBuf = NULL; return; } @@ -3962,7 +3969,7 @@ public: err = clEnqueueUnmapMemObject(kEnv.mpkCmdQueue, mpCLResBuf, mpResBuf, 0, NULL, NULL); if (err != CL_SUCCESS) { - SAL_WARN("sc.opencl", "Dynamic formula compiler: OpenCL error: " << err); + SAL_WARN("sc.opencl", "clEnqueueUnmapMemObject failed: " << ::opencl::errorString(err)); mpResBuf = NULL; } } @@ -4023,7 +4030,7 @@ public: } catch (const OpenCLError& oce) { - SAL_WARN("sc.opencl", "Dynamic formula compiler: OpenCL error: " << oce.mError << " at " << oce.mFile << ":" << oce.mLineNumber); + SAL_WARN("sc.opencl", "Dynamic formula compiler: OpenCL error from " << oce.mFunction << ": " << ::opencl::errorString(oce.mError) << " at " << oce.mFile << ":" << oce.mLineNumber); return CLInterpreterResult(); } catch (const Unhandled& uh) @@ -4088,6 +4095,9 @@ bool waitForResults() ::opencl::setKernelEnv(&kEnv); cl_int err = clFinish(kEnv.mpkCmdQueue); + if (err != CL_SUCCESS) + SAL_WARN("sc.opencl", "clFinish failed: " << ::opencl::errorString(err)); + return err == CL_SUCCESS; } diff --git a/sc/source/core/opencl/opbase.cxx b/sc/source/core/opencl/opbase.cxx index 1c24c12691f7..671987e103ae 100644 --- a/sc/source/core/opencl/opbase.cxx +++ b/sc/source/core/opencl/opbase.cxx @@ -19,10 +19,14 @@ UnhandledToken::UnhandledToken( formula::FormulaToken* t, const char* m, const std::string& fn, int ln ) : mToken(t), mMessage(m), mFile(fn), mLineNumber(ln) {} -OpenCLError::OpenCLError( cl_int err, const std::string& fn, int ln ) : - mError(err), mFile(fn), mLineNumber(ln) +OpenCLError::OpenCLError( const std::string function, cl_int error, const std::string& file, int line ) : + mFunction(function), mError(error), mFile(file), mLineNumber(line) { - SAL_INFO("sc.opencl", "OpenCL error: " << ::opencl::errorString(mError)); + // Not sure if this SAL_INFO() is useful; the place in + // CLInterpreterContext::launchKernel() where OpenCLError is + // caught already uses SAL_WARN() to display it. + + // SAL_INFO("sc.opencl", "OpenCL error: " << ::opencl::errorString(mError)); } Unhandled::Unhandled( const std::string& fn, int ln ) : diff --git a/sc/source/core/opencl/opbase.hxx b/sc/source/core/opencl/opbase.hxx index b4c4493869fb..5c7228d93975 100644 --- a/sc/source/core/opencl/opbase.hxx +++ b/sc/source/core/opencl/opbase.hxx @@ -42,8 +42,9 @@ public: class OpenCLError { public: - OpenCLError( cl_int err, const std::string& fn, int ln ); + OpenCLError( const std::string function, cl_int error, const std::string& file, int line ); + std::string mFunction; cl_int mError; std::string mFile; int mLineNumber; -- cgit