test in a separate helper process if OpenCL crashes (tdf#112252)

Some OpenCL implementations may be broken, e.g. pocl simply
asserts and aborts if it can't find Clang. In order to protect
against crashes caused by faulty OpenCL drivers, when testing OpenCL
functionality on OpenCL setup change, first do a simple test
in a separate helper.

Change-Id: I1cf328e731c48f47745b27c7130e7521254209f5
Reviewed-on: https://gerrit.libreoffice.org/71080
Tested-by: Jenkins
Reviewed-by: Luboš Luňák <l.lunak@collabora.com>

8 files changed, 328 insertions, 1 deletions

diff --git a/Repository.mk b/Repository.mk
index b97a68e371e5..d9654e717025 100644
--- a/Repository.mk
+++ b/Repository.mk
@@ -209,6 +209,7 @@ $(eval $(call gb_Helper_register_executables_for_install,OOO,ooo, \
 	$(if $(filter WNT,$(OS)), \
 		senddoc \
 	) \
+	$(if $(filter OPENCL,$(BUILD_TYPE)),opencltest) \
 ))
 
 ifeq ($(OS),WNT)
diff --git a/desktop/source/app/opencl.cxx b/desktop/source/app/opencl.cxx
index f62c93e32e1f..a728ba2b8b7d 100644
--- a/desktop/source/app/opencl.cxx
+++ b/desktop/source/app/opencl.cxx
@@ -38,6 +38,7 @@
 #include <opencl/OpenCLZone.hxx>
 
 #include <osl/file.hxx>
+#include <osl/process.h>
 
 using namespace ::osl;
 using namespace ::com::sun::star::uno;
@@ -47,6 +48,61 @@ namespace desktop {
 
 #if HAVE_FEATURE_OPENCL
 
+static bool testOpenCLDriver()
+{
+    // A simple OpenCL test run in a separate process in order to test
+    // whether the driver crashes (asserts,etc.) when trying to use OpenCL.
+    SAL_INFO("opencl", "Starting CL driver test");
+
+    OUString testerURL("$BRAND_BASE_DIR/" LIBO_BIN_FOLDER "/opencltest");
+    rtl::Bootstrap::expandMacros(testerURL); //TODO: detect failure
+
+    OUString deviceName, platformName;
+    openclwrapper::getOpenCLDeviceName( deviceName, platformName );
+    rtl_uString* args[] = { deviceName.pData, platformName.pData };
+    sal_Int32 numArgs = 2;
+
+    oslProcess process;
+    oslSecurity security = osl_getCurrentSecurity();
+    oslProcessError error = osl_executeProcess(testerURL.pData, args, numArgs,
+        osl_Process_SEARCHPATH | osl_Process_HIDDEN, security,
+        nullptr, nullptr, 0, &process );
+    osl_freeSecurityHandle( security );
+    if( error != osl_Process_E_None )
+    {
+        SAL_WARN( "opencl", "failed to start CL driver test: " << error );
+        return false;
+    }
+    // If the driver takes more than 10 seconds, it's probably broken/useless.
+    TimeValue timeout( 10, 0 );
+    error = osl_joinProcessWithTimeout( process, &timeout );
+    if( error == osl_Process_E_None )
+    {
+        oslProcessInfo info;
+        info.Size = sizeof( info );
+        error = osl_getProcessInfo( process, osl_Process_EXITCODE, &info );
+        if( error == osl_Process_E_None )
+        {
+            if( info.Code == 0 )
+            {
+                SAL_INFO( "opencl", "CL driver test passed" );
+                osl_freeProcessHandle( process );
+                return true;
+            }
+            else
+            {
+                SAL_WARN( "opencl", "CL driver test failed - disabling: " << info.Code );
+                osl_freeProcessHandle( process );
+                return false;
+            }
+        }
+    }
+    SAL_WARN( "opencl", "CL driver test did not finish - disabling: " << error );
+    osl_terminateProcess( process );
+    osl_freeProcessHandle( process );
+    return false;
+}
+
 static bool testOpenCLCompute(const Reference< XDesktop2 > &xDesktop, const OUString &rURL)
 {
     bool bSuccess = false;
@@ -178,7 +234,10 @@ void Desktop::CheckOpenCLCompute(const Reference< XDesktop2 > &xDesktop)
             xBatch->commit();
         }
 
-        bool bSucceeded = testOpenCLCompute(xDesktop, aURL);
+        // Hopefully at least basic functionality always works and broken OpenCL implementations break
+        // only when they are used to compute something. If this assumptions turns out to be not true,
+        // the driver check needs to be moved sooner.
+        bool bSucceeded = testOpenCLDriver() && testOpenCLCompute(xDesktop, aURL);
 
         { // restore the minimum group size
             std::shared_ptr<comphelper::ConfigurationChanges> xBatch(comphelper::ConfigurationChanges::create());
diff --git a/include/opencl/openclwrapper.hxx b/include/opencl/openclwrapper.hxx
index 67dfc8798416..81138ced24e8 100644
--- a/include/opencl/openclwrapper.hxx
+++ b/include/opencl/openclwrapper.hxx
@@ -77,6 +77,8 @@ OPENCL_DLLPUBLIC bool switchOpenCLDevice(const OUString* pDeviceId, bool bAutoSe
 
 OPENCL_DLLPUBLIC void getOpenCLDeviceInfo(size_t& rDeviceId, size_t& rPlatformId);
 
+OPENCL_DLLPUBLIC void getOpenCLDeviceName(OUString& rDeviceName, OUString& rPlatformName);
+
 /**
  * Set the current command queue position in case of multiple command queues
  * for a given device.
diff --git a/opencl/Executable_opencltest.mk b/opencl/Executable_opencltest.mk
new file mode 100644
index 000000000000..4a8ef280e8e1
--- /dev/null
+++ b/opencl/Executable_opencltest.mk
@@ -0,0 +1,30 @@
+# -*- Mode: makefile-gmake; tab-width: 4; indent-tabs-mode: t -*-
+#
+# This file is part of the LibreOffice project.
+#
+# This Source Code Form is subject to the terms of the Mozilla Public
+# License, v. 2.0. If a copy of the MPL was not distributed with this
+# file, You can obtain one at http://mozilla.org/MPL/2.0/.
+#
+
+$(eval $(call gb_Executable_Executable,opencltest))
+
+$(eval $(call gb_Executable_set_include,opencltest,\
+    -I$(SRCDIR)/opencl/inc \
+    $$(INCLUDE) \
+))
+
+
+$(eval $(call gb_Executable_add_exception_objects,opencltest,\
+    opencl/opencltest/main \
+))
+
+$(eval $(call gb_Executable_use_externals,opencltest,\
+    clew \
+))
+
+$(eval $(call gb_Executable_use_libraries,opencltest,\
+    sal \
+))
+
+# vim: set noet sw=4 ts=4:
diff --git a/opencl/Module_opencl.mk b/opencl/Module_opencl.mk
index 92a80160a1bb..5db04e530788 100644
--- a/opencl/Module_opencl.mk
+++ b/opencl/Module_opencl.mk
@@ -10,6 +10,7 @@
 $(eval $(call gb_Module_Module,opencl))
 
 $(eval $(call gb_Module_add_targets,opencl,\
+    Executable_opencltest \
     Library_opencl \
 ))
 
diff --git a/opencl/inc/clew_setup.hxx b/opencl/inc/clew_setup.hxx
new file mode 100644
index 000000000000..58571faad463
--- /dev/null
+++ b/opencl/inc/clew_setup.hxx
@@ -0,0 +1,25 @@
+/* -*- Mode: C++; tab-width: 4; indent-tabs-mode: nil; c-basic-offset: 4 -*- */
+/*
+ * This file is part of the LibreOffice project.
+ *
+ * This Source Code Form is subject to the terms of the Mozilla Public
+ * License, v. 2.0. If a copy of the MPL was not distributed with this
+ * file, You can obtain one at http://mozilla.org/MPL/2.0/.
+ */
+
+#ifndef INCLUDED_OPENCL_INC_CLEW_SETUP_HXX
+#define INCLUDED_OPENCL_INC_CLEW_SETUP_HXX
+
+#ifdef _WIN32
+#include <prewin.h>
+#include <postwin.h>
+#define OPENCL_DLL_NAME "OpenCL.dll"
+#elif defined(MACOSX)
+#define OPENCL_DLL_NAME nullptr
+#else
+#define OPENCL_DLL_NAME "libOpenCL.so.1"
+#endif
+
+#endif
+
+/* vim:set shiftwidth=4 softtabstop=4 expandtab: */
diff --git a/opencl/opencltest/main.cxx b/opencl/opencltest/main.cxx
new file mode 100644
index 000000000000..0b1292e3e7a3
--- /dev/null
+++ b/opencl/opencltest/main.cxx
@@ -0,0 +1,185 @@
+/* -*- Mode: C++; tab-width: 4; indent-tabs-mode: nil; c-basic-offset: 4 -*- */
+/*
+ * This file is part of the LibreOffice project.
+ *
+ * This Source Code Form is subject to the terms of the Mozilla Public
+ * License, v. 2.0. If a copy of the MPL was not distributed with this
+ * file, You can obtain one at http://mozilla.org/MPL/2.0/.
+ */
+
+#include <clew/clew.h>
+
+#include <vector>
+#include <cassert>
+#include <cstring>
+#include <iostream>
+
+#include <sal/log.hxx>
+
+#include <clew_setup.hxx>
+
+using namespace std;
+
+// The purpose of this separate executable is to check whether OpenCL works
+// without crashing (asserting, etc.). Other checks can be done by LO core itself.
+
+#define check(value, expected)                                                                     \
+    do                                                                                             \
+    {                                                                                              \
+        auto val = (value);                                                                        \
+        if (val != (expected))                                                                     \
+        {                                                                                          \
+            SAL_WARN("opencl",                                                                     \
+                     "OpenCL driver check failed: " << val << "(line " << __LINE__ << ")");        \
+            return;                                                                                \
+        }                                                                                          \
+    } while (false);
+#define openclcheck(value) check(value, CL_SUCCESS)
+
+static void runTest(const char* deviceName, const char* devicePlatform)
+{
+    int status = clewInit(OPENCL_DLL_NAME);
+    check(status, CLEW_SUCCESS);
+
+    // Find the given OpenCL device (in order to use the same one as LO core).
+    cl_uint numPlatforms;
+    openclcheck(clGetPlatformIDs(0, nullptr, &numPlatforms));
+    vector<cl_platform_id> platforms(numPlatforms);
+    openclcheck(clGetPlatformIDs(numPlatforms, platforms.data(), nullptr));
+    cl_platform_id platformId = nullptr;
+    for (cl_uint i = 0; i < numPlatforms; ++i)
+    {
+        char platformName[64];
+        if (clGetPlatformInfo(platforms[i], CL_PLATFORM_NAME, 64, platformName, nullptr)
+                == CL_SUCCESS
+            && strcmp(devicePlatform, platformName) == 0)
+        {
+            platformId = platforms[i];
+            break;
+        }
+    }
+    if (platformId == nullptr)
+    {
+        SAL_WARN("opencl", "Device platform not found: " << devicePlatform);
+        assert(false);
+        return;
+    }
+
+    cl_uint numDevices;
+    openclcheck(clGetDeviceIDs(platformId, CL_DEVICE_TYPE_ALL, 0, nullptr, &numDevices));
+    vector<cl_device_id> devices(numDevices);
+    openclcheck(
+        clGetDeviceIDs(platformId, CL_DEVICE_TYPE_ALL, numDevices, devices.data(), nullptr));
+    cl_device_id deviceId = nullptr;
+    for (cl_uint i = 0; i < numDevices; ++i)
+    {
+        char name[1024];
+        if (clGetDeviceInfo(devices[i], CL_DEVICE_NAME, 64, name, nullptr) == CL_SUCCESS
+            && strcmp(deviceName, name) == 0)
+        {
+            deviceId = devices[i];
+            break;
+        }
+    }
+    if (deviceId == nullptr)
+    {
+        SAL_WARN("opencl", "Device not found: " << deviceName);
+        assert(false);
+        return;
+    }
+
+    cl_context context;
+    cl_int state;
+    cl_context_properties cps[3];
+    cps[0] = CL_CONTEXT_PLATFORM;
+    cps[1] = reinterpret_cast<cl_context_properties>(platformId);
+    cps[2] = 0;
+    context = clCreateContext(cps, 1, &deviceId, nullptr, nullptr, &state);
+    openclcheck(state);
+    cl_command_queue queue = clCreateCommandQueue(context, deviceId, 0, &state);
+    openclcheck(state);
+
+    // Just a simple OpenCL program, the functionality or results do not really matter.
+    const char* source[] = { "__kernel void testFunction( __global float* input1, __global float* "
+                             "input2, __global float* output )"
+                             "{"
+                             "    int gid0 = get_global_id( 0 );"
+                             "    output[ gid0 ] = input1[ gid0 ] * input2[ gid0 ];"
+                             "}" };
+    size_t sourceSize[] = { strlen(source[0]) };
+    cl_program program = clCreateProgramWithSource(context, 1, source, sourceSize, &state);
+    openclcheck(state);
+    state = clBuildProgram(program, 1, &deviceId, nullptr, nullptr, nullptr);
+    if (state != CL_SUCCESS)
+    {
+#ifdef DBG_UTIL
+        size_t length;
+        status
+            = clGetProgramBuildInfo(program, deviceId, CL_PROGRAM_BUILD_LOG, 0, nullptr, &length);
+        vector<char> error(length + 1);
+        status = clGetProgramBuildInfo(program, deviceId, CL_PROGRAM_BUILD_LOG, length,
+                                       error.data(), nullptr);
+        error[length] = '\0';
+        cerr << "OpenCL driver check build error:" << error.data() << endl;
+        abort();
+#else
+        openclcheck(state);
+#endif
+    }
+    cl_kernel kernel = clCreateKernel(program, "testFunction", &state);
+    openclcheck(state);
+
+    // Some random data for the program.
+    constexpr int dataSize = 1000;
+    cl_float inputData1[dataSize];
+    cl_float inputData2[dataSize];
+    cl_float outputData[dataSize];
+    for (int i = 0; i < dataSize; ++i)
+    {
+        inputData1[i] = i * 2;
+        inputData2[i] = i % 100;
+    }
+    cl_mem input1 = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
+                                   sizeof(cl_float) * dataSize, inputData1, &state);
+    openclcheck(state);
+    cl_mem input2 = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
+                                   sizeof(cl_float) * dataSize, inputData2, &state);
+    openclcheck(state);
+    cl_mem output = clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR,
+                                   sizeof(cl_float) * dataSize, outputData, &state);
+    openclcheck(state);
+    state = clSetKernelArg(kernel, 0, sizeof(cl_mem), &input1);
+    openclcheck(state);
+    state = clSetKernelArg(kernel, 1, sizeof(cl_mem), &input2);
+    openclcheck(state);
+    state = clSetKernelArg(kernel, 2, sizeof(cl_mem), &output);
+    openclcheck(state);
+
+    const size_t globalWorkSize[] = { dataSize };
+    const size_t localSize[1] = { 64 };
+    state = clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, globalWorkSize, localSize, 0, nullptr,
+                                   nullptr);
+    openclcheck(state);
+    openclcheck(clFinish(queue));
+    openclcheck(clEnqueueReadBuffer(queue, output, CL_TRUE, 0, sizeof(cl_float) * dataSize,
+                                    outputData, 0, nullptr, nullptr));
+    clReleaseMemObject(input1);
+    clReleaseMemObject(input2);
+    clReleaseMemObject(output);
+    clReleaseKernel(kernel);
+    clReleaseProgram(program);
+    clReleaseCommandQueue(queue);
+    clReleaseContext(context);
+}
+
+int main(int argc, char* argv[])
+{
+    if (argc != 3)
+        return 1;
+    runTest(argv[1], argv[2]);
+    // Always return exit code 0, LO itself can do error checking better, we just care
+    // if this helper crashes or not.
+    return 0;
+}
+
+/* vim:set shiftwidth=4 softtabstop=4 expandtab: */
diff --git a/opencl/source/openclwrapper.cxx b/opencl/source/openclwrapper.cxx
index 1194758f4f4d..26a5e6b7f1b8 100644
--- a/opencl/source/openclwrapper.cxx
+++ b/opencl/source/openclwrapper.cxx
@@ -856,6 +856,30 @@ void getOpenCLDeviceInfo(size_t& rDeviceId, size_t& rPlatformId)
     findDeviceInfoFromDeviceId(id, rDeviceId, rPlatformId);
 }
 
+void getOpenCLDeviceName(OUString& rDeviceName, OUString& rPlatformName)
+{
+    if (!canUseOpenCL())
+        return;
+
+    int status = clewInit(OPENCL_DLL_NAME);
+    if (status < 0)
+        return;
+
+    cl_device_id deviceId = gpuEnv.mpDevID;
+    cl_platform_id platformId;
+    if( clGetDeviceInfo(deviceId, CL_DEVICE_PLATFORM, sizeof(platformId), &platformId, nullptr) != CL_SUCCESS )
+        return;
+
+    char deviceName[DEVICE_NAME_LENGTH] = {0};
+    if( clGetDeviceInfo(deviceId, CL_DEVICE_NAME, sizeof(deviceName), deviceName, nullptr) != CL_SUCCESS )
+        return;
+    char platformName[64];
+    if( clGetPlatformInfo(platformId, CL_PLATFORM_NAME, 64, platformName, nullptr) != CL_SUCCESS )
+        return;
+    rDeviceName = OUString::createFromAscii(deviceName);
+    rPlatformName = OUString::createFromAscii(platformName);
+}
+
 void setOpenCLCmdQueuePosition( int nPos )
 {
     if (nPos < 0 || nPos >= OPENCL_CMDQUEUE_SIZE)