diff options
author | Dima Panov <fluffy@FreeBSD.org> | 2020-07-07 03:14:58 +0000 |
---|---|---|
committer | Dima Panov <fluffy@FreeBSD.org> | 2020-07-07 03:14:58 +0000 |
commit | 2d98f1c14da2bc2da0ea8e65c4a5b93b4ab49d4d (patch) | |
tree | 16394ca8b04abdbe1393df7041d1ffa359e32354 /editors/libreoffice | |
parent | 22f2f91d5c019a5a8f9843636439c055741db9ce (diff) | |
download | ports-2d98f1c14da2bc2da0ea8e65c4a5b93b4ab49d4d.tar.gz ports-2d98f1c14da2bc2da0ea8e65c4a5b93b4ab49d4d.zip |
Update LibreOffice suite to 6.4.5 release
QT5 VCL fixes:
- Reduce startup flickering
- Always use cairo font renderer
Special thanks to: tijl
PR: 247444
Notes
Notes:
svn path=/head/; revision=541383
Diffstat (limited to 'editors/libreoffice')
-rw-r--r-- | editors/libreoffice/Makefile | 1 | ||||
-rw-r--r-- | editors/libreoffice/Makefile.common | 2 | ||||
-rw-r--r-- | editors/libreoffice/distinfo | 10 | ||||
-rw-r--r-- | editors/libreoffice/files/patch-sc_source_core_opencl_formulagroupcl.cxx | 780 | ||||
-rw-r--r-- | editors/libreoffice/files/patch-sdext_source_pdfimport_xpdfwrapper_pdfioutdev__gpl.cxx | 18 | ||||
-rw-r--r-- | editors/libreoffice/files/patch-vcl_qt5_Qt5Frame.cxx | 11 | ||||
-rw-r--r-- | editors/libreoffice/files/patch-vcl_qt5_Qt5Instance.cxx | 35 |
7 files changed, 52 insertions, 805 deletions
diff --git a/editors/libreoffice/Makefile b/editors/libreoffice/Makefile index 23987a705609..819f6179aa36 100644 --- a/editors/libreoffice/Makefile +++ b/editors/libreoffice/Makefile @@ -1,5 +1,4 @@ # $FreeBSD$ -PORTREVISION= 1 .include "${.CURDIR}/Makefile.common" diff --git a/editors/libreoffice/Makefile.common b/editors/libreoffice/Makefile.common index 7eafaa7959a7..4fa98e5823cd 100644 --- a/editors/libreoffice/Makefile.common +++ b/editors/libreoffice/Makefile.common @@ -1,7 +1,7 @@ # $FreeBSD$ PORTNAME= libreoffice -PORTVERSION= 6.4.4 +PORTVERSION= 6.4.5 CATEGORIES+= editors MAINTAINER= office@FreeBSD.org diff --git a/editors/libreoffice/distinfo b/editors/libreoffice/distinfo index 82b8655cc00a..a864333c427b 100644 --- a/editors/libreoffice/distinfo +++ b/editors/libreoffice/distinfo @@ -1,8 +1,8 @@ -TIMESTAMP = 1590779230 -SHA256 (libreoffice/libreoffice-6.4.4.2.tar.xz) = 54388597dffc9c32f81446e6e634f7af76ca0e0e5a0d27bc3fe89033a011c078 -SIZE (libreoffice/libreoffice-6.4.4.2.tar.xz) = 231325236 -SHA256 (libreoffice/libreoffice-help-6.4.4.2.tar.xz) = c61ac7b9ceac9aa0813bce405ce25b13dc1b698509ce33827ff7c65dffde25ac -SIZE (libreoffice/libreoffice-help-6.4.4.2.tar.xz) = 88296340 +TIMESTAMP = 1593246155 +SHA256 (libreoffice/libreoffice-6.4.5.2.tar.xz) = 6a4cc6b9ca838ca997b83181d8d163b1981de2d4d1268387741d342453a491ec +SIZE (libreoffice/libreoffice-6.4.5.2.tar.xz) = 236255044 +SHA256 (libreoffice/libreoffice-help-6.4.5.2.tar.xz) = 66675ccb023a8ed88cdad6e877cfef55ea139630cdf8df17516835c7456e2fa2 +SIZE (libreoffice/libreoffice-help-6.4.5.2.tar.xz) = 88376816 SHA256 (libreoffice/884ed41809687c3e168fc7c19b16585149ff058eca79acbf3ee784f6630704cc-opens___.ttf) = 884ed41809687c3e168fc7c19b16585149ff058eca79acbf3ee784f6630704cc SIZE (libreoffice/884ed41809687c3e168fc7c19b16585149ff058eca79acbf3ee784f6630704cc-opens___.ttf) = 207544 SHA256 (libreoffice/17410483b5b5f267aa18b7e00b65e6e0-hsqldb_1_8_0.zip) = d30b13f4ba2e3b6a2d4f020c0dee0a9fb9fc6fbcc2d561f36b78da4bf3802370 diff --git a/editors/libreoffice/files/patch-sc_source_core_opencl_formulagroupcl.cxx b/editors/libreoffice/files/patch-sc_source_core_opencl_formulagroupcl.cxx deleted file mode 100644 index d995f1d25eb7..000000000000 --- a/editors/libreoffice/files/patch-sc_source_core_opencl_formulagroupcl.cxx +++ /dev/null @@ -1,780 +0,0 @@ -https://bugs.gentoo.org/713574 -https://bugs.documentfoundation.org/show_bug.cgi?id=131591 - ---- sc/source/core/opencl/formulagroupcl.cxx.orig 2020-03-11 16:18:35 UTC -+++ sc/source/core/opencl/formulagroupcl.cxx -@@ -1026,8 +1026,6 @@ class DynamicKernelMixedArgument : public VectorRef (p - /// Handling a Double Vector that is used as a sliding window input - /// to either a sliding window average or sum-of-products - /// Generate a sequential loop for reductions --class OpAverage; --class OpCount; - - template<class Base> - class DynamicKernelSlidingArgument : public Base -@@ -1335,186 +1333,8 @@ class ParallelReductionVectorRef : public Base (public - } - - /// Emit the definition for the auxiliary reduction kernel -- virtual void GenSlidingWindowFunction( std::stringstream& ss ) -- { -- if (!dynamic_cast<OpAverage*>(mpCodeGen.get())) -- { -- std::string name = Base::GetName(); -- ss << "__kernel void " << name; -- ss << "_reduction(__global double* A, " -- "__global double *result,int arrayLength,int windowSize){\n"; -- ss << " double tmp, current_result =" << -- mpCodeGen->GetBottom(); -- ss << ";\n"; -- ss << " int writePos = get_group_id(1);\n"; -- ss << " int lidx = get_local_id(0);\n"; -- ss << " __local double shm_buf[256];\n"; -- if (mpDVR->IsStartFixed()) -- ss << " int offset = 0;\n"; -- else // if (!mpDVR->IsStartFixed()) -- ss << " int offset = get_group_id(1);\n"; -- if (mpDVR->IsStartFixed() && mpDVR->IsEndFixed()) -- ss << " int end = windowSize;\n"; -- else if (!mpDVR->IsStartFixed() && !mpDVR->IsEndFixed()) -- ss << " int end = offset + windowSize;\n"; -- else if (mpDVR->IsStartFixed() && !mpDVR->IsEndFixed()) -- ss << " int end = windowSize + get_group_id(1);\n"; -- else if (!mpDVR->IsStartFixed() && mpDVR->IsEndFixed()) -- ss << " int end = windowSize;\n"; -- ss << " end = min(end, arrayLength);\n"; -+ virtual void GenSlidingWindowFunction( std::stringstream& ss ); - -- ss << " barrier(CLK_LOCAL_MEM_FENCE);\n"; -- ss << " int loop = arrayLength/512 + 1;\n"; -- ss << " for (int l=0; l<loop; l++){\n"; -- ss << " tmp = " << mpCodeGen->GetBottom() << ";\n"; -- ss << " int loopOffset = l*512;\n"; -- ss << " if((loopOffset + lidx + offset + 256) < end) {\n"; -- ss << " tmp = legalize(" << mpCodeGen->Gen2( -- "A[loopOffset + lidx + offset]", "tmp") << ", tmp);\n"; -- ss << " tmp = legalize(" << mpCodeGen->Gen2( -- "A[loopOffset + lidx + offset + 256]", "tmp") << ", tmp);\n"; -- ss << " } else if ((loopOffset + lidx + offset) < end)\n"; -- ss << " tmp = legalize(" << mpCodeGen->Gen2( -- "A[loopOffset + lidx + offset]", "tmp") << ", tmp);\n"; -- ss << " shm_buf[lidx] = tmp;\n"; -- ss << " barrier(CLK_LOCAL_MEM_FENCE);\n"; -- ss << " for (int i = 128; i >0; i/=2) {\n"; -- ss << " if (lidx < i)\n"; -- ss << " shm_buf[lidx] = "; -- // Special case count -- if (dynamic_cast<OpCount*>(mpCodeGen.get())) -- ss << "shm_buf[lidx] + shm_buf[lidx + i];\n"; -- else -- ss << mpCodeGen->Gen2("shm_buf[lidx]", "shm_buf[lidx + i]") << ";\n"; -- ss << " barrier(CLK_LOCAL_MEM_FENCE);\n"; -- ss << " }\n"; -- ss << " if (lidx == 0)\n"; -- ss << " current_result ="; -- if (dynamic_cast<OpCount*>(mpCodeGen.get())) -- ss << "current_result + shm_buf[0]"; -- else -- ss << mpCodeGen->Gen2("current_result", "shm_buf[0]"); -- ss << ";\n"; -- ss << " barrier(CLK_LOCAL_MEM_FENCE);\n"; -- ss << " }\n"; -- ss << " if (lidx == 0)\n"; -- ss << " result[writePos] = current_result;\n"; -- ss << "}\n"; -- } -- else -- { -- std::string name = Base::GetName(); -- /*sum reduction*/ -- ss << "__kernel void " << name << "_sum"; -- ss << "_reduction(__global double* A, " -- "__global double *result,int arrayLength,int windowSize){\n"; -- ss << " double tmp, current_result =" << -- mpCodeGen->GetBottom(); -- ss << ";\n"; -- ss << " int writePos = get_group_id(1);\n"; -- ss << " int lidx = get_local_id(0);\n"; -- ss << " __local double shm_buf[256];\n"; -- if (mpDVR->IsStartFixed()) -- ss << " int offset = 0;\n"; -- else // if (!mpDVR->IsStartFixed()) -- ss << " int offset = get_group_id(1);\n"; -- if (mpDVR->IsStartFixed() && mpDVR->IsEndFixed()) -- ss << " int end = windowSize;\n"; -- else if (!mpDVR->IsStartFixed() && !mpDVR->IsEndFixed()) -- ss << " int end = offset + windowSize;\n"; -- else if (mpDVR->IsStartFixed() && !mpDVR->IsEndFixed()) -- ss << " int end = windowSize + get_group_id(1);\n"; -- else if (!mpDVR->IsStartFixed() && mpDVR->IsEndFixed()) -- ss << " int end = windowSize;\n"; -- ss << " end = min(end, arrayLength);\n"; -- ss << " barrier(CLK_LOCAL_MEM_FENCE);\n"; -- ss << " int loop = arrayLength/512 + 1;\n"; -- ss << " for (int l=0; l<loop; l++){\n"; -- ss << " tmp = " << mpCodeGen->GetBottom() << ";\n"; -- ss << " int loopOffset = l*512;\n"; -- ss << " if((loopOffset + lidx + offset + 256) < end) {\n"; -- ss << " tmp = legalize("; -- ss << "(A[loopOffset + lidx + offset]+ tmp)"; -- ss << ", tmp);\n"; -- ss << " tmp = legalize((A[loopOffset + lidx + offset + 256]+ tmp)"; -- ss << ", tmp);\n"; -- ss << " } else if ((loopOffset + lidx + offset) < end)\n"; -- ss << " tmp = legalize((A[loopOffset + lidx + offset] + tmp)"; -- ss << ", tmp);\n"; -- ss << " shm_buf[lidx] = tmp;\n"; -- ss << " barrier(CLK_LOCAL_MEM_FENCE);\n"; -- ss << " for (int i = 128; i >0; i/=2) {\n"; -- ss << " if (lidx < i)\n"; -- ss << " shm_buf[lidx] = "; -- ss << "shm_buf[lidx] + shm_buf[lidx + i];\n"; -- ss << " barrier(CLK_LOCAL_MEM_FENCE);\n"; -- ss << " }\n"; -- ss << " if (lidx == 0)\n"; -- ss << " current_result ="; -- ss << "current_result + shm_buf[0]"; -- ss << ";\n"; -- ss << " barrier(CLK_LOCAL_MEM_FENCE);\n"; -- ss << " }\n"; -- ss << " if (lidx == 0)\n"; -- ss << " result[writePos] = current_result;\n"; -- ss << "}\n"; -- /*count reduction*/ -- ss << "__kernel void " << name << "_count"; -- ss << "_reduction(__global double* A, " -- "__global double *result,int arrayLength,int windowSize){\n"; -- ss << " double tmp, current_result =" << -- mpCodeGen->GetBottom(); -- ss << ";\n"; -- ss << " int writePos = get_group_id(1);\n"; -- ss << " int lidx = get_local_id(0);\n"; -- ss << " __local double shm_buf[256];\n"; -- if (mpDVR->IsStartFixed()) -- ss << " int offset = 0;\n"; -- else // if (!mpDVR->IsStartFixed()) -- ss << " int offset = get_group_id(1);\n"; -- if (mpDVR->IsStartFixed() && mpDVR->IsEndFixed()) -- ss << " int end = windowSize;\n"; -- else if (!mpDVR->IsStartFixed() && !mpDVR->IsEndFixed()) -- ss << " int end = offset + windowSize;\n"; -- else if (mpDVR->IsStartFixed() && !mpDVR->IsEndFixed()) -- ss << " int end = windowSize + get_group_id(1);\n"; -- else if (!mpDVR->IsStartFixed() && mpDVR->IsEndFixed()) -- ss << " int end = windowSize;\n"; -- ss << " end = min(end, arrayLength);\n"; -- ss << " barrier(CLK_LOCAL_MEM_FENCE);\n"; -- ss << " int loop = arrayLength/512 + 1;\n"; -- ss << " for (int l=0; l<loop; l++){\n"; -- ss << " tmp = " << mpCodeGen->GetBottom() << ";\n"; -- ss << " int loopOffset = l*512;\n"; -- ss << " if((loopOffset + lidx + offset + 256) < end) {\n"; -- ss << " tmp = legalize((isnan(A[loopOffset + lidx + offset])?tmp:tmp+1.0)"; -- ss << ", tmp);\n"; -- ss << " tmp = legalize((isnan(A[loopOffset + lidx + offset+256])?tmp:tmp+1.0)"; -- ss << ", tmp);\n"; -- ss << " } else if ((loopOffset + lidx + offset) < end)\n"; -- ss << " tmp = legalize((isnan(A[loopOffset + lidx + offset])?tmp:tmp+1.0)"; -- ss << ", tmp);\n"; -- ss << " shm_buf[lidx] = tmp;\n"; -- ss << " barrier(CLK_LOCAL_MEM_FENCE);\n"; -- ss << " for (int i = 128; i >0; i/=2) {\n"; -- ss << " if (lidx < i)\n"; -- ss << " shm_buf[lidx] = "; -- ss << "shm_buf[lidx] + shm_buf[lidx + i];\n"; -- ss << " barrier(CLK_LOCAL_MEM_FENCE);\n"; -- ss << " }\n"; -- ss << " if (lidx == 0)\n"; -- ss << " current_result ="; -- ss << "current_result + shm_buf[0];"; -- ss << ";\n"; -- ss << " barrier(CLK_LOCAL_MEM_FENCE);\n"; -- ss << " }\n"; -- ss << " if (lidx == 0)\n"; -- ss << " result[writePos] = current_result;\n"; -- ss << "}\n"; -- } -- -- } -- - virtual std::string GenSlidingWindowDeclRef( bool ) const - { - std::stringstream ss; -@@ -1527,195 +1347,10 @@ class ParallelReductionVectorRef : public Base (public - - /// Controls how the elements in the DoubleVectorRef are traversed - size_t GenReductionLoopHeader( -- std::stringstream& ss, int nResultSize, bool& needBody ) -- { -- assert(mpDVR); -- size_t nCurWindowSize = mpDVR->GetRefRowSize(); -- std::string temp = Base::GetName() + "[gid0]"; -- ss << "tmp = "; -- // Special case count -- if (dynamic_cast<OpAverage*>(mpCodeGen.get())) -- { -- ss << mpCodeGen->Gen2(temp, "tmp") << ";\n"; -- ss << "nCount = nCount-1;\n"; -- ss << "nCount = nCount +"; /*re-assign nCount from count reduction*/ -- ss << Base::GetName() << "[gid0+" << nResultSize << "]" << ";\n"; -- } -- else if (dynamic_cast<OpCount*>(mpCodeGen.get())) -- ss << temp << "+ tmp"; -- else -- ss << mpCodeGen->Gen2(temp, "tmp"); -- ss << ";\n\t"; -- needBody = false; -- return nCurWindowSize; -- } -+ std::stringstream& ss, int nResultSize, bool& needBody ); - -- virtual size_t Marshal( cl_kernel k, int argno, int w, cl_program mpProgram ) -- { -- assert(Base::mpClmem == nullptr); -+ virtual size_t Marshal( cl_kernel k, int argno, int w, cl_program mpProgram ); - -- openclwrapper::KernelEnv kEnv; -- openclwrapper::setKernelEnv(&kEnv); -- cl_int err; -- size_t nInput = mpDVR->GetArrayLength(); -- size_t nCurWindowSize = mpDVR->GetRefRowSize(); -- // create clmem buffer -- if (mpDVR->GetArrays()[Base::mnIndex].mpNumericArray == nullptr) -- throw Unhandled(__FILE__, __LINE__); -- double* pHostBuffer = const_cast<double*>( -- mpDVR->GetArrays()[Base::mnIndex].mpNumericArray); -- size_t szHostBuffer = nInput * sizeof(double); -- Base::mpClmem = clCreateBuffer(kEnv.mpkContext, -- 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 << " using host buffer " << pHostBuffer); -- -- mpClmem2 = clCreateBuffer(kEnv.mpkContext, -- CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, -- sizeof(double) * w, nullptr, nullptr); -- if (CL_SUCCESS != err) -- 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 -- std::string kernelName; -- if (!dynamic_cast<OpAverage*>(mpCodeGen.get())) -- kernelName = Base::GetName() + "_reduction"; -- else -- kernelName = Base::GetName() + "_sum_reduction"; -- cl_kernel redKernel = clCreateKernel(mpProgram, kernelName.c_str(), &err); -- if (err != CL_SUCCESS) -- throw OpenCLError("clCreateKernel", err, __FILE__, __LINE__); -- SAL_INFO("sc.opencl", "Created kernel " << redKernel << " with name " << kernelName << " in program " << mpProgram); -- -- // 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), -- static_cast<void*>(&buf)); -- if (CL_SUCCESS != err) -- throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); -- -- SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 1 << ": cl_mem: " << mpClmem2); -- err = clSetKernelArg(redKernel, 1, sizeof(cl_mem), &mpClmem2); -- if (CL_SUCCESS != err) -- throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); -- -- SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 2 << ": cl_int: " << nInput); -- err = clSetKernelArg(redKernel, 2, sizeof(cl_int), static_cast<void*>(&nInput)); -- if (CL_SUCCESS != err) -- throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); -- -- SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 3 << ": cl_int: " << nCurWindowSize); -- err = clSetKernelArg(redKernel, 3, sizeof(cl_int), static_cast<void*>(&nCurWindowSize)); -- if (CL_SUCCESS != err) -- throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); -- -- // set work group size and execute -- size_t global_work_size[] = { 256, static_cast<size_t>(w) }; -- size_t const local_work_size[] = { 256, 1 }; -- SAL_INFO("sc.opencl", "Enqueing kernel " << redKernel); -- err = clEnqueueNDRangeKernel(kEnv.mpkCmdQueue, redKernel, 2, nullptr, -- global_work_size, local_work_size, 0, nullptr, nullptr); -- if (CL_SUCCESS != err) -- throw OpenCLError("clEnqueueNDRangeKernel", err, __FILE__, __LINE__); -- err = clFinish(kEnv.mpkCmdQueue); -- if (CL_SUCCESS != err) -- throw OpenCLError("clFinish", err, __FILE__, __LINE__); -- if (dynamic_cast<OpAverage*>(mpCodeGen.get())) -- { -- /*average need more reduction kernel for count computing*/ -- std::unique_ptr<double[]> pAllBuffer(new double[2 * w]); -- double* resbuf = static_cast<double*>(clEnqueueMapBuffer(kEnv.mpkCmdQueue, -- mpClmem2, -- CL_TRUE, CL_MAP_READ, 0, -- sizeof(double) * w, 0, nullptr, nullptr, -- &err)); -- if (err != CL_SUCCESS) -- throw OpenCLError("clEnqueueMapBuffer", err, __FILE__, __LINE__); -- -- for (int i = 0; i < w; i++) -- pAllBuffer[i] = resbuf[i]; -- err = clEnqueueUnmapMemObject(kEnv.mpkCmdQueue, mpClmem2, resbuf, 0, nullptr, nullptr); -- if (err != CL_SUCCESS) -- throw OpenCLError("clEnqueueUnmapMemObject", err, __FILE__, __LINE__); -- -- kernelName = Base::GetName() + "_count_reduction"; -- redKernel = clCreateKernel(mpProgram, kernelName.c_str(), &err); -- if (err != CL_SUCCESS) -- throw OpenCLError("clCreateKernel", err, __FILE__, __LINE__); -- SAL_INFO("sc.opencl", "Created kernel " << redKernel << " with name " << kernelName << " in program " << mpProgram); -- -- // 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), -- static_cast<void*>(&buf)); -- if (CL_SUCCESS != err) -- throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); -- -- SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 1 << ": cl_mem: " << mpClmem2); -- err = clSetKernelArg(redKernel, 1, sizeof(cl_mem), &mpClmem2); -- if (CL_SUCCESS != err) -- throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); -- -- SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 2 << ": cl_int: " << nInput); -- err = clSetKernelArg(redKernel, 2, sizeof(cl_int), static_cast<void*>(&nInput)); -- if (CL_SUCCESS != err) -- throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); -- -- SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 3 << ": cl_int: " << nCurWindowSize); -- err = clSetKernelArg(redKernel, 3, sizeof(cl_int), static_cast<void*>(&nCurWindowSize)); -- if (CL_SUCCESS != err) -- throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); -- -- // set work group size and execute -- size_t global_work_size1[] = { 256, static_cast<size_t>(w) }; -- size_t const local_work_size1[] = { 256, 1 }; -- SAL_INFO("sc.opencl", "Enqueing kernel " << redKernel); -- err = clEnqueueNDRangeKernel(kEnv.mpkCmdQueue, redKernel, 2, nullptr, -- global_work_size1, local_work_size1, 0, nullptr, nullptr); -- if (CL_SUCCESS != err) -- throw OpenCLError("clEnqueueNDRangeKernel", err, __FILE__, __LINE__); -- err = clFinish(kEnv.mpkCmdQueue); -- if (CL_SUCCESS != err) -- throw OpenCLError("clFinish", err, __FILE__, __LINE__); -- resbuf = static_cast<double*>(clEnqueueMapBuffer(kEnv.mpkCmdQueue, -- mpClmem2, -- CL_TRUE, CL_MAP_READ, 0, -- sizeof(double) * w, 0, nullptr, nullptr, -- &err)); -- if (err != CL_SUCCESS) -- 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, nullptr, nullptr); -- // FIXME: Is it intentional to not throw an OpenCLError even if the clEnqueueUnmapMemObject() fails? -- if (CL_SUCCESS != err) -- SAL_WARN("sc.opencl", "clEnqueueUnmapMemObject failed: " << openclwrapper::errorString(err)); -- if (mpClmem2) -- { -- err = clReleaseMemObject(mpClmem2); -- SAL_WARN_IF(err != CL_SUCCESS, "sc.opencl", "clReleaseMemObject failed: " << openclwrapper::errorString(err)); -- mpClmem2 = nullptr; -- } -- mpClmem2 = clCreateBuffer(kEnv.mpkContext, -- 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("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), &mpClmem2); -- if (CL_SUCCESS != err) -- throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); -- return 1; -- } -- - ~ParallelReductionVectorRef() - { - if (mpClmem2) -@@ -2324,6 +1959,379 @@ struct SumIfsArgs - cl_mem mCLMem; - double mConst; - }; -+} -+ -+template<class Base> -+void ParallelReductionVectorRef<Base>::GenSlidingWindowFunction( std::stringstream& ss ) -+{ -+ if (!dynamic_cast<OpAverage*>(mpCodeGen.get())) -+ { -+ std::string name = Base::GetName(); -+ ss << "__kernel void " << name; -+ ss << "_reduction(__global double* A, " -+ "__global double *result,int arrayLength,int windowSize){\n"; -+ ss << " double tmp, current_result =" << -+ mpCodeGen->GetBottom(); -+ ss << ";\n"; -+ ss << " int writePos = get_group_id(1);\n"; -+ ss << " int lidx = get_local_id(0);\n"; -+ ss << " __local double shm_buf[256];\n"; -+ if (mpDVR->IsStartFixed()) -+ ss << " int offset = 0;\n"; -+ else // if (!mpDVR->IsStartFixed()) -+ ss << " int offset = get_group_id(1);\n"; -+ if (mpDVR->IsStartFixed() && mpDVR->IsEndFixed()) -+ ss << " int end = windowSize;\n"; -+ else if (!mpDVR->IsStartFixed() && !mpDVR->IsEndFixed()) -+ ss << " int end = offset + windowSize;\n"; -+ else if (mpDVR->IsStartFixed() && !mpDVR->IsEndFixed()) -+ ss << " int end = windowSize + get_group_id(1);\n"; -+ else if (!mpDVR->IsStartFixed() && mpDVR->IsEndFixed()) -+ ss << " int end = windowSize;\n"; -+ ss << " end = min(end, arrayLength);\n"; -+ -+ ss << " barrier(CLK_LOCAL_MEM_FENCE);\n"; -+ ss << " int loop = arrayLength/512 + 1;\n"; -+ ss << " for (int l=0; l<loop; l++){\n"; -+ ss << " tmp = " << mpCodeGen->GetBottom() << ";\n"; -+ ss << " int loopOffset = l*512;\n"; -+ ss << " if((loopOffset + lidx + offset + 256) < end) {\n"; -+ ss << " tmp = legalize(" << mpCodeGen->Gen2( -+ "A[loopOffset + lidx + offset]", "tmp") << ", tmp);\n"; -+ ss << " tmp = legalize(" << mpCodeGen->Gen2( -+ "A[loopOffset + lidx + offset + 256]", "tmp") << ", tmp);\n"; -+ ss << " } else if ((loopOffset + lidx + offset) < end)\n"; -+ ss << " tmp = legalize(" << mpCodeGen->Gen2( -+ "A[loopOffset + lidx + offset]", "tmp") << ", tmp);\n"; -+ ss << " shm_buf[lidx] = tmp;\n"; -+ ss << " barrier(CLK_LOCAL_MEM_FENCE);\n"; -+ ss << " for (int i = 128; i >0; i/=2) {\n"; -+ ss << " if (lidx < i)\n"; -+ ss << " shm_buf[lidx] = "; -+ // Special case count -+ if (dynamic_cast<OpCount*>(mpCodeGen.get())) -+ ss << "shm_buf[lidx] + shm_buf[lidx + i];\n"; -+ else -+ ss << mpCodeGen->Gen2("shm_buf[lidx]", "shm_buf[lidx + i]") << ";\n"; -+ ss << " barrier(CLK_LOCAL_MEM_FENCE);\n"; -+ ss << " }\n"; -+ ss << " if (lidx == 0)\n"; -+ ss << " current_result ="; -+ if (dynamic_cast<OpCount*>(mpCodeGen.get())) -+ ss << "current_result + shm_buf[0]"; -+ else -+ ss << mpCodeGen->Gen2("current_result", "shm_buf[0]"); -+ ss << ";\n"; -+ ss << " barrier(CLK_LOCAL_MEM_FENCE);\n"; -+ ss << " }\n"; -+ ss << " if (lidx == 0)\n"; -+ ss << " result[writePos] = current_result;\n"; -+ ss << "}\n"; -+ } -+ else -+ { -+ std::string name = Base::GetName(); -+ /*sum reduction*/ -+ ss << "__kernel void " << name << "_sum"; -+ ss << "_reduction(__global double* A, " -+ "__global double *result,int arrayLength,int windowSize){\n"; -+ ss << " double tmp, current_result =" << -+ mpCodeGen->GetBottom(); -+ ss << ";\n"; -+ ss << " int writePos = get_group_id(1);\n"; -+ ss << " int lidx = get_local_id(0);\n"; -+ ss << " __local double shm_buf[256];\n"; -+ if (mpDVR->IsStartFixed()) -+ ss << " int offset = 0;\n"; -+ else // if (!mpDVR->IsStartFixed()) -+ ss << " int offset = get_group_id(1);\n"; -+ if (mpDVR->IsStartFixed() && mpDVR->IsEndFixed()) -+ ss << " int end = windowSize;\n"; -+ else if (!mpDVR->IsStartFixed() && !mpDVR->IsEndFixed()) -+ ss << " int end = offset + windowSize;\n"; -+ else if (mpDVR->IsStartFixed() && !mpDVR->IsEndFixed()) -+ ss << " int end = windowSize + get_group_id(1);\n"; -+ else if (!mpDVR->IsStartFixed() && mpDVR->IsEndFixed()) -+ ss << " int end = windowSize;\n"; -+ ss << " end = min(end, arrayLength);\n"; -+ ss << " barrier(CLK_LOCAL_MEM_FENCE);\n"; -+ ss << " int loop = arrayLength/512 + 1;\n"; -+ ss << " for (int l=0; l<loop; l++){\n"; -+ ss << " tmp = " << mpCodeGen->GetBottom() << ";\n"; -+ ss << " int loopOffset = l*512;\n"; -+ ss << " if((loopOffset + lidx + offset + 256) < end) {\n"; -+ ss << " tmp = legalize("; -+ ss << "(A[loopOffset + lidx + offset]+ tmp)"; -+ ss << ", tmp);\n"; -+ ss << " tmp = legalize((A[loopOffset + lidx + offset + 256]+ tmp)"; -+ ss << ", tmp);\n"; -+ ss << " } else if ((loopOffset + lidx + offset) < end)\n"; -+ ss << " tmp = legalize((A[loopOffset + lidx + offset] + tmp)"; -+ ss << ", tmp);\n"; -+ ss << " shm_buf[lidx] = tmp;\n"; -+ ss << " barrier(CLK_LOCAL_MEM_FENCE);\n"; -+ ss << " for (int i = 128; i >0; i/=2) {\n"; -+ ss << " if (lidx < i)\n"; -+ ss << " shm_buf[lidx] = "; -+ ss << "shm_buf[lidx] + shm_buf[lidx + i];\n"; -+ ss << " barrier(CLK_LOCAL_MEM_FENCE);\n"; -+ ss << " }\n"; -+ ss << " if (lidx == 0)\n"; -+ ss << " current_result ="; -+ ss << "current_result + shm_buf[0]"; -+ ss << ";\n"; -+ ss << " barrier(CLK_LOCAL_MEM_FENCE);\n"; -+ ss << " }\n"; -+ ss << " if (lidx == 0)\n"; -+ ss << " result[writePos] = current_result;\n"; -+ ss << "}\n"; -+ /*count reduction*/ -+ ss << "__kernel void " << name << "_count"; -+ ss << "_reduction(__global double* A, " -+ "__global double *result,int arrayLength,int windowSize){\n"; -+ ss << " double tmp, current_result =" << -+ mpCodeGen->GetBottom(); -+ ss << ";\n"; -+ ss << " int writePos = get_group_id(1);\n"; -+ ss << " int lidx = get_local_id(0);\n"; -+ ss << " __local double shm_buf[256];\n"; -+ if (mpDVR->IsStartFixed()) -+ ss << " int offset = 0;\n"; -+ else // if (!mpDVR->IsStartFixed()) -+ ss << " int offset = get_group_id(1);\n"; -+ if (mpDVR->IsStartFixed() && mpDVR->IsEndFixed()) -+ ss << " int end = windowSize;\n"; -+ else if (!mpDVR->IsStartFixed() && !mpDVR->IsEndFixed()) -+ ss << " int end = offset + windowSize;\n"; -+ else if (mpDVR->IsStartFixed() && !mpDVR->IsEndFixed()) -+ ss << " int end = windowSize + get_group_id(1);\n"; -+ else if (!mpDVR->IsStartFixed() && mpDVR->IsEndFixed()) -+ ss << " int end = windowSize;\n"; -+ ss << " end = min(end, arrayLength);\n"; -+ ss << " barrier(CLK_LOCAL_MEM_FENCE);\n"; -+ ss << " int loop = arrayLength/512 + 1;\n"; -+ ss << " for (int l=0; l<loop; l++){\n"; -+ ss << " tmp = " << mpCodeGen->GetBottom() << ";\n"; -+ ss << " int loopOffset = l*512;\n"; -+ ss << " if((loopOffset + lidx + offset + 256) < end) {\n"; -+ ss << " tmp = legalize((isnan(A[loopOffset + lidx + offset])?tmp:tmp+1.0)"; -+ ss << ", tmp);\n"; -+ ss << " tmp = legalize((isnan(A[loopOffset + lidx + offset+256])?tmp:tmp+1.0)"; -+ ss << ", tmp);\n"; -+ ss << " } else if ((loopOffset + lidx + offset) < end)\n"; -+ ss << " tmp = legalize((isnan(A[loopOffset + lidx + offset])?tmp:tmp+1.0)"; -+ ss << ", tmp);\n"; -+ ss << " shm_buf[lidx] = tmp;\n"; -+ ss << " barrier(CLK_LOCAL_MEM_FENCE);\n"; -+ ss << " for (int i = 128; i >0; i/=2) {\n"; -+ ss << " if (lidx < i)\n"; -+ ss << " shm_buf[lidx] = "; -+ ss << "shm_buf[lidx] + shm_buf[lidx + i];\n"; -+ ss << " barrier(CLK_LOCAL_MEM_FENCE);\n"; -+ ss << " }\n"; -+ ss << " if (lidx == 0)\n"; -+ ss << " current_result ="; -+ ss << "current_result + shm_buf[0];"; -+ ss << ";\n"; -+ ss << " barrier(CLK_LOCAL_MEM_FENCE);\n"; -+ ss << " }\n"; -+ ss << " if (lidx == 0)\n"; -+ ss << " result[writePos] = current_result;\n"; -+ ss << "}\n"; -+ } -+ -+} -+ -+template<class Base> -+size_t ParallelReductionVectorRef<Base>::GenReductionLoopHeader( -+ std::stringstream& ss, int nResultSize, bool& needBody ) -+{ -+ assert(mpDVR); -+ size_t nCurWindowSize = mpDVR->GetRefRowSize(); -+ std::string temp = Base::GetName() + "[gid0]"; -+ ss << "tmp = "; -+ // Special case count -+ if (dynamic_cast<OpAverage*>(mpCodeGen.get())) -+ { -+ ss << mpCodeGen->Gen2(temp, "tmp") << ";\n"; -+ ss << "nCount = nCount-1;\n"; -+ ss << "nCount = nCount +"; /*re-assign nCount from count reduction*/ -+ ss << Base::GetName() << "[gid0+" << nResultSize << "]" << ";\n"; -+ } -+ else if (dynamic_cast<OpCount*>(mpCodeGen.get())) -+ ss << temp << "+ tmp"; -+ else -+ ss << mpCodeGen->Gen2(temp, "tmp"); -+ ss << ";\n\t"; -+ needBody = false; -+ return nCurWindowSize; -+} -+ -+template<class Base> -+size_t ParallelReductionVectorRef<Base>::Marshal( cl_kernel k, int argno, int w, cl_program mpProgram ) -+{ -+ assert(Base::mpClmem == nullptr); -+ -+ openclwrapper::KernelEnv kEnv; -+ openclwrapper::setKernelEnv(&kEnv); -+ cl_int err; -+ size_t nInput = mpDVR->GetArrayLength(); -+ size_t nCurWindowSize = mpDVR->GetRefRowSize(); -+ // create clmem buffer -+ if (mpDVR->GetArrays()[Base::mnIndex].mpNumericArray == nullptr) -+ throw Unhandled(__FILE__, __LINE__); -+ double* pHostBuffer = const_cast<double*>( -+ mpDVR->GetArrays()[Base::mnIndex].mpNumericArray); -+ size_t szHostBuffer = nInput * sizeof(double); -+ Base::mpClmem = clCreateBuffer(kEnv.mpkContext, -+ 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 << " using host buffer " << pHostBuffer); -+ -+ mpClmem2 = clCreateBuffer(kEnv.mpkContext, -+ CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, -+ sizeof(double) * w, nullptr, nullptr); -+ if (CL_SUCCESS != err) -+ 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 -+ std::string kernelName; -+ if (!dynamic_cast<OpAverage*>(mpCodeGen.get())) -+ kernelName = Base::GetName() + "_reduction"; -+ else -+ kernelName = Base::GetName() + "_sum_reduction"; -+ cl_kernel redKernel = clCreateKernel(mpProgram, kernelName.c_str(), &err); -+ if (err != CL_SUCCESS) -+ throw OpenCLError("clCreateKernel", err, __FILE__, __LINE__); -+ SAL_INFO("sc.opencl", "Created kernel " << redKernel << " with name " << kernelName << " in program " << mpProgram); -+ -+ // 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), -+ static_cast<void*>(&buf)); -+ if (CL_SUCCESS != err) -+ throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); -+ -+ SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 1 << ": cl_mem: " << mpClmem2); -+ err = clSetKernelArg(redKernel, 1, sizeof(cl_mem), &mpClmem2); -+ if (CL_SUCCESS != err) -+ throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); -+ -+ SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 2 << ": cl_int: " << nInput); -+ err = clSetKernelArg(redKernel, 2, sizeof(cl_int), static_cast<void*>(&nInput)); -+ if (CL_SUCCESS != err) -+ throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); -+ -+ SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 3 << ": cl_int: " << nCurWindowSize); -+ err = clSetKernelArg(redKernel, 3, sizeof(cl_int), static_cast<void*>(&nCurWindowSize)); -+ if (CL_SUCCESS != err) -+ throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); -+ -+ // set work group size and execute -+ size_t global_work_size[] = { 256, static_cast<size_t>(w) }; -+ size_t const local_work_size[] = { 256, 1 }; -+ SAL_INFO("sc.opencl", "Enqueing kernel " << redKernel); -+ err = clEnqueueNDRangeKernel(kEnv.mpkCmdQueue, redKernel, 2, nullptr, -+ global_work_size, local_work_size, 0, nullptr, nullptr); -+ if (CL_SUCCESS != err) -+ throw OpenCLError("clEnqueueNDRangeKernel", err, __FILE__, __LINE__); -+ err = clFinish(kEnv.mpkCmdQueue); -+ if (CL_SUCCESS != err) -+ throw OpenCLError("clFinish", err, __FILE__, __LINE__); -+ if (dynamic_cast<OpAverage*>(mpCodeGen.get())) -+ { -+ /*average need more reduction kernel for count computing*/ -+ std::unique_ptr<double[]> pAllBuffer(new double[2 * w]); -+ double* resbuf = static_cast<double*>(clEnqueueMapBuffer(kEnv.mpkCmdQueue, -+ mpClmem2, -+ CL_TRUE, CL_MAP_READ, 0, -+ sizeof(double) * w, 0, nullptr, nullptr, -+ &err)); -+ if (err != CL_SUCCESS) -+ throw OpenCLError("clEnqueueMapBuffer", err, __FILE__, __LINE__); -+ -+ for (int i = 0; i < w; i++) -+ pAllBuffer[i] = resbuf[i]; -+ err = clEnqueueUnmapMemObject(kEnv.mpkCmdQueue, mpClmem2, resbuf, 0, nullptr, nullptr); -+ if (err != CL_SUCCESS) -+ throw OpenCLError("clEnqueueUnmapMemObject", err, __FILE__, __LINE__); -+ -+ kernelName = Base::GetName() + "_count_reduction"; -+ redKernel = clCreateKernel(mpProgram, kernelName.c_str(), &err); -+ if (err != CL_SUCCESS) -+ throw OpenCLError("clCreateKernel", err, __FILE__, __LINE__); -+ SAL_INFO("sc.opencl", "Created kernel " << redKernel << " with name " << kernelName << " in program " << mpProgram); -+ -+ // 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), -+ static_cast<void*>(&buf)); -+ if (CL_SUCCESS != err) -+ throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); -+ -+ SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 1 << ": cl_mem: " << mpClmem2); -+ err = clSetKernelArg(redKernel, 1, sizeof(cl_mem), &mpClmem2); -+ if (CL_SUCCESS != err) -+ throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); -+ -+ SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 2 << ": cl_int: " << nInput); -+ err = clSetKernelArg(redKernel, 2, sizeof(cl_int), static_cast<void*>(&nInput)); -+ if (CL_SUCCESS != err) -+ throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); -+ -+ SAL_INFO("sc.opencl", "Kernel " << redKernel << " arg " << 3 << ": cl_int: " << nCurWindowSize); -+ err = clSetKernelArg(redKernel, 3, sizeof(cl_int), static_cast<void*>(&nCurWindowSize)); -+ if (CL_SUCCESS != err) -+ throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); -+ -+ // set work group size and execute -+ size_t global_work_size1[] = { 256, static_cast<size_t>(w) }; -+ size_t const local_work_size1[] = { 256, 1 }; -+ SAL_INFO("sc.opencl", "Enqueing kernel " << redKernel); -+ err = clEnqueueNDRangeKernel(kEnv.mpkCmdQueue, redKernel, 2, nullptr, -+ global_work_size1, local_work_size1, 0, nullptr, nullptr); -+ if (CL_SUCCESS != err) -+ throw OpenCLError("clEnqueueNDRangeKernel", err, __FILE__, __LINE__); -+ err = clFinish(kEnv.mpkCmdQueue); -+ if (CL_SUCCESS != err) -+ throw OpenCLError("clFinish", err, __FILE__, __LINE__); -+ resbuf = static_cast<double*>(clEnqueueMapBuffer(kEnv.mpkCmdQueue, -+ mpClmem2, -+ CL_TRUE, CL_MAP_READ, 0, -+ sizeof(double) * w, 0, nullptr, nullptr, -+ &err)); -+ if (err != CL_SUCCESS) -+ 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, nullptr, nullptr); -+ // FIXME: Is it intentional to not throw an OpenCLError even if the clEnqueueUnmapMemObject() fails? -+ if (CL_SUCCESS != err) -+ SAL_WARN("sc.opencl", "clEnqueueUnmapMemObject failed: " << openclwrapper::errorString(err)); -+ if (mpClmem2) -+ { -+ err = clReleaseMemObject(mpClmem2); -+ SAL_WARN_IF(err != CL_SUCCESS, "sc.opencl", "clReleaseMemObject failed: " << openclwrapper::errorString(err)); -+ mpClmem2 = nullptr; -+ } -+ mpClmem2 = clCreateBuffer(kEnv.mpkContext, -+ 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("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), &mpClmem2); -+ if (CL_SUCCESS != err) -+ throw OpenCLError("clSetKernelArg", err, __FILE__, __LINE__); -+ return 1; - } - - /// Helper functions that have multiple buffers diff --git a/editors/libreoffice/files/patch-sdext_source_pdfimport_xpdfwrapper_pdfioutdev__gpl.cxx b/editors/libreoffice/files/patch-sdext_source_pdfimport_xpdfwrapper_pdfioutdev__gpl.cxx deleted file mode 100644 index 5c9586c81662..000000000000 --- a/editors/libreoffice/files/patch-sdext_source_pdfimport_xpdfwrapper_pdfioutdev__gpl.cxx +++ /dev/null @@ -1,18 +0,0 @@ -From c1adc0a7559340213d754db950fd68bfe9e5a651 Mon Sep 17 00:00:00 2001 -From: Martin Whitaker <foss@martin-whitaker.me.uk> -Date: Sun, 15 Mar 2020 13:14:34 +0000 -Subject: [PATCH] Fix build with poppler 0.86.0. - ---- sdext/source/pdfimport/xpdfwrapper/pdfioutdev_gpl.cxx.orig 2020-03-11 16:18:35 UTC -+++ sdext/source/pdfimport/xpdfwrapper/pdfioutdev_gpl.cxx -@@ -563,7 +563,9 @@ void PDFOutDev::processLink(Link* link, Catalog*) - if (!(pAction && pAction->getKind() == actionURI)) - return; - --#if POPPLER_CHECK_VERSION(0, 72, 0) -+#if POPPLER_CHECK_VERSION(0, 86, 0) -+ const char* pURI = static_cast<LinkURI*>(pAction)->getURI().c_str(); -+#elif POPPLER_CHECK_VERSION(0, 72, 0) - const char* pURI = static_cast<LinkURI*>(pAction)->getURI()->c_str(); - #else - const char* pURI = static_cast<LinkURI*>(pAction)->getURI()->getCString(); diff --git a/editors/libreoffice/files/patch-vcl_qt5_Qt5Frame.cxx b/editors/libreoffice/files/patch-vcl_qt5_Qt5Frame.cxx new file mode 100644 index 000000000000..43fc420c44d4 --- /dev/null +++ b/editors/libreoffice/files/patch-vcl_qt5_Qt5Frame.cxx @@ -0,0 +1,11 @@ +--- vcl/qt5/Qt5Frame.cxx.orig 2020-05-13 11:19:20 UTC ++++ vcl/qt5/Qt5Frame.cxx +@@ -318,7 +318,7 @@ SalGraphics* Qt5Frame::AcquireGraphics() + + if (m_bUseCairo) + { +- if (!m_pOurSvpGraphics.get() || m_bGraphicsInvalid) ++ if (!m_pOurSvpGraphics.get()) + { + m_pOurSvpGraphics.reset(new Qt5SvpGraphics(this)); + InitQt5SvpGraphics(m_pOurSvpGraphics.get()); diff --git a/editors/libreoffice/files/patch-vcl_qt5_Qt5Instance.cxx b/editors/libreoffice/files/patch-vcl_qt5_Qt5Instance.cxx new file mode 100644 index 000000000000..aad13b769fb4 --- /dev/null +++ b/editors/libreoffice/files/patch-vcl_qt5_Qt5Instance.cxx @@ -0,0 +1,35 @@ +--- vcl/qt5/Qt5Instance.cxx.orig 2020-05-13 11:19:20 UTC ++++ vcl/qt5/Qt5Instance.cxx +@@ -261,7 +261,13 @@ SalFrame* Qt5Instance::CreateChildFrame(SystemParentDa + SalFrame* Qt5Instance::CreateFrame(SalFrame* pParent, SalFrameStyleFlags nStyle) + { + assert(!pParent || dynamic_cast<Qt5Frame*>(pParent)); +- return new Qt5Frame(static_cast<Qt5Frame*>(pParent), nStyle, m_bUseCairo); ++ SalFrame* pRet(nullptr); ++ bool bUseCairo = m_bUseCairo; ++ RunInMainThread([&pRet, pParent, nStyle, bUseCairo]() { ++ pRet = new Qt5Frame(static_cast<Qt5Frame*>(pParent), nStyle, bUseCairo); ++ }); ++ assert(pRet); ++ return pRet; + } + + void Qt5Instance::DestroyFrame(SalFrame* pFrame) +@@ -420,7 +426,7 @@ Qt5Instance::createPicker(css::uno::Reference<css::uno + { + SolarMutexGuard g; + Qt5FilePicker* pPicker; +- RunInMainThread([&, this]() { pPicker = createPicker(context, eMode); }); ++ RunInMainThread([&pPicker, this, context, eMode]() { pPicker = createPicker(context, eMode); }); + assert(pPicker); + return pPicker; + } +@@ -624,7 +630,7 @@ std::unique_ptr<QApplication> Qt5Instance::CreateQAppl + extern "C" { + VCLPLUG_QT5_PUBLIC SalInstance* create_SalInstance() + { +- static const bool bUseCairo = (nullptr != getenv("SAL_VCL_QT5_USE_CAIRO")); ++ static const bool bUseCairo = true; // (nullptr != getenv("SAL_VCL_QT5_USE_CAIRO")); + + std::unique_ptr<char* []> pFakeArgv; + std::unique_ptr<int> pFakeArgc; |