Source-Changes-HG archive
[Date Prev][Date Next][Thread Prev][Thread Next][Date Index][Thread Index][Old Index]
[pkgsrc/trunk]: pkgsrc/misc/libreoffice Declare classes before dyncasting them
details: https://anonhg.NetBSD.org/pkgsrc/rev/9c94646daa88
branches: trunk
changeset: 429456:9c94646daa88
user: joerg <joerg%pkgsrc.org@localhost>
date: Mon Apr 20 13:00:48 2020 +0000
description:
Declare classes before dyncasting them
diffstat:
misc/libreoffice/distinfo | 3 +-
misc/libreoffice/patches/patch-sc_source_core_opencl_formulagroupcl.cxx | 798 ++++++++++
2 files changed, 800 insertions(+), 1 deletions(-)
diffs (truncated from 819 to 300 lines):
diff -r 5c21a4ef9da7 -r 9c94646daa88 misc/libreoffice/distinfo
--- a/misc/libreoffice/distinfo Mon Apr 20 12:50:02 2020 +0000
+++ b/misc/libreoffice/distinfo Mon Apr 20 13:00:48 2020 +0000
@@ -1,4 +1,4 @@
-$NetBSD: distinfo,v 1.97 2020/04/18 04:18:05 ryoon Exp $
+$NetBSD: distinfo,v 1.98 2020/04/20 13:00:48 joerg Exp $
SHA1 (libreoffice/0168229624cfac409e766913506961a8-ucpp-1.3.2.tar.gz) = 452eba922e4f41603539c9dc39947d2271e47093
RMD160 (libreoffice/0168229624cfac409e766913506961a8-ucpp-1.3.2.tar.gz) = dbeb7a7f8c89961ca2e544b810345d025561866b
@@ -235,6 +235,7 @@
SHA1 (patch-i18nutil_source_utility_unicode.cxx) = 014d48574a379f92f2d0afe3c91d2c2956708dd8
SHA1 (patch-sal_osl_unx_system.hxx) = 478dd47dc512d283a554600951010bbaa10c2c0e
SHA1 (patch-sal_rtl_cipher.cxx) = e5e46d0b96e25572be86bf26f0053436859736cd
+SHA1 (patch-sc_source_core_opencl_formulagroupcl.cxx) = 1a589c8190a28f603671faf1d84cac022425849e
SHA1 (patch-sc_source_core_tool_math.cxx) = dd665c6bbe3c18fca0b5fbf847a6ec1ff50d0bb9
SHA1 (patch-solenv_gbuild_Module.mk) = cbd6ca3acae187458e49fe76d973e6475ed5fe1d
SHA1 (patch-solenv_gbuild_platform_NETBSD_AARCH64_GCC.mk) = eda8dbd0e9394b2dde2f98c0df39e7e9888f49dc
diff -r 5c21a4ef9da7 -r 9c94646daa88 misc/libreoffice/patches/patch-sc_source_core_opencl_formulagroupcl.cxx
--- /dev/null Thu Jan 01 00:00:00 1970 +0000
+++ b/misc/libreoffice/patches/patch-sc_source_core_opencl_formulagroupcl.cxx Mon Apr 20 13:00:48 2020 +0000
@@ -0,0 +1,798 @@
+$NetBSD: patch-sc_source_core_opencl_formulagroupcl.cxx,v 1.1 2020/04/20 13:00:48 joerg Exp $
+
+--- sc/source/core/opencl/formulagroupcl.cxx.orig 2020-04-20 00:48:10.479759827 +0000
++++ sc/source/core/opencl/formulagroupcl.cxx
+@@ -1335,185 +1335,7 @@ 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";
+-
+- 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 void GenSlidingWindowFunction( std::stringstream& ss );
+
+ virtual std::string GenSlidingWindowDeclRef( bool ) const
+ {
+@@ -1527,213 +1349,28 @@ public:
+
+ /// Controls how the elements in the DoubleVectorRef are traversed
+ size_t GenReductionLoopHeader(
+- std::stringstream& ss, int nResultSize, bool& needBody )
++ std::stringstream& ss, int nResultSize, bool& needBody );
++
++ virtual size_t Marshal( cl_kernel k, int argno, int w, cl_program mpProgram );
++
++ ~ParallelReductionVectorRef()
+ {
+- assert(mpDVR);
+- size_t nCurWindowSize = mpDVR->GetRefRowSize();
+- std::string temp = Base::GetName() + "[gid0]";
+- ss << "tmp = ";
+- // Special case count
+- if (dynamic_cast<OpAverage*>(mpCodeGen.get()))
++ if (mpClmem2)
+ {
+- 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";
++ cl_int err;
++ err = clReleaseMemObject(mpClmem2);
++ SAL_WARN_IF(err != CL_SUCCESS, "sc.opencl", "clReleaseMemObject failed: " << openclwrapper::errorString(err));
++ mpClmem2 = nullptr;
+ }
+- else if (dynamic_cast<OpCount*>(mpCodeGen.get()))
+- ss << temp << "+ tmp";
+- else
+- ss << mpCodeGen->Gen2(temp, "tmp");
+- ss << ";\n\t";
+- needBody = false;
+- return nCurWindowSize;
+ }
+
+- virtual size_t 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));
++ size_t GetArrayLength() const { return mpDVR->GetArrayLength(); }
+
+- // 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);
++ size_t GetWindowSize() const { return mpDVR->GetRefRowSize(); }
+
+- // 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__);
++ bool GetStartFixed() const { return bIsStartFixed; }
+
Home |
Main Index |
Thread Index |
Old Index