aboutsummaryrefslogtreecommitdiff
path: root/editors/libreoffice
diff options
context:
space:
mode:
authorDima Panov <fluffy@FreeBSD.org>2020-07-07 03:14:58 +0000
committerDima Panov <fluffy@FreeBSD.org>2020-07-07 03:14:58 +0000
commit2d98f1c14da2bc2da0ea8e65c4a5b93b4ab49d4d (patch)
tree16394ca8b04abdbe1393df7041d1ffa359e32354 /editors/libreoffice
parent22f2f91d5c019a5a8f9843636439c055741db9ce (diff)
downloadports-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/Makefile1
-rw-r--r--editors/libreoffice/Makefile.common2
-rw-r--r--editors/libreoffice/distinfo10
-rw-r--r--editors/libreoffice/files/patch-sc_source_core_opencl_formulagroupcl.cxx780
-rw-r--r--editors/libreoffice/files/patch-sdext_source_pdfimport_xpdfwrapper_pdfioutdev__gpl.cxx18
-rw-r--r--editors/libreoffice/files/patch-vcl_qt5_Qt5Frame.cxx11
-rw-r--r--editors/libreoffice/files/patch-vcl_qt5_Qt5Instance.cxx35
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;