summaryrefslogtreecommitdiff
path: root/sc
diff options
context:
space:
mode:
authorTor Lillqvist <tml@collabora.com>2014-11-27 15:13:12 +0200
committerTor Lillqvist <tml@collabora.com>2014-11-28 10:24:20 +0200
commit87c18d6f4cb9bcb8c4c9e2c80e4d779f01675ff6 (patch)
treeb4e1442dfe3c2b45b93113f7f4ba653e9c2647a3 /sc
parentd4f63f5e767d23675240e9dd6d95fc9496ad9df5 (diff)
Move more Calc-independent OpenCL stuff from the sc to the opencl module
No cleanups yet. Just removed the "sc" namespace parts now when this stuff is no longer Calc-specific. There is still horribly confusing use of the same OpenCLDevice name for both a class and as a namespace, for instance. And the OpenCLDevice class has only public static members even, so effectively it acts as just a namespace anyway... Etc. Change-Id: Idc5f30a721df0101426c676f04a85e02c5dc8443
Diffstat (limited to 'sc')
-rw-r--r--sc/Library_sc.mk2
-rw-r--r--sc/source/core/inc/openclwrapper.hxx108
-rw-r--r--sc/source/core/opencl/formulagroupcl.cxx42
-rw-r--r--sc/source/core/opencl/opencl_device.cxx597
-rw-r--r--sc/source/core/opencl/opencl_device.hxx26
-rw-r--r--sc/source/core/opencl/opencl_device_selection.h642
-rw-r--r--sc/source/core/opencl/openclwrapper.cxx806
-rw-r--r--sc/source/core/tool/formulagroup.cxx8
8 files changed, 25 insertions, 2206 deletions
diff --git a/sc/Library_sc.mk b/sc/Library_sc.mk
index 51ad36ea6004..7c9aab0da51e 100644
--- a/sc/Library_sc.mk
+++ b/sc/Library_sc.mk
@@ -663,8 +663,6 @@ $(eval $(call gb_Library_add_exception_objects,sc,\
$(eval $(call gb_Helper_optional,OPENCL,\
$(call gb_Library_add_exception_objects,sc,\
sc/source/core/opencl/formulagroupcl \
- sc/source/core/opencl/openclwrapper \
- sc/source/core/opencl/opencl_device \
sc/source/core/opencl/opbase \
sc/source/core/opencl/op_financial \
sc/source/core/opencl/op_database \
diff --git a/sc/source/core/inc/openclwrapper.hxx b/sc/source/core/inc/openclwrapper.hxx
deleted file mode 100644
index 9dad74749bb6..000000000000
--- a/sc/source/core/inc/openclwrapper.hxx
+++ /dev/null
@@ -1,108 +0,0 @@
-/* -*- 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_SC_SOURCE_CORE_OPENCL_OPENCLWRAPPER_HXX
-#define INCLUDED_SC_SOURCE_CORE_OPENCL_OPENCLWRAPPER_HXX
-
-#include <config_features.h>
-#include <sal/detail/log.h>
-#include <opencl/platforminfo.hxx>
-#include <osl/file.hxx>
-#include <vector>
-#include <boost/shared_ptr.hpp>
-#include <cassert>
-
-#include <rtl/string.hxx>
-
-#include <clew.h>
-
-#define CHECK_OPENCL(status,name) \
-if( status != CL_SUCCESS ) \
-{ \
- printf ("OpenCL error code is %d at " SAL_DETAIL_WHERE " when %s .\n", status, name); \
- return false; \
-}
-
-#define MAX_CLFILE_NUM 50
-
-#include <cstdio>
-
-struct KernelEnv
-{
- cl_context mpkContext;
- cl_command_queue mpkCmdQueue;
- cl_program mpkProgram;
-};
-
-namespace sc { namespace opencl {
-
-typedef unsigned int uint;
-
-struct OpenCLEnv
-{
- cl_platform_id mpOclPlatformID;
- cl_context mpOclContext;
- cl_device_id mpOclDevsID;
- cl_command_queue mpOclCmdQueue;
-};
-
-struct GPUEnv
-{
- //share vb in all modules in hb library
- cl_platform_id mpPlatformID;
- cl_device_type mDevType;
- cl_context mpContext;
- cl_device_id *mpArryDevsID;
- cl_device_id mpDevID;
- cl_command_queue mpCmdQueue;
- cl_program mpArryPrograms[MAX_CLFILE_NUM]; //one program object maps one kernel source file
- int mnIsUserCreated; // 1: created , 0:no create and needed to create by opencl wrapper
- bool mnKhrFp64Flag;
- bool mnAmdFp64Flag;
-};
-
-class OpenCLDevice
-{
-public:
- static GPUEnv gpuEnv;
- static bool bIsInited;
- static OString maCacheFolder;
-
- static bool initOpenCLRunEnv( GPUEnv *gpu );
- static void releaseOpenCLEnv( GPUEnv *gpuInfo );
- static bool initOpenCLRunEnv( int argc );
- static bool generatBinFromKernelSource( cl_program program, const char * clFileName );
- static bool writeBinaryToFile( const OString& rName, const char* birary, size_t numBytes );
- static std::vector<boost::shared_ptr<osl::File> > binaryGenerated( const char * clFileName, cl_context context);
- static bool buildProgramFromBinary(const char* buildOption, GPUEnv* gpuEnv, const char* filename, int idx);
-
- static bool initOpenCLAttr( OpenCLEnv * env );
- static void setKernelEnv( KernelEnv *envInfo );
-};
-
-const std::vector<OpenCLPlatformInfo>& fillOpenCLInfo();
-
-/**
- * Used to set or switch between OpenCL devices.
- *
- * @param pDeviceId the id of the opencl device of type cl_device_id, NULL means use software calculation
- * @param bAutoSelect use the algorithm to select the best OpenCL device
- *
- * @return returns true if there is a valid opencl device that has been set up
- */
-bool switchOpenCLDevice(const OUString* pDeviceId, bool bAutoSelect,
- bool bForceEvaluation);
-
-void getOpenCLDeviceInfo(size_t& rDeviceId, size_t& rPlatformId);
-
-}}
-
-#endif
-
-/* vim:set shiftwidth=4 softtabstop=4 expandtab: */
diff --git a/sc/source/core/opencl/formulagroupcl.cxx b/sc/source/core/opencl/formulagroupcl.cxx
index 3dfcc10d012a..0615410c0ad3 100644
--- a/sc/source/core/opencl/formulagroupcl.cxx
+++ b/sc/source/core/opencl/formulagroupcl.cxx
@@ -19,7 +19,7 @@
#include <formula/vectortoken.hxx>
#include "scmatrix.hxx"
-#include "openclwrapper.hxx"
+#include <opencl/openclwrapper.hxx>
#include "op_financial.hxx"
#include "op_database.hxx"
@@ -114,7 +114,7 @@ size_t VectorRef::Marshal( cl_kernel k, int argno, int, cl_program )
}
// Obtain cl context
KernelEnv kEnv;
- OpenCLDevice::setKernelEnv(&kEnv);
+ ::opencl::OpenCLDevice::setKernelEnv(&kEnv);
cl_int err;
if (pHostBuffer)
{
@@ -205,7 +205,7 @@ public:
// marshaling
// Obtain cl context
KernelEnv kEnv;
- OpenCLDevice::setKernelEnv(&kEnv);
+ ::opencl::OpenCLDevice::setKernelEnv(&kEnv);
// Pass the scalar result back to the rest of the formula kernel
cl_int err = clSetKernelArg(k, argno, sizeof(cl_uint), (void*)&hashCode);
if (CL_SUCCESS != err)
@@ -389,7 +389,7 @@ size_t DynamicKernelStringArgument::Marshal( cl_kernel k, int argno, int, cl_pro
FormulaToken* ref = mFormulaTree->GetFormulaToken();
// Obtain cl context
KernelEnv kEnv;
- OpenCLDevice::setKernelEnv(&kEnv);
+ ::opencl::OpenCLDevice::setKernelEnv(&kEnv);
cl_int err;
formula::VectorRefArray vRef;
size_t nStrings = 0;
@@ -1132,7 +1132,7 @@ public:
assert(Base::mpClmem == NULL);
// Obtain cl context
KernelEnv kEnv;
- OpenCLDevice::setKernelEnv(&kEnv);
+ ::opencl::OpenCLDevice::setKernelEnv(&kEnv);
cl_int err;
size_t nInput = mpDVR->GetArrayLength();
size_t nCurWindowSize = mpDVR->GetRefRowSize();
@@ -1913,7 +1913,7 @@ public:
{
// Obtain cl context
KernelEnv kEnv;
- OpenCLDevice::setKernelEnv(&kEnv);
+ ::opencl::OpenCLDevice::setKernelEnv(&kEnv);
cl_int err;
cl_mem pClmem2;
@@ -1971,7 +1971,7 @@ public:
{
// Obtain cl context
KernelEnv kEnv;
- OpenCLDevice::setKernelEnv(&kEnv);
+ ::opencl::OpenCLDevice::setKernelEnv(&kEnv);
cl_int err;
DynamicKernelArgument* Arg = mvSubArguments[0].get();
DynamicKernelSlidingArgument<VectorRef>* slidingArgPtr =
@@ -3282,11 +3282,11 @@ public:
DynamicKernelSoPArguments>(mpRoot, new OpNop);
std::stringstream decl;
- if (OpenCLDevice::gpuEnv.mnKhrFp64Flag)
+ if (::opencl::OpenCLDevice::gpuEnv.mnKhrFp64Flag)
{
decl << "#pragma OPENCL EXTENSION cl_khr_fp64: enable\n";
}
- else if (OpenCLDevice::gpuEnv.mnAmdFp64Flag)
+ else if (::opencl::OpenCLDevice::gpuEnv.mnAmdFp64Flag)
{
decl << "#pragma OPENCL EXTENSION cl_amd_fp64: enable\n";
}
@@ -3355,7 +3355,7 @@ public:
{
// Obtain cl context
KernelEnv kEnv;
- OpenCLDevice::setKernelEnv(&kEnv);
+ ::opencl::OpenCLDevice::setKernelEnv(&kEnv);
cl_int err;
// The results
mpResClmem = clCreateBuffer(kEnv.mpkContext,
@@ -3417,7 +3417,7 @@ void DynamicKernel::CreateKernel()
// Compile kernel here!!!
// Obtain cl context
KernelEnv kEnv;
- OpenCLDevice::setKernelEnv(&kEnv);
+ ::opencl::OpenCLDevice::setKernelEnv(&kEnv);
const char* src = mFullProgramSrc.c_str();
static std::string lastOneKernelHash = "";
static std::string lastSecondKernelHash = "";
@@ -3439,11 +3439,11 @@ void DynamicKernel::CreateKernel()
{
clReleaseProgram(lastSecondProgram);
}
- if (OpenCLDevice::buildProgramFromBinary("",
- &OpenCLDevice::gpuEnv, KernelHash.c_str(), 0))
+ if (::opencl::OpenCLDevice::buildProgramFromBinary("",
+ &::opencl::OpenCLDevice::gpuEnv, KernelHash.c_str(), 0))
{
- mpProgram = OpenCLDevice::gpuEnv.mpArryPrograms[0];
- OpenCLDevice::gpuEnv.mpArryPrograms[0] = NULL;
+ mpProgram = ::opencl::OpenCLDevice::gpuEnv.mpArryPrograms[0];
+ ::opencl::OpenCLDevice::gpuEnv.mpArryPrograms[0] = NULL;
}
else
{
@@ -3452,7 +3452,7 @@ void DynamicKernel::CreateKernel()
if (err != CL_SUCCESS)
throw OpenCLError(err, __FILE__, __LINE__);
err = clBuildProgram(mpProgram, 1,
- OpenCLDevice::gpuEnv.mpArryDevsID, "", NULL, NULL);
+ ::opencl::OpenCLDevice::gpuEnv.mpArryDevsID, "", NULL, NULL);
if (err != CL_SUCCESS)
{
#if OSL_DEBUG_LEVEL > 0
@@ -3460,7 +3460,7 @@ void DynamicKernel::CreateKernel()
{
cl_build_status stat;
cl_int e = clGetProgramBuildInfo(
- mpProgram, OpenCLDevice::gpuEnv.mpArryDevsID[0],
+ mpProgram, ::opencl::OpenCLDevice::gpuEnv.mpArryDevsID[0],
CL_PROGRAM_BUILD_STATUS, sizeof(cl_build_status),
&stat, 0);
SAL_WARN_IF(
@@ -3472,7 +3472,7 @@ void DynamicKernel::CreateKernel()
{
size_t n;
e = clGetProgramBuildInfo(
- mpProgram, OpenCLDevice::gpuEnv.mpArryDevsID[0],
+ mpProgram, ::opencl::OpenCLDevice::gpuEnv.mpArryDevsID[0],
CL_PROGRAM_BUILD_LOG, 0, 0, &n);
SAL_WARN_IF(
e != CL_SUCCESS || n == 0, "sc.opencl",
@@ -3483,7 +3483,7 @@ void DynamicKernel::CreateKernel()
{
std::vector<char> log(n);
e = clGetProgramBuildInfo(
- mpProgram, OpenCLDevice::gpuEnv.mpArryDevsID[0],
+ mpProgram, ::opencl::OpenCLDevice::gpuEnv.mpArryDevsID[0],
CL_PROGRAM_BUILD_LOG, n, &log[0], 0);
SAL_WARN_IF(
e != CL_SUCCESS || n == 0, "sc.opencl",
@@ -3502,7 +3502,7 @@ void DynamicKernel::CreateKernel()
throw OpenCLError(err, __FILE__, __LINE__);
}
// Generate binary out of compiled kernel.
- OpenCLDevice::generatBinFromKernelSource(mpProgram,
+ ::opencl::OpenCLDevice::generatBinFromKernelSource(mpProgram,
(mKernelSignature + GetMD5()).c_str());
}
lastSecondKernelHash = lastOneKernelHash;
@@ -3668,7 +3668,7 @@ bool FormulaGroupInterpreterOpenCL::interpret( ScDocument& rDoc,
{
// Obtain cl context
KernelEnv kEnv;
- OpenCLDevice::setKernelEnv(&kEnv);
+ ::opencl::OpenCLDevice::setKernelEnv(&kEnv);
// Run the kernel.
pKernel->Launch(xGroup->mnLength);
// Map results back
diff --git a/sc/source/core/opencl/opencl_device.cxx b/sc/source/core/opencl/opencl_device.cxx
deleted file mode 100644
index f845f423d0e1..000000000000
--- a/sc/source/core/opencl/opencl_device.cxx
+++ /dev/null
@@ -1,597 +0,0 @@
-/* -*- 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/.
- */
-
-#ifdef _WIN32
-#include <prewin.h>
-#include <postwin.h>
-#elif defined __MACH__
-#include <mach/mach_time.h>
-#else
-#include <sys/time.h>
-#endif
-#include <time.h>
-#include <math.h>
-#include <float.h>
-#include <iostream>
-#include <sstream>
-#include <vector>
-
-#include <boost/scoped_ptr.hpp>
-
-#include <comphelper/random.hxx>
-#include <opencl/openclconfig.hxx>
-#include <opencl/platforminfo.hxx>
-#include <sal/log.hxx>
-
-#include "opencl_device.hxx"
-#include "openclwrapper.hxx"
-
-#define INPUTSIZE 15360
-#define OUTPUTSIZE 15360
-
-#define STRINGIFY(...) #__VA_ARGS__"\n"
-
-#define DS_CHECK_STATUS(status, name) \
- if (CL_SUCCESS != status) \
- { \
- SAL_INFO("sc.opencl.device", "Error code is " << status << " at " name); \
- }
-
-namespace sc { namespace OpenCLDevice {
-
-bool bIsInited = false;
-bool bIsDeviceSelected = false;
-ds_device selectedDevice;
-
-struct LibreOfficeDeviceScore
-{
- double fTime; // small time means faster device
- bool bNoCLErrors; // were there any opencl errors
-};
-
-struct LibreOfficeDeviceEvaluationIO
-{
- std::vector<double> input0;
- std::vector<double> input1;
- std::vector<double> input2;
- std::vector<double> input3;
- std::vector<double> output;
- unsigned long inputSize;
- unsigned long outputSize;
-};
-
-struct timer
-{
-#ifdef _WIN32
- LARGE_INTEGER start;
-#else
- long long start;
-#endif
-};
-
-const char* source = STRINGIFY(
-\n#if defined(KHR_DP_EXTENSION)
-\n#pragma OPENCL EXTENSION cl_khr_fp64 : enable
-\n#elif defined(AMD_DP_EXTENSION)
-\n#pragma OPENCL EXTENSION cl_amd_fp64 : enable
-\n#endif
- \n
- int isNan(fp_t a) { return a != a; }
- fp_t fsum(fp_t a, fp_t b) { return a + b; }
-
- fp_t fAverage(__global fp_t* input)
-{
- fp_t sum = 0;
- int count = 0;
- for (int i = 0; i < INPUTSIZE; i++)
- {
- if (!isNan(input[i]))
- {
- sum = fsum(input[i], sum);
- count += 1;
- }
- }
- return sum / (fp_t)count;
-}
- fp_t fMin(__global fp_t* input)
-{
- fp_t min = MAXFLOAT;
- for (int i = 0; i < INPUTSIZE; i++)
- {
- if (!isNan(input[i]))
- {
- min = fmin(input[i], min);
- }
- }
- return min;
-}
- fp_t fSoP(__global fp_t* input0, __global fp_t* input1)
-{
- fp_t sop = 0.0;
- for (int i = 0; i < INPUTSIZE; i++)
- {
- sop += (isNan(input0[i]) ? 0 : input0[i]) * (isNan(input1[i]) ? 0 : input1[i]);
- }
- return sop;
-}
- __kernel void DynamicKernel(
- __global fp_t* result, __global fp_t* input0, __global fp_t* input1, __global fp_t* input2, __global fp_t* input3)
-{
- int gid0 = get_global_id(0);
- fp_t tmp0 = fAverage(input0);
- fp_t tmp1 = fMin(input1) * fSoP(input2, input3);
- result[gid0] = fsum(tmp0, tmp1);
-}
- );
-
-size_t sourceSize[] = { strlen(source) };
-
-/*************************************************************************/
-/* INTERNAL FUNCTIONS */
-/*************************************************************************/
-/* Timer functions - start timer */
-void timerStart(timer* mytimer)
-{
-#ifdef _WIN32
- QueryPerformanceCounter(&mytimer->start);
-#elif defined __MACH__
- mytimer->start = mach_absolute_time();
-#else
- struct timespec s;
- clock_gettime(CLOCK_MONOTONIC, &s);
- mytimer->start = (long long)s.tv_sec * (long long)1.0E6 + (long long)s.tv_nsec / (long long)1.0E3;
-#endif
-}
-
-/* Timer functions - get current value */
-double timerCurrent(timer* mytimer)
-{
-#ifdef _WIN32
- LARGE_INTEGER stop, frequency;
- QueryPerformanceCounter(&stop);
- QueryPerformanceFrequency(&frequency);
- double time = ((double)(stop.QuadPart - mytimer->start.QuadPart) / frequency.QuadPart);
-#elif defined __MACH__
- static mach_timebase_info_data_t info = { 0, 0 };
- if (info.numer == 0)
- mach_timebase_info(&info);
- long long stop = mach_absolute_time();
- double time = ((stop - mytimer->start) * (double) info.numer / info.denom) / 1.0E9;
-#else
- struct timespec s;
- long long stop;
- clock_gettime(CLOCK_MONOTONIC, &s);
- stop = (long long)s.tv_sec * (long long)1.0E6 + (long long)s.tv_nsec / (long long)1.0E3;
- double time = ((double)(stop - mytimer->start) / 1.0E6);
-#endif
- return time;
-}
-
-/* Random number generator */
-double random(double min, double max)
-{
- if (min == max)
- return min;
- return comphelper::rng::uniform_real_distribution(min, max);
-}
-
-/* Populate input */
-void populateInput(LibreOfficeDeviceEvaluationIO* testData)
-{
- double* input0 = &testData->input0[0];
- double* input1 = &testData->input1[0];
- double* input2 = &testData->input2[0];
- double* input3 = &testData->input3[0];
- for (unsigned long i = 0; i < testData->inputSize; i++)
- {
- input0[i] = random(0, i);
- input1[i] = random(0, i);
- input2[i] = random(0, i);
- input3[i] = random(0, i);
- }
-}
-/* Encode score object as byte string */
-ds_status serializeScore(ds_device* device, void** serializedScore, unsigned int* serializedScoreSize)
-{
- *serializedScoreSize = sizeof(LibreOfficeDeviceScore);
- *serializedScore = (void*)new unsigned char[*serializedScoreSize];
- memcpy(*serializedScore, device->score, *serializedScoreSize);
- return DS_SUCCESS;
-}
-
-/* Parses byte string and stores in score object */
-ds_status deserializeScore(ds_device* device, const unsigned char* serializedScore, unsigned int serializedScoreSize)
-{
- // check that serializedScoreSize == sizeof(LibreOfficeDeviceScore);
- device->score = new LibreOfficeDeviceScore;
- memcpy(device->score, serializedScore, serializedScoreSize);
- return DS_SUCCESS;
-}
-
-/* Releases memory held by score */
-ds_status releaseScore(void* score)
-{
- if (NULL != score)
- {
- delete (LibreOfficeDeviceScore*)score;
- }
- return DS_SUCCESS;
-}
-
-/* Evaluate devices */
-ds_status evaluateScoreForDevice(ds_device* device, void* evalData)
-{
- if (DS_DEVICE_OPENCL_DEVICE == device->type)
- {
- /* Evaluating an OpenCL device */
- SAL_INFO("sc.opencl.device", "Device: \"" << device->oclDeviceName << "\" (OpenCL) evaluation...");
- cl_int clStatus;
- /* Check for 64-bit float extensions */
- size_t aDevExtInfoSize = 0;
- clStatus = clGetDeviceInfo(device->oclDeviceID, CL_DEVICE_EXTENSIONS, 0, NULL, &aDevExtInfoSize);
- DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clGetDeviceInfo");
-
- char* aExtInfo = new char[aDevExtInfoSize];
- clStatus = clGetDeviceInfo(device->oclDeviceID, CL_DEVICE_EXTENSIONS, sizeof(char) * aDevExtInfoSize, aExtInfo, NULL);
- DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clGetDeviceInfo");
- bool bKhrFp64Flag = false;
- bool bAmdFp64Flag = false;
- const char* buildOption = NULL;
- std::string tmpStr("-Dfp_t=double -Dfp_t4=double4 -Dfp_t16=double16 -DINPUTSIZE=");
- std::ostringstream tmpOStrStr;
- tmpOStrStr << std::dec << INPUTSIZE;
- tmpStr.append(tmpOStrStr.str());
-
- if ((std::string(aExtInfo)).find("cl_khr_fp64") != std::string::npos)
- {
- bKhrFp64Flag = true;
- //buildOption = "-D KHR_DP_EXTENSION -Dfp_t=double -Dfp_t4=double4 -Dfp_t16=double16";
- tmpStr.append(" -DKHR_DP_EXTENSION");
- buildOption = tmpStr.c_str();
- SAL_INFO("sc.opencl.device", "... has cl_khr_fp64");
- }
- else if ((std::string(aExtInfo)).find("cl_amd_fp64") != std::string::npos)
- {
- bAmdFp64Flag = true;
- //buildOption = "-D AMD_DP_EXTENSION -Dfp_t=double -Dfp_t4=double4 -Dfp_t16=double16";
- tmpStr.append(" -DAMD_DP_EXTENSION");
- buildOption = tmpStr.c_str();
- SAL_INFO("sc.opencl.device", "... has cl_amd_fp64");
- }
- delete[] aExtInfo;
-
- if (!bKhrFp64Flag && !bAmdFp64Flag)
- {
- /* No 64-bit float support */
- device->score = (void*)new LibreOfficeDeviceScore;
- ((LibreOfficeDeviceScore*)device->score)->fTime = DBL_MAX;
- ((LibreOfficeDeviceScore*)device->score)->bNoCLErrors = true;
- SAL_INFO("sc.opencl.device", "... no fp64 support");
- }
- else
- {
- /* 64-bit float support present */
-
- /* Create context and command queue */
- cl_context clContext = clCreateContext(NULL, 1, &device->oclDeviceID, NULL, NULL, &clStatus);
- DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clCreateContext");
- cl_command_queue clQueue = clCreateCommandQueue(clContext, device->oclDeviceID, 0, &clStatus);
- DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clCreateCommandQueue");
-
- /* Build program */
- cl_program clProgram = clCreateProgramWithSource(clContext, 1, &source, sourceSize, &clStatus);
- DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clCreateProgramWithSource");
- clStatus = clBuildProgram(clProgram, 1, &device->oclDeviceID, buildOption, NULL, NULL);
- DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clBuildProgram");
- if (CL_SUCCESS != clStatus)
- {
- /* Build program failed */
- size_t length;
- char* buildLog;
- clStatus = clGetProgramBuildInfo(clProgram, device->oclDeviceID, CL_PROGRAM_BUILD_LOG, 0, NULL, &length);
- buildLog = (char*)malloc(length);
- clGetProgramBuildInfo(clProgram, device->oclDeviceID, CL_PROGRAM_BUILD_LOG, length, buildLog, &length);
- SAL_INFO("sc.opencl.device", "Build Errors:\n" << buildLog);
- free(buildLog);
-
- device->score = (void*)new LibreOfficeDeviceScore;
- ((LibreOfficeDeviceScore*)device->score)->fTime = DBL_MAX;
- ((LibreOfficeDeviceScore*)device->score)->bNoCLErrors = false;
- }
- else
- {
- /* Build program succeeded */
- timer kernelTime;
- timerStart(&kernelTime);
-
- /* Run kernel */
- LibreOfficeDeviceEvaluationIO* testData = (LibreOfficeDeviceEvaluationIO*)evalData;
- cl_kernel clKernel = clCreateKernel(clProgram, "DynamicKernel", &clStatus);
- DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clCreateKernel");
- cl_mem clResult = clCreateBuffer(clContext, CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR, sizeof(cl_double) * testData->outputSize, &testData->output[0], &clStatus);
- DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clCreateBuffer::clResult");
- cl_mem clInput0 = clCreateBuffer(clContext, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, sizeof(cl_double) * testData->inputSize, &testData->input0[0], &clStatus);
- DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clCreateBuffer::clInput0");
- cl_mem clInput1 = clCreateBuffer(clContext, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, sizeof(cl_double) * testData->inputSize, &testData->input1[0], &clStatus);
- DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clCreateBuffer::clInput1");
- cl_mem clInput2 = clCreateBuffer(clContext, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, sizeof(cl_double) * testData->inputSize, &testData->input2[0], &clStatus);
- DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clCreateBuffer::clInput2");
- cl_mem clInput3 = clCreateBuffer(clContext, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, sizeof(cl_double) * testData->inputSize, &testData->input3[0], &clStatus);
- DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clCreateBuffer::clInput3");
- clStatus = clSetKernelArg(clKernel, 0, sizeof(cl_mem), (void*)&clResult);
- DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clSetKernelArg::clResult");
- clStatus = clSetKernelArg(clKernel, 1, sizeof(cl_mem), (void*)&clInput0);
- DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clSetKernelArg::clInput0");
- clStatus = clSetKernelArg(clKernel, 2, sizeof(cl_mem), (void*)&clInput1);
- DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clSetKernelArg::clInput1");
- clStatus = clSetKernelArg(clKernel, 3, sizeof(cl_mem), (void*)&clInput2);
- DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clSetKernelArg::clInput2");
- clStatus = clSetKernelArg(clKernel, 4, sizeof(cl_mem), (void*)&clInput3);
- DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clSetKernelArg::clInput3");
- size_t globalWS[1] = { testData->outputSize };
- size_t localSize[1] = { 64 };
- clStatus = clEnqueueNDRangeKernel(clQueue, clKernel, 1, 0, globalWS, localSize, 0, NULL, NULL);
- DS_CHECK_STATUS(clStatus, "evaluateScoreForDevice::clEnqueueNDRangeKernel");
- clFinish(clQueue);
- clReleaseMemObject(clInput3);
- clReleaseMemObject(clInput2);
- clReleaseMemObject(clInput1);
- clReleaseMemObject(clInput0);
- clReleaseMemObject(clResult);
- clReleaseKernel(clKernel);
-
- device->score = (void*)new LibreOfficeDeviceScore;
- ((LibreOfficeDeviceScore*)device->score)->fTime = timerCurrent(&kernelTime);
- ((LibreOfficeDeviceScore*)device->score)->bNoCLErrors = true;
- }
-
- clReleaseProgram(clProgram);
- clReleaseCommandQueue(clQueue);
- clReleaseContext(clContext);
- }
- }
- else
- {
- /* Evaluating an Native CPU device */
- SAL_INFO("sc.opencl.device", "Device: \"CPU\" (Native) evaluation...");
- timer kernelTime;
- timerStart(&kernelTime);
-
- LibreOfficeDeviceEvaluationIO* testData = (LibreOfficeDeviceEvaluationIO*)evalData;
- for (unsigned long j = 0; j < testData->outputSize; j++)
- {
- double fAverage = 0.0f;
- double fMin = DBL_MAX;
- double fSoP = 0.0f;
- for (unsigned long i = 0; i < testData->inputSize; i++)
- {
- fAverage += testData->input0[i];
- fMin = ((fMin < testData->input1[i]) ? fMin : testData->input1[i]);
- fSoP += testData->input2[i] * testData->input3[i];
- }
- fAverage /= testData->inputSize;
- testData->output[j] = fAverage + (fMin * fSoP);
- }
-
- // InterpretTail - the S/W fallback is nothing like as efficient
- // as any good openCL implementation: no SIMD, tons of branching
- // in the inner loops etc. Generously characterise it as only 10x
- // slower than the above.
- float fInterpretTailFactor = 10.0;
-
- device->score = (void*)new LibreOfficeDeviceScore;
- ((LibreOfficeDeviceScore*)device->score)->fTime = timerCurrent(&kernelTime);
- ((LibreOfficeDeviceScore*)device->score)->bNoCLErrors = true;
-
- ((LibreOfficeDeviceScore*)device->score)->fTime *= fInterpretTailFactor;
- }
- return DS_SUCCESS;
-}
-
-/* Pick best device */
-ds_status pickBestDevice(ds_profile* profile, int* bestDeviceIdx)
-{
- double bestScore = DBL_MAX;
- *bestDeviceIdx = -1;
-
- for (unsigned int d = 0; d < profile->numDevices; d++)
- {
- ds_device device = profile->devices[d];
- LibreOfficeDeviceScore *pScore = (LibreOfficeDeviceScore*)device.score;
-
- // Check blacklist and whitelist for actual devices
- if (device.type == DS_DEVICE_OPENCL_DEVICE)
- {
- // There is a silly impedance mismatch here. Why do we
- // need two different ways to describe an OpenCL platform
- // and an OpenCL device driver?
-
- OpenCLPlatformInfo aPlatform;
- OpenCLDeviceInfo aDevice;
-
- // We know that only the below fields are used by checkForKnownBadCompilers()
- aPlatform.maVendor = OUString(device.oclPlatformVendor, strlen(device.oclPlatformVendor), RTL_TEXTENCODING_UTF8);
- aDevice.maName = OUString(device.oclDeviceName, strlen(device.oclDeviceName), RTL_TEXTENCODING_UTF8);
- aDevice.maDriver = OUString(device.oclDriverVersion, strlen(device.oclDriverVersion), RTL_TEXTENCODING_UTF8);
-
- // If blacklisted or not whitelisted, ignore it
- if (OpenCLConfig::get().checkImplementation(aPlatform, aDevice))
- {
- SAL_INFO("sc.opencl.device", "Device[" << d << "] " << device.oclDeviceName << " is blacklisted or not whitelisted");
- pScore->fTime = DBL_MAX;
- pScore->bNoCLErrors = true;
- }
- }
-
- double fScore = DBL_MAX;
- if (pScore)
- {
- fScore = pScore->fTime;
- }
- else
- {
- SAL_INFO("sc.opencl.device", "Unusual null score");
- }
-
- if (DS_DEVICE_OPENCL_DEVICE == device.type)
- {
- SAL_INFO("sc.opencl.device", "Device[" << d << "] " << device.oclDeviceName << " (OpenCL) score is " << fScore);
- }
- else
- {
- SAL_INFO("sc.opencl.device", "Device[" << d << "] CPU (Native) score is " << fScore);
- }
- if (fScore < bestScore)
- {
- bestScore = fScore;
- *bestDeviceIdx = d;
- }
- }
- if (DS_DEVICE_OPENCL_DEVICE == profile->devices[*bestDeviceIdx].type)
- {
- SAL_INFO("sc.opencl.device", "Selected Device[" << *bestDeviceIdx << "]: " << profile->devices[*bestDeviceIdx].oclDeviceName << "(OpenCL).");
- }
- else
- {
- SAL_INFO("sc.opencl.device", "Selected Device[" << *bestDeviceIdx << "]: CPU (Native).");
- }
-
- return DS_SUCCESS;
-}
-
-/* Return device ID for matching device name */
-int matchDevice(ds_profile* profile, char* deviceName)
-{
- int deviceMatch = -1;
- for (unsigned int d = 0; d < profile->numDevices - 1; d++)
- {
- if ((std::string(profile->devices[d].oclDeviceName)).find(deviceName) != std::string::npos) deviceMatch = d;
- }
- if (std::string("NATIVE_CPU").find(deviceName) != std::string::npos) deviceMatch = profile->numDevices - 1;
- return deviceMatch;
-}
-
-/*************************************************************************/
-/* EXTERNAL FUNCTIONS */
-/*************************************************************************/
-ds_device getDeviceSelection(const char* sProfilePath, bool bForceSelection)
-{
- /* Run only if device is not yet selected */
- if (!bIsDeviceSelected || bForceSelection)
- {
- /* Setup */
- ds_status status;
- ds_profile* profile = NULL;
- status = initDSProfile(&profile, "LibreOffice v0.1");
-
- if (!profile)
- {
- // failed to initialize profile.
- selectedDevice.type = DS_DEVICE_NATIVE_CPU;
- return selectedDevice;
- }
-
- /* Try reading scores from file */
- std::string tmpStr(sProfilePath);
- const char* fileName = tmpStr.append("sc_opencl_device_profile.dat").c_str();
- if (!bForceSelection)
- {
- status = readProfileFromFile(profile, deserializeScore, fileName);
- }
- else
- {
- status = DS_INVALID_PROFILE;
- SAL_INFO("sc.opencl.device", "Performing forced profiling.");
- }
- if (DS_SUCCESS != status)
- {
- if (!bForceSelection)
- {
- SAL_INFO("sc.opencl.device", "Profile file not available (" << fileName << "); performing profiling.");
- }
-
- /* Populate input data for micro-benchmark */
- boost::scoped_ptr<LibreOfficeDeviceEvaluationIO> testData(new LibreOfficeDeviceEvaluationIO);
- testData->inputSize = INPUTSIZE;
- testData->outputSize = OUTPUTSIZE;
- testData->input0.resize(testData->inputSize);
- testData->input1.resize(testData->inputSize);
- testData->input2.resize(testData->inputSize);
- testData->input3.resize(testData->inputSize);
- testData->output.resize(testData->outputSize);
- populateInput(testData.get());
-
- /* Perform evaluations */
- unsigned int numUpdates;
- status = profileDevices(profile, DS_EVALUATE_ALL, evaluateScoreForDevice, (void*)testData.get(), &numUpdates);
-
- if (DS_SUCCESS == status)
- {
- /* Write scores to file */
- status = writeProfileToFile(profile, serializeScore, fileName);
- if (DS_SUCCESS == status)
- {
- SAL_INFO("sc.opencl.device", "Scores written to file (" << fileName << ").");
- }
- else
- {
- SAL_INFO("sc.opencl.device", "Error saving scores to file (" << fileName << "); scores not written to file.");
- }
- }
- else
- {
- SAL_INFO("sc.opencl.device", "Unable to evaluate performance; scores not written to file.");
- }
- }
- else
- {
- SAL_INFO("sc.opencl.device", "Profile read from file (" << fileName << ").");
- }
-
- /* Pick best device */
- int bestDeviceIdx;
- pickBestDevice(profile, &bestDeviceIdx);
-
- /* Overide if necessary */
- char* overrideDeviceStr = getenv("SC_OPENCL_DEVICE_OVERRIDE");
- if (NULL != overrideDeviceStr)
- {
- int overrideDeviceIdx = matchDevice(profile, overrideDeviceStr);
- if (-1 != overrideDeviceIdx)
- {
- SAL_INFO("sc.opencl.device", "Overriding Device Selection (SC_OPENCL_DEVICE_OVERRIDE=" << overrideDeviceStr << ").");
- bestDeviceIdx = overrideDeviceIdx;
- if (DS_DEVICE_OPENCL_DEVICE == profile->devices[bestDeviceIdx].type)
- {
- SAL_INFO("sc.opencl.device", "Selected Device[" << bestDeviceIdx << "]: " << profile->devices[bestDeviceIdx].oclDeviceName << " (OpenCL).");
- }
- else
- {
- SAL_INFO("sc.opencl.device", "Selected Device[" << bestDeviceIdx << "]: CPU (Native).");
- }
- }
- else
- {
- SAL_INFO("sc.opencl.device", "Ignoring invalid SC_OPENCL_DEVICE_OVERRIDE=" << overrideDeviceStr << ").");
- }
- }
-
- /* Final device selection */
- selectedDevice = profile->devices[bestDeviceIdx];
- bIsDeviceSelected = true;
-
- /* Release profile */
- releaseDSProfile(profile, releaseScore);
- }
- return selectedDevice;
-}
-
-}}
-
-/* vim:set shiftwidth=4 softtabstop=4 expandtab: */
diff --git a/sc/source/core/opencl/opencl_device.hxx b/sc/source/core/opencl/opencl_device.hxx
deleted file mode 100644
index c5367f24f20c..000000000000
--- a/sc/source/core/opencl/opencl_device.hxx
+++ /dev/null
@@ -1,26 +0,0 @@
-/* -*- 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_SC_SOURCE_CORE_OPENCL_OPENCL_DEVICE_HXX
-#define INCLUDED_SC_SOURCE_CORE_OPENCL_OPENCL_DEVICE_HXX
-
-#pragma once
-#include "opencl_device_selection.h"
-
-namespace sc { namespace OpenCLDevice {
-
-ds_device getDeviceSelection(const char* pFileName, bool bForceSelection = false);
-bool selectedDeviceIsOpenCL(ds_device device);
-bool selectedDeviceIsNativeCPU(ds_device device);
-
-}}
-
-#endif
-
-/* vim:set shiftwidth=4 softtabstop=4 expandtab: */
diff --git a/sc/source/core/opencl/opencl_device_selection.h b/sc/source/core/opencl/opencl_device_selection.h
deleted file mode 100644
index 30e947a8c01d..000000000000
--- a/sc/source/core/opencl/opencl_device_selection.h
+++ /dev/null
@@ -1,642 +0,0 @@
-/* -*- 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_SC_SOURCE_CORE_OPENCL_OPENCL_DEVICE_SELECTION_H
-#define INCLUDED_SC_SOURCE_CORE_OPENCL_OPENCL_DEVICE_SELECTION_H
-
-
-#ifdef _MSC_VER
-#define _CRT_SECURE_NO_WARNINGS
-#endif
-
-#include <stdlib.h>
-#include <stdio.h>
-#include <string.h>
-#include <clew.h>
-
-#define DS_DEVICE_NAME_LENGTH 256
-
-enum ds_status
-{
- DS_SUCCESS = 0
- ,DS_INVALID_PROFILE = 1000
- ,DS_MEMORY_ERROR
- , DS_INVALID_PERF_EVALUATOR_TYPE
- , DS_INVALID_PERF_EVALUATOR
- , DS_PERF_EVALUATOR_ERROR
- , DS_FILE_ERROR
- , DS_UNKNOWN_DEVICE_TYPE
- , DS_PROFILE_FILE_ERROR
- , DS_SCORE_SERIALIZER_ERROR
- , DS_SCORE_DESERIALIZER_ERROR
-};
-
-// device type
-enum ds_device_type
-{
- DS_DEVICE_NATIVE_CPU = 0
- ,DS_DEVICE_OPENCL_DEVICE
-};
-
-
-struct ds_device
-{
- ds_device_type type;
- cl_device_id oclDeviceID;
- char* oclPlatformVendor;
- char* oclDeviceName;
- char* oclDriverVersion;
- void* score; // a pointer to the score data, the content/format is application defined
-};
-
-struct ds_profile
-{
- unsigned int numDevices;
- ds_device* devices;
- const char* version;
-};
-
-// deallocate memory used by score
-typedef ds_status(* ds_score_release)(void* score);
-inline ds_status releaseDSProfile(ds_profile* profile, ds_score_release sr)
-{
- ds_status status = DS_SUCCESS;
- if (profile != NULL)
- {
- if (profile->devices != NULL && sr != NULL)
- {
- unsigned int i;
- for (i = 0; i < profile->numDevices; i++)
- {
- free(profile->devices[i].oclPlatformVendor);
- free(profile->devices[i].oclDeviceName);
- free(profile->devices[i].oclDriverVersion);
- status = sr(profile->devices[i].score);
- if (status != DS_SUCCESS) break;
- }
- free(profile->devices);
- }
- free(profile);
- }
- return status;
-}
-
-
-inline ds_status initDSProfile(ds_profile** p, const char* version)
-{
- int numDevices;
- cl_uint numPlatforms;
- cl_platform_id* platforms = NULL;
- cl_device_id* devices = NULL;
- ds_status status = DS_SUCCESS;
- ds_profile* profile = NULL;
- unsigned int next;
- unsigned int i;
-
- if (p == NULL) return DS_INVALID_PROFILE;
-
- profile = (ds_profile*)malloc(sizeof(ds_profile));
- if (profile == NULL) return DS_MEMORY_ERROR;
-
- memset(profile, 0, sizeof(ds_profile));
-
- clGetPlatformIDs(0, NULL, &numPlatforms);
- if (numPlatforms != 0)
- {
- platforms = (cl_platform_id*)malloc(numPlatforms * sizeof(cl_platform_id));
- if (platforms == NULL)
- {
- status = DS_MEMORY_ERROR;
- goto cleanup;
- }
- clGetPlatformIDs(numPlatforms, platforms, NULL);
- }
-
- numDevices = 0;
- for (i = 0; i < (unsigned int)numPlatforms; i++)
- {
- cl_uint num;
- clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, 0, NULL, &num);
- numDevices += num;
- }
- if (numDevices != 0)
- {
- devices = (cl_device_id*)malloc(numDevices * sizeof(cl_device_id));
- if (devices == NULL)
- {
- status = DS_MEMORY_ERROR;
- goto cleanup;
- }
- }
-
- profile->numDevices = numDevices + 1; // +1 to numDevices to include the native CPU
- profile->devices = (ds_device*)malloc(profile->numDevices * sizeof(ds_device));
- if (profile->devices == NULL)
- {
- profile->numDevices = 0;
- status = DS_MEMORY_ERROR;
- goto cleanup;
- }
- memset(profile->devices, 0, profile->numDevices * sizeof(ds_device));
-
- next = 0;
- for (i = 0; i < (unsigned int)numPlatforms; i++)
- {
- cl_uint num;
- unsigned j;
- char vendor[256];
- if (clGetPlatformInfo(platforms[i], CL_PLATFORM_VENDOR, sizeof(vendor), vendor, NULL) != CL_SUCCESS)
- vendor[0] = '\0';
- clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, numDevices, devices, &num);
- for (j = 0; j < num; j++, next++)
- {
- char buffer[DS_DEVICE_NAME_LENGTH];
- size_t length;
-
- profile->devices[next].type = DS_DEVICE_OPENCL_DEVICE;
- profile->devices[next].oclDeviceID = devices[j];
-
- profile->devices[next].oclPlatformVendor = strdup(vendor);
-
- clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DEVICE_NAME
- , DS_DEVICE_NAME_LENGTH, &buffer, NULL);
- length = strlen(buffer);
- profile->devices[next].oclDeviceName = (char*)malloc(length + 1);
- memcpy(profile->devices[next].oclDeviceName, buffer, length + 1);
-
- clGetDeviceInfo(profile->devices[next].oclDeviceID, CL_DRIVER_VERSION
- , DS_DEVICE_NAME_LENGTH, &buffer, NULL);
- length = strlen(buffer);
- profile->devices[next].oclDriverVersion = (char*)malloc(length + 1);
- memcpy(profile->devices[next].oclDriverVersion, buffer, length + 1);
- }
- }
- profile->devices[next].type = DS_DEVICE_NATIVE_CPU;
- profile->version = version;
-
-cleanup:
- if (platforms) free(platforms);
- if (devices) free(devices);
- if (status == DS_SUCCESS)
- {
- *p = profile;
- }
- else
- {
- if (profile)
- {
- if (profile->devices) free(profile->devices);
- free(profile);
- }
- }
- return status;
-}
-
-// Pointer to a function that calculates the score of a device (ex: device->score)
-// update the data size of score. The encoding and the format of the score data
-// is implementation defined. The function should return DS_SUCCESS if there's no error to be reported.
-typedef ds_status(* ds_perf_evaluator)(ds_device* device, void* data);
-
-typedef enum {
- DS_EVALUATE_ALL
- , DS_EVALUATE_NEW_ONLY
-} ds_evaluation_type;
-
-inline ds_status profileDevices(ds_profile* profile, const ds_evaluation_type type,
- ds_perf_evaluator evaluator, void* evaluatorData, unsigned int* numUpdates)
-{
- ds_status status = DS_SUCCESS;
- unsigned int i;
- unsigned int updates = 0;
-
- if (profile == NULL)
- {
- return DS_INVALID_PROFILE;
- }
- if (evaluator == NULL)
- {
- return DS_INVALID_PERF_EVALUATOR;
- }
-
- for (i = 0; i < profile->numDevices; i++)
- {
- ds_status evaluatorStatus;
-
- switch (type)
- {
- case DS_EVALUATE_NEW_ONLY:
- if (profile->devices[i].score != NULL) break;
- // else fall through
- case DS_EVALUATE_ALL:
- evaluatorStatus = evaluator(profile->devices + i, evaluatorData);
- if (evaluatorStatus != DS_SUCCESS)
- {
- status = evaluatorStatus;
- return status;
- }
- updates++;
- break;
- default:
- return DS_INVALID_PERF_EVALUATOR_TYPE;
- break;
- };
- }
- if (numUpdates) *numUpdates = updates;
- return status;
-}
-
-
-#define DS_TAG_VERSION "<version>"
-#define DS_TAG_VERSION_END "</version>"
-#define DS_TAG_DEVICE "<device>"
-#define DS_TAG_DEVICE_END "</device>"
-#define DS_TAG_SCORE "<score>"
-#define DS_TAG_SCORE_END "</score>"
-#define DS_TAG_DEVICE_TYPE "<type>"
-#define DS_TAG_DEVICE_TYPE_END "</type>"
-#define DS_TAG_DEVICE_NAME "<name>"
-#define DS_TAG_DEVICE_NAME_END "</name>"
-#define DS_TAG_DEVICE_DRIVER_VERSION "<driver>"
-#define DS_TAG_DEVICE_DRIVER_VERSION_END "</driver>"
-
-#define DS_DEVICE_NATIVE_CPU_STRING "native_cpu"
-
-typedef ds_status(* ds_score_serializer)(ds_device* device, void** serializedScore, unsigned int* serializedScoreSize);
-inline ds_status writeProfileToFile(ds_profile* profile, ds_score_serializer serializer, const char* file)
-{
- ds_status status = DS_SUCCESS;
- FILE* profileFile = NULL;
-
-
- if (profile == NULL) return DS_INVALID_PROFILE;
-
- profileFile = fopen(file, "wb");
- if (profileFile == NULL)
- {
- status = DS_FILE_ERROR;
- }
- else
- {
- unsigned int i;
-
- // write version string
- fwrite(DS_TAG_VERSION, sizeof(char), strlen(DS_TAG_VERSION), profileFile);
- fwrite(profile->version, sizeof(char), strlen(profile->version), profileFile);
- fwrite(DS_TAG_VERSION_END, sizeof(char), strlen(DS_TAG_VERSION_END), profileFile);
- fwrite("\n", sizeof(char), 1, profileFile);
-
- for (i = 0; i < profile->numDevices && status == DS_SUCCESS; i++)
- {
- void* serializedScore;
- unsigned int serializedScoreSize;
-
- fwrite(DS_TAG_DEVICE, sizeof(char), strlen(DS_TAG_DEVICE), profileFile);
-
- fwrite(DS_TAG_DEVICE_TYPE, sizeof(char), strlen(DS_TAG_DEVICE_TYPE), profileFile);
- fwrite(&profile->devices[i].type, sizeof(ds_device_type), 1, profileFile);
- fwrite(DS_TAG_DEVICE_TYPE_END, sizeof(char), strlen(DS_TAG_DEVICE_TYPE_END), profileFile);
-
- switch (profile->devices[i].type)
- {
- case DS_DEVICE_NATIVE_CPU:
- {
- // There's no need to emit a device name for the native CPU device.
- /*
- fwrite(DS_TAG_DEVICE_NAME, sizeof(char), strlen(DS_TAG_DEVICE_NAME), profileFile);
- fwrite(DS_DEVICE_NATIVE_CPU_STRING,sizeof(char),strlen(DS_DEVICE_NATIVE_CPU_STRING), profileFile);
- fwrite(DS_TAG_DEVICE_NAME_END, sizeof(char), strlen(DS_TAG_DEVICE_NAME_END), profileFile);
- */
- }
- break;
- case DS_DEVICE_OPENCL_DEVICE:
- {
- fwrite(DS_TAG_DEVICE_NAME, sizeof(char), strlen(DS_TAG_DEVICE_NAME), profileFile);
- fwrite(profile->devices[i].oclDeviceName, sizeof(char), strlen(profile->devices[i].oclDeviceName), profileFile);
- fwrite(DS_TAG_DEVICE_NAME_END, sizeof(char), strlen(DS_TAG_DEVICE_NAME_END), profileFile);
-
- fwrite(DS_TAG_DEVICE_DRIVER_VERSION, sizeof(char), strlen(DS_TAG_DEVICE_DRIVER_VERSION), profileFile);
- fwrite(profile->devices[i].oclDriverVersion, sizeof(char), strlen(profile->devices[i].oclDriverVersion), profileFile);
- fwrite(DS_TAG_DEVICE_DRIVER_VERSION_END, sizeof(char), strlen(DS_TAG_DEVICE_DRIVER_VERSION_END), profileFile);
- }
- break;
- default:
- break;
- };
-
- fwrite(DS_TAG_SCORE, sizeof(char), strlen(DS_TAG_SCORE), profileFile);
- status = serializer(profile->devices + i, &serializedScore, &serializedScoreSize);
- if (status == DS_SUCCESS && serializedScore != NULL && serializedScoreSize > 0)
- {
- fwrite(serializedScore, sizeof(char), serializedScoreSize, profileFile);
- free(serializedScore);
- }
- fwrite(DS_TAG_SCORE_END, sizeof(char), strlen(DS_TAG_SCORE_END), profileFile);
- fwrite(DS_TAG_DEVICE_END, sizeof(char), strlen(DS_TAG_DEVICE_END), profileFile);
- fwrite("\n", sizeof(char), 1, profileFile);
- }
- fclose(profileFile);
- }
- return status;
-}
-
-
-inline ds_status readProFile(const char* fileName, char** content, size_t* contentSize)
-{
- FILE* input = NULL;
- size_t size = 0;
- char* binary = NULL;
- long pos = -1;
-
- *contentSize = 0;
- *content = NULL;
-
- input = fopen(fileName, "rb");
- if (input == NULL)
- {
- return DS_FILE_ERROR;
- }
-
- fseek(input, 0L, SEEK_END);
- pos = ftell(input);
- if (pos < 0)
- {
- fclose(input);
- return DS_FILE_ERROR;
- }
-
- size = pos;
- rewind(input);
- binary = (char*)malloc(size);
- if (binary == NULL)
- {
- fclose(input);
- return DS_FILE_ERROR;
- }
- size_t bytesRead = fread(binary, sizeof(char), size, input);
- (void) bytesRead; // avoid warning
- fclose(input);
-
- *contentSize = size;
- *content = binary;
- return DS_SUCCESS;
-}
-
-
-inline const char* findString(const char* contentStart, const char* contentEnd, const char* string)
-{
- size_t stringLength;
- const char* currentPosition;
- const char* found;
- found = NULL;
- stringLength = strlen(string);
- currentPosition = contentStart;
- for (currentPosition = contentStart; currentPosition < contentEnd; currentPosition++)
- {
- if (*currentPosition == string[0])
- {
- if (currentPosition + stringLength < contentEnd)
- {
- if (strncmp(currentPosition, string, stringLength) == 0)
- {
- found = currentPosition;
- break;
- }
- }
- }
- }
- return found;
-}
-
-
-typedef ds_status(* ds_score_deserializer)(ds_device* device, const unsigned char* serializedScore, unsigned int serializedScoreSize);
-inline ds_status readProfileFromFile(ds_profile* profile, ds_score_deserializer deserializer, const char* file)
-{
-
- ds_status status = DS_SUCCESS;
- char* contentStart = NULL;
- const char* contentEnd = NULL;
- size_t contentSize;
-
- if (profile == NULL) return DS_INVALID_PROFILE;
-
- status = readProFile(file, &contentStart, &contentSize);
- if (status == DS_SUCCESS)
- {
- const char* currentPosition;
- const char* dataStart;
- const char* dataEnd;
- size_t versionStringLength;
-
- contentEnd = contentStart + contentSize;
- currentPosition = contentStart;
-
-
- // parse the version string
- dataStart = findString(currentPosition, contentEnd, DS_TAG_VERSION);
- if (dataStart == NULL)
- {
- status = DS_PROFILE_FILE_ERROR;
- goto cleanup;
- }
- dataStart += strlen(DS_TAG_VERSION);
-
- dataEnd = findString(dataStart, contentEnd, DS_TAG_VERSION_END);
- if (dataEnd == NULL)
- {
- status = DS_PROFILE_FILE_ERROR;
- goto cleanup;
- }
-
- versionStringLength = strlen(profile->version);
- if (versionStringLength != static_cast<size_t>(dataEnd - dataStart)
- || strncmp(profile->version, dataStart, versionStringLength) != 0)
- {
- // version mismatch
- status = DS_PROFILE_FILE_ERROR;
- goto cleanup;
- }
- currentPosition = dataEnd + strlen(DS_TAG_VERSION_END);
-
- // parse the device information
- while (true)
- {
- unsigned int i;
-
- const char* deviceTypeStart;
- const char* deviceTypeEnd;
- ds_device_type deviceType;
-
- const char* deviceNameStart;
- const char* deviceNameEnd;
-
- const char* deviceScoreStart;
- const char* deviceScoreEnd;
-
- const char* deviceDriverStart;
- const char* deviceDriverEnd;
-
- dataStart = findString(currentPosition, contentEnd, DS_TAG_DEVICE);
- if (dataStart == NULL)
- {
- // nothing useful remain, quit...
- break;
- }
- dataStart += strlen(DS_TAG_DEVICE);
- dataEnd = findString(dataStart, contentEnd, DS_TAG_DEVICE_END);
- if (dataEnd == NULL)
- {
- status = DS_PROFILE_FILE_ERROR;
- goto cleanup;
- }
-
- // parse the device type
- deviceTypeStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_TYPE);
- if (deviceTypeStart == NULL)
- {
- status = DS_PROFILE_FILE_ERROR;
- goto cleanup;
- }
- deviceTypeStart += strlen(DS_TAG_DEVICE_TYPE);
- deviceTypeEnd = findString(deviceTypeStart, contentEnd, DS_TAG_DEVICE_TYPE_END);
- if (deviceTypeEnd == NULL)
- {
- status = DS_PROFILE_FILE_ERROR;
- goto cleanup;
- }
- memcpy(&deviceType, deviceTypeStart, sizeof(ds_device_type));
-
-
- // parse the device name
- if (deviceType == DS_DEVICE_OPENCL_DEVICE)
- {
-
- deviceNameStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_NAME);
- if (deviceNameStart == NULL)
- {
- status = DS_PROFILE_FILE_ERROR;
- goto cleanup;
- }
- deviceNameStart += strlen(DS_TAG_DEVICE_NAME);
- deviceNameEnd = findString(deviceNameStart, contentEnd, DS_TAG_DEVICE_NAME_END);
- if (deviceNameEnd == NULL)
- {
- status = DS_PROFILE_FILE_ERROR;
- goto cleanup;
- }
-
-
- deviceDriverStart = findString(dataStart, contentEnd, DS_TAG_DEVICE_DRIVER_VERSION);
- if (deviceDriverStart == NULL)
- {
- status = DS_PROFILE_FILE_ERROR;
- goto cleanup;
- }
- deviceDriverStart += strlen(DS_TAG_DEVICE_DRIVER_VERSION);
- deviceDriverEnd = findString(deviceDriverStart, contentEnd, DS_TAG_DEVICE_DRIVER_VERSION_END);
- if (deviceDriverEnd == NULL)
- {
- status = DS_PROFILE_FILE_ERROR;
- goto cleanup;
- }
-
-
- // check if this device is on the system
- for (i = 0; i < profile->numDevices; i++)
- {
- if (profile->devices[i].type == DS_DEVICE_OPENCL_DEVICE)
- {
- size_t actualDeviceNameLength;
- size_t driverVersionLength;
-
- actualDeviceNameLength = strlen(profile->devices[i].oclDeviceName);
- driverVersionLength = strlen(profile->devices[i].oclDriverVersion);
- if (actualDeviceNameLength == static_cast<size_t>(deviceNameEnd - deviceNameStart)
- && driverVersionLength == static_cast<size_t>(deviceDriverEnd - deviceDriverStart)
- && strncmp(profile->devices[i].oclDeviceName, deviceNameStart, actualDeviceNameLength) == 0
- && strncmp(profile->devices[i].oclDriverVersion, deviceDriverStart, driverVersionLength) == 0)
- {
-
- deviceScoreStart = findString(dataStart, contentEnd, DS_TAG_SCORE);
- if (deviceScoreStart == NULL)
- {
- status = DS_PROFILE_FILE_ERROR;
- goto cleanup;
- }
- deviceScoreStart += strlen(DS_TAG_SCORE);
- deviceScoreEnd = findString(deviceScoreStart, contentEnd, DS_TAG_SCORE_END);
- status = deserializer(profile->devices + i, (const unsigned char*)deviceScoreStart, deviceScoreEnd - deviceScoreStart);
- if (status != DS_SUCCESS)
- {
- goto cleanup;
- }
- }
- }
- }
-
- }
- else if (deviceType == DS_DEVICE_NATIVE_CPU)
- {
- for (i = 0; i < profile->numDevices; i++)
- {
- if (profile->devices[i].type == DS_DEVICE_NATIVE_CPU)
- {
- deviceScoreStart = findString(dataStart, contentEnd, DS_TAG_SCORE);
- if (deviceScoreStart == NULL)
- {
- status = DS_PROFILE_FILE_ERROR;
- goto cleanup;
- }
- deviceScoreStart += strlen(DS_TAG_SCORE);
- deviceScoreEnd = findString(deviceScoreStart, contentEnd, DS_TAG_SCORE_END);
- status = deserializer(profile->devices + i, (const unsigned char*)deviceScoreStart, deviceScoreEnd - deviceScoreStart);
- if (status != DS_SUCCESS)
- {
- goto cleanup;
- }
- }
- }
- }
-
- // skip over the current one to find the next device
- currentPosition = dataEnd + strlen(DS_TAG_DEVICE_END);
- }
- }
-cleanup:
- if (contentStart != NULL) free(contentStart);
- if (status != DS_SUCCESS)
- return status;
-
- // Check that all the devices present had valid cached scores. If
- // not, return DS_INVALID_PROFILE and let the caller re-evaluate
- // scores for present devices, and write a new profile file.
- for (unsigned int i = 0; i < profile->numDevices; i++)
- if (profile->devices[i].score == NULL)
- return DS_INVALID_PROFILE;
-
- return DS_SUCCESS;
-}
-
-inline ds_status getNumDeviceWithEmptyScore(ds_profile* profile, unsigned int* num)
-{
- unsigned int i;
- if (profile == NULL || num == NULL) return DS_MEMORY_ERROR;
- *num = 0;
- for (i = 0; i < profile->numDevices; i++)
- {
- if (profile->devices[i].score == NULL)
- {
- (*num)++;
- }
- }
- return DS_SUCCESS;
-}
-
-#endif
-
-/* vim:set shiftwidth=4 softtabstop=4 expandtab: */
diff --git a/sc/source/core/opencl/openclwrapper.cxx b/sc/source/core/opencl/openclwrapper.cxx
deleted file mode 100644
index 20f4919c88c0..000000000000
--- a/sc/source/core/opencl/openclwrapper.cxx
+++ /dev/null
@@ -1,806 +0,0 @@
-/* -*- 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 <config_folders.h>
-
-#include "calcconfig.hxx"
-#include "interpre.hxx"
-#include "opencl_device.hxx"
-#include "openclwrapper.hxx"
-
-#include <comphelper/string.hxx>
-#include <opencl/openclconfig.hxx>
-#include <osl/file.hxx>
-#include <rtl/bootstrap.hxx>
-#include <rtl/digest.h>
-#include <rtl/strbuf.hxx>
-#include <rtl/ustring.hxx>
-#include <sal/config.h>
-
-#include <boost/scoped_array.hpp>
-#include <unicode/regex.h>
-
-#include <stdio.h>
-#include <stdlib.h>
-#include <string.h>
-
-#include <cmath>
-
-#ifdef _WIN32
-#include <prewin.h>
-#include <postwin.h>
-#define OPENCL_DLL_NAME "OpenCL.dll"
-#elif defined(MACOSX)
-#define OPENCL_DLL_NAME NULL
-#else
-#define OPENCL_DLL_NAME "libOpenCL.so"
-#endif
-
-#define DEVICE_NAME_LENGTH 1024
-#define DRIVER_VERSION_LENGTH 1024
-#define PLATFORM_VERSION_LENGTH 1024
-
-using namespace std;
-
-namespace sc { namespace opencl {
-
-GPUEnv OpenCLDevice::gpuEnv;
-bool OpenCLDevice::bIsInited = false;
-
-namespace {
-
-OString generateMD5(const void* pData, size_t length)
-{
- sal_uInt8 pBuffer[RTL_DIGEST_LENGTH_MD5];
- rtlDigestError aError = rtl_digest_MD5(pData, length,
- pBuffer, RTL_DIGEST_LENGTH_MD5);
- SAL_WARN_IF(aError != rtl_Digest_E_None, "sc", "md5 generation failed");
-
- OStringBuffer aBuffer;
- const char* pString = "0123456789ABCDEF";
- for(size_t i = 0; i < RTL_DIGEST_LENGTH_MD5; ++i)
- {
- sal_uInt8 val = pBuffer[i];
- aBuffer.append(pString[val/16]);
- aBuffer.append(pString[val%16]);
- }
- return aBuffer.makeStringAndClear();
-}
-
-OString getCacheFolder()
-{
- OUString url("${$BRAND_BASE_DIR/" LIBO_ETC_FOLDER "/" SAL_CONFIGFILE("bootstrap") ":UserInstallation}/cache/");
- rtl::Bootstrap::expandMacros(url);
-
- osl::Directory::create(url);
-
- return rtl::OUStringToOString(url, RTL_TEXTENCODING_UTF8);
-}
-
-}
-
-OString OpenCLDevice::maCacheFolder = getCacheFolder();
-
-void OpenCLDevice::setKernelEnv( KernelEnv *envInfo )
-{
- envInfo->mpkContext = gpuEnv.mpContext;
- envInfo->mpkCmdQueue = gpuEnv.mpCmdQueue;
- envInfo->mpkProgram = gpuEnv.mpArryPrograms[0];
-}
-
-namespace {
-
-OString createFileName(cl_device_id deviceId, const char* clFileName)
-{
- OString fileName(clFileName);
- sal_Int32 nIndex = fileName.lastIndexOf(".cl");
- if(nIndex > 0)
- fileName = fileName.copy(0, nIndex);
-
- char deviceName[DEVICE_NAME_LENGTH] = {0};
- clGetDeviceInfo(deviceId, CL_DEVICE_NAME,
- sizeof(deviceName), deviceName, NULL);
-
- char driverVersion[DRIVER_VERSION_LENGTH] = {0};
- clGetDeviceInfo(deviceId, CL_DRIVER_VERSION,
- sizeof(driverVersion), driverVersion, NULL);
-
- cl_platform_id platformId;
- clGetDeviceInfo(deviceId, CL_DEVICE_PLATFORM,
- sizeof(platformId), &platformId, NULL);
-
- char platformVersion[PLATFORM_VERSION_LENGTH] = {0};
- clGetPlatformInfo(platformId, CL_PLATFORM_VERSION, sizeof(platformVersion),
- platformVersion, NULL);
-
- // create hash for deviceName + driver version + platform version
- OString aString = OString(deviceName) + driverVersion + platformVersion;
- OString aHash = generateMD5(aString.getStr(), aString.getLength());
-
- return OpenCLDevice::maCacheFolder + fileName + "-" +
- aHash + ".bin";
-}
-
-}
-
-std::vector<boost::shared_ptr<osl::File> > OpenCLDevice::binaryGenerated( const char * clFileName, cl_context context )
-{
- size_t numDevices=0;
-
- std::vector<boost::shared_ptr<osl::File> > aGeneratedFiles;
- cl_int clStatus = clGetContextInfo( context, CL_CONTEXT_DEVICES,
- 0, NULL, &numDevices );
- numDevices /= sizeof(numDevices);
-
- if(clStatus != CL_SUCCESS)
- return aGeneratedFiles;
-
- // grab the handles to all of the devices in the context.
- boost::scoped_array<cl_device_id> pArryDevsID(new cl_device_id[numDevices]);
- clStatus = clGetContextInfo( context, CL_CONTEXT_DEVICES,
- sizeof( cl_device_id ) * numDevices, pArryDevsID.get(), NULL );
-
- if(clStatus != CL_SUCCESS)
- return aGeneratedFiles;
-
- for ( size_t i = 0; i < numDevices; i++ )
- {
- if ( pArryDevsID[i] != 0 )
- {
- OString fileName = createFileName(gpuEnv.mpArryDevsID[i], clFileName);
- osl::File* pNewFile = new osl::File(rtl::OStringToOUString(fileName, RTL_TEXTENCODING_UTF8));
- if(pNewFile->open(osl_File_OpenFlag_Read) == osl::FileBase::E_None)
- {
- aGeneratedFiles.push_back(boost::shared_ptr<osl::File>(pNewFile));
- SAL_INFO("sc.opencl.file", "Opening binary file '" << fileName << "' for reading: success");
- }
- else
- {
- SAL_INFO("sc.opencl.file", "Opening binary file '" << fileName << "' for reading: FAIL");
- delete pNewFile;
- break;
- }
- }
- }
-
- return aGeneratedFiles;
-}
-
-bool OpenCLDevice::writeBinaryToFile( const OString& rFileName, const char* binary, size_t numBytes )
-{
- osl::File file(rtl::OStringToOUString(rFileName, RTL_TEXTENCODING_UTF8));
- osl::FileBase::RC status = file.open(
- osl_File_OpenFlag_Write | osl_File_OpenFlag_Create );
-
- if(status != osl::FileBase::E_None)
- return false;
-
- sal_uInt64 nBytesWritten = 0;
- file.write( binary, numBytes, nBytesWritten );
-
- assert(numBytes == nBytesWritten);
-
- return true;
-}
-
-bool OpenCLDevice::generatBinFromKernelSource( cl_program program, const char * clFileName )
-{
- cl_uint numDevices;
-
- cl_int clStatus = clGetProgramInfo( program, CL_PROGRAM_NUM_DEVICES,
- sizeof(numDevices), &numDevices, NULL );
- CHECK_OPENCL( clStatus, "clGetProgramInfo" );
-
- std::vector<cl_device_id> pArryDevsID(numDevices);
- /* grab the handles to all of the devices in the program. */
- clStatus = clGetProgramInfo( program, CL_PROGRAM_DEVICES,
- sizeof(cl_device_id) * numDevices, &pArryDevsID[0], NULL );
- CHECK_OPENCL( clStatus, "clGetProgramInfo" );
-
- /* figure out the sizes of each of the binaries. */
- std::vector<size_t> binarySizes(numDevices);
-
- clStatus = clGetProgramInfo( program, CL_PROGRAM_BINARY_SIZES,
- sizeof(size_t) * numDevices, &binarySizes[0], NULL );
- CHECK_OPENCL( clStatus, "clGetProgramInfo" );
-
- /* copy over all of the generated binaries. */
- boost::scoped_array<char*> binaries(new char*[numDevices]);
-
- for ( size_t i = 0; i < numDevices; i++ )
- {
- if ( binarySizes[i] != 0 )
- {
- binaries[i] = new char[binarySizes[i]];
- }
- else
- {
- binaries[i] = NULL;
- }
- }
-
- clStatus = clGetProgramInfo( program, CL_PROGRAM_BINARIES,
- sizeof(char *) * numDevices, binaries.get(), NULL );
- CHECK_OPENCL(clStatus,"clGetProgramInfo");
-
- /* dump out each binary into its own separate file. */
- for ( size_t i = 0; i < numDevices; i++ )
- {
-
- if ( binarySizes[i] != 0 )
- {
- OString fileName = createFileName(pArryDevsID[i], clFileName);
- if ( !writeBinaryToFile( fileName,
- binaries[i], binarySizes[i] ) )
- SAL_INFO("sc.opencl.file", "Writing binary file '" << fileName << "': FAIL");
- else
- SAL_INFO("sc.opencl.file", "Writing binary file '" << fileName << "': success");
- }
- }
-
- // Release all resouces and memory
- for ( size_t i = 0; i < numDevices; i++ )
- {
- delete[] binaries[i];
- }
-
- return true;
-}
-
-bool OpenCLDevice::initOpenCLAttr( OpenCLEnv * env )
-{
- if ( gpuEnv.mnIsUserCreated )
- return true;
-
- gpuEnv.mpContext = env->mpOclContext;
- gpuEnv.mpPlatformID = env->mpOclPlatformID;
- gpuEnv.mpDevID = env->mpOclDevsID;
- gpuEnv.mpCmdQueue = env->mpOclCmdQueue;
-
- gpuEnv.mnIsUserCreated = 1;
-
- return false;
-}
-
-void OpenCLDevice::releaseOpenCLEnv( GPUEnv *gpuInfo )
-{
- if ( !bIsInited )
- {
- return;
- }
-
- if ( gpuEnv.mpCmdQueue )
- {
- clReleaseCommandQueue( gpuEnv.mpCmdQueue );
- gpuEnv.mpCmdQueue = NULL;
- }
- if ( gpuEnv.mpContext )
- {
- clReleaseContext( gpuEnv.mpContext );
- gpuEnv.mpContext = NULL;
- }
- bIsInited = false;
- gpuInfo->mnIsUserCreated = 0;
- free( gpuInfo->mpArryDevsID );
-
- return;
-}
-
-namespace {
-
-bool buildProgram(const char* buildOption, GPUEnv* gpuInfo, int idx)
-{
- cl_int clStatus;
- //char options[512];
- // create a cl program executable for all the devices specified
- if (!gpuInfo->mnIsUserCreated)
- {
- clStatus = clBuildProgram(gpuInfo->mpArryPrograms[idx], 1, gpuInfo->mpArryDevsID,
- buildOption, NULL, NULL);
- }
- else
- {
- clStatus = clBuildProgram(gpuInfo->mpArryPrograms[idx], 1, &(gpuInfo->mpDevID),
- buildOption, NULL, NULL);
- }
-
- if ( clStatus != CL_SUCCESS )
- {
- size_t length;
- if ( !gpuInfo->mnIsUserCreated )
- {
- clStatus = clGetProgramBuildInfo( gpuInfo->mpArryPrograms[idx], gpuInfo->mpArryDevsID[0],
- CL_PROGRAM_BUILD_LOG, 0, NULL, &length );
- }
- else
- {
- clStatus = clGetProgramBuildInfo( gpuInfo->mpArryPrograms[idx], gpuInfo->mpDevID,
- CL_PROGRAM_BUILD_LOG, 0, NULL, &length);
- }
- if ( clStatus != CL_SUCCESS )
- {
- return false;
- }
-
- boost::scoped_array<char> buildLog(new char[length]);
- if ( !gpuInfo->mnIsUserCreated )
- {
- clStatus = clGetProgramBuildInfo( gpuInfo->mpArryPrograms[idx], gpuInfo->mpArryDevsID[0],
- CL_PROGRAM_BUILD_LOG, length, buildLog.get(), &length );
- }
- else
- {
- clStatus = clGetProgramBuildInfo( gpuInfo->mpArryPrograms[idx], gpuInfo->mpDevID,
- CL_PROGRAM_BUILD_LOG, length, buildLog.get(), &length );
- }
- if ( clStatus != CL_SUCCESS )
- {
- return false;
- }
-
- OString aBuildLogFileURL = OpenCLDevice::maCacheFolder + "kernel-build.log";
- osl::File aBuildLogFile(rtl::OStringToOUString(aBuildLogFileURL, RTL_TEXTENCODING_UTF8));
- osl::FileBase::RC status = aBuildLogFile.open(
- osl_File_OpenFlag_Write | osl_File_OpenFlag_Create );
-
- if(status != osl::FileBase::E_None)
- return false;
-
- sal_uInt64 nBytesWritten = 0;
- aBuildLogFile.write( buildLog.get(), length, nBytesWritten );
-
- return false;
- }
-
- return true;
-}
-
-}
-
-bool OpenCLDevice::buildProgramFromBinary(const char* buildOption, GPUEnv* gpuInfo, const char* filename, int idx)
-{
- size_t numDevices;
- cl_int clStatus = clGetContextInfo( gpuInfo->mpContext, CL_CONTEXT_DEVICES,
- 0, NULL, &numDevices );
- numDevices /= sizeof(numDevices);
- CHECK_OPENCL( clStatus, "clGetContextInfo" );
-
- std::vector<boost::shared_ptr<osl::File> > aGeneratedFiles = binaryGenerated(
- filename, gpuInfo->mpContext );
-
- if (aGeneratedFiles.size() == numDevices)
- {
- boost::scoped_array<size_t> length(new size_t[numDevices]);
- boost::scoped_array<unsigned char*> pBinary(new unsigned char*[numDevices]);
- for(size_t i = 0; i < numDevices; ++i)
- {
- sal_uInt64 nSize;
- aGeneratedFiles[i]->getSize(nSize);
- unsigned char* binary = new unsigned char[nSize];
- sal_uInt64 nBytesRead;
- aGeneratedFiles[i]->read(binary, nSize, nBytesRead);
- if(nSize != nBytesRead)
- assert(false);
-
- length[i] = nBytesRead;
-
- pBinary[i] = binary;
- }
-
- // grab the handles to all of the devices in the context.
- boost::scoped_array<cl_device_id> pArryDevsID(new cl_device_id[numDevices]);
- clStatus = clGetContextInfo( gpuInfo->mpContext, CL_CONTEXT_DEVICES,
- sizeof( cl_device_id ) * numDevices, pArryDevsID.get(), NULL );
-
- if(clStatus != CL_SUCCESS)
- {
- for(size_t i = 0; i < numDevices; ++i)
- {
- delete[] pBinary[i];
- }
- return false;
- }
-
- cl_int binary_status;
-
- gpuInfo->mpArryPrograms[idx] = clCreateProgramWithBinary( gpuInfo->mpContext,numDevices,
- pArryDevsID.get(), length.get(), (const unsigned char**) pBinary.get(),
- &binary_status, &clStatus );
- if(clStatus != CL_SUCCESS)
- {
- // something went wrong, fall back to compiling from source
- return false;
- }
- for(size_t i = 0; i < numDevices; ++i)
- {
- delete[] pBinary[i];
- }
- }
-
- if ( !gpuInfo->mpArryPrograms[idx] )
- {
- return false;
- }
- return buildProgram(buildOption, gpuInfo, idx);
-}
-
-bool OpenCLDevice::initOpenCLRunEnv( int argc )
-{
- if ( ( argc > MAX_CLFILE_NUM ) || ( argc < 0 ) )
- return true;
-
- if ( !bIsInited )
- {
- if ( !gpuEnv.mnIsUserCreated )
- memset( &gpuEnv, 0, sizeof(gpuEnv) );
-
- //initialize devices, context, command_queue
- bool status = initOpenCLRunEnv( &gpuEnv );
- if ( status )
- {
- return true;
- }
- //initialize program, kernelName, kernelCount
- if( getenv( "SC_FLOAT" ) )
- {
- gpuEnv.mnKhrFp64Flag = false;
- gpuEnv.mnAmdFp64Flag = false;
- }
- if( gpuEnv.mnKhrFp64Flag )
- {
- SAL_INFO("sc.opencl", "Use Khr double");
- }
- else if( gpuEnv.mnAmdFp64Flag )
- {
- SAL_INFO("sc.opencl", "Use AMD double type");
- }
- else
- {
- SAL_INFO("sc.opencl", "USE float type");
- }
- bIsInited = true;
- }
- return false;
-}
-
-namespace {
-
-void checkDeviceForDoubleSupport(cl_device_id deviceId, bool& bKhrFp64, bool& bAmdFp64)
-{
- bKhrFp64 = false;
- bAmdFp64 = false;
-
- // Check device extensions for double type
- size_t aDevExtInfoSize = 0;
-
- cl_uint clStatus = clGetDeviceInfo( deviceId, CL_DEVICE_EXTENSIONS, 0, NULL, &aDevExtInfoSize );
- if( clStatus != CL_SUCCESS )
- return;
-
- boost::scoped_array<char> pExtInfo(new char[aDevExtInfoSize]);
-
- clStatus = clGetDeviceInfo( deviceId, CL_DEVICE_EXTENSIONS,
- sizeof(char) * aDevExtInfoSize, pExtInfo.get(), NULL);
-
- if( clStatus != CL_SUCCESS )
- return;
-
- if ( strstr( pExtInfo.get(), "cl_khr_fp64" ) )
- {
- bKhrFp64 = true;
- }
- else
- {
- // Check if cl_amd_fp64 extension is supported
- if ( strstr( pExtInfo.get(), "cl_amd_fp64" ) )
- bAmdFp64 = true;
- }
-}
-
-}
-
-bool OpenCLDevice::initOpenCLRunEnv( GPUEnv *gpuInfo )
-{
- bool bKhrFp64 = false;
- bool bAmdFp64 = false;
-
- checkDeviceForDoubleSupport(gpuInfo->mpArryDevsID[0], bKhrFp64, bAmdFp64);
-
- gpuInfo->mnKhrFp64Flag = bKhrFp64;
- gpuInfo->mnAmdFp64Flag = bAmdFp64;
-
- return false;
-}
-
-namespace {
-
-// based on crashes and hanging during kernel compilation
-void createDeviceInfo(cl_device_id aDeviceId, OpenCLPlatformInfo& rPlatformInfo)
-{
- OpenCLDeviceInfo aDeviceInfo;
- aDeviceInfo.device = aDeviceId;
-
- char pName[DEVICE_NAME_LENGTH];
- cl_int nState = clGetDeviceInfo(aDeviceId, CL_DEVICE_NAME, DEVICE_NAME_LENGTH, pName, NULL);
- if(nState != CL_SUCCESS)
- return;
-
- aDeviceInfo.maName = OUString::createFromAscii(pName);
-
- char pVendor[DEVICE_NAME_LENGTH];
- nState = clGetDeviceInfo(aDeviceId, CL_DEVICE_VENDOR, DEVICE_NAME_LENGTH, pVendor, NULL);
- if(nState != CL_SUCCESS)
- return;
-
- aDeviceInfo.maVendor = OUString::createFromAscii(pVendor);
-
- cl_ulong nMemSize;
- nState = clGetDeviceInfo(aDeviceId, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(nMemSize), &nMemSize, NULL);
- if(nState != CL_SUCCESS)
- return;
-
- aDeviceInfo.mnMemory = nMemSize;
-
- cl_uint nClockFrequency;
- nState = clGetDeviceInfo(aDeviceId, CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof(nClockFrequency), &nClockFrequency, NULL);
- if(nState != CL_SUCCESS)
- return;
-
- aDeviceInfo.mnFrequency = nClockFrequency;
-
- cl_uint nComputeUnits;
- nState = clGetDeviceInfo(aDeviceId, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(nComputeUnits), &nComputeUnits, NULL);
- if(nState != CL_SUCCESS)
- return;
-
- char pDriver[DEVICE_NAME_LENGTH];
- nState = clGetDeviceInfo(aDeviceId, CL_DRIVER_VERSION, DEVICE_NAME_LENGTH, pDriver, NULL);
-
- if(nState != CL_SUCCESS)
- return;
-
- aDeviceInfo.maDriver = OUString::createFromAscii(pDriver);
-
- bool bKhrFp64 = false;
- bool bAmdFp64 = false;
- checkDeviceForDoubleSupport(aDeviceId, bKhrFp64, bAmdFp64);
-
- // only list devices that support double
- if(!bKhrFp64 && !bAmdFp64)
- return;
-
- aDeviceInfo.mnComputeUnits = nComputeUnits;
-
- if(!OpenCLConfig::get().checkImplementation(rPlatformInfo, aDeviceInfo))
- rPlatformInfo.maDevices.push_back(aDeviceInfo);
-}
-
-bool createPlatformInfo(cl_platform_id nPlatformId, OpenCLPlatformInfo& rPlatformInfo)
-{
- rPlatformInfo.platform = nPlatformId;
- char pName[64];
- cl_int nState = clGetPlatformInfo(nPlatformId, CL_PLATFORM_NAME, 64,
- pName, NULL);
- if(nState != CL_SUCCESS)
- return false;
- rPlatformInfo.maName = OUString::createFromAscii(pName);
-
- char pVendor[64];
- nState = clGetPlatformInfo(nPlatformId, CL_PLATFORM_VENDOR, 64,
- pVendor, NULL);
- if(nState != CL_SUCCESS)
- return false;
-
- rPlatformInfo.maVendor = OUString::createFromAscii(pVendor);
-
- cl_uint nDevices;
- nState = clGetDeviceIDs(nPlatformId, CL_DEVICE_TYPE_ALL, 0, NULL, &nDevices);
- if(nState != CL_SUCCESS)
- return false;
-
- // memory leak that does not matter
- // memory is stored in static variable that lives through the whole program
- cl_device_id* pDevices = new cl_device_id[nDevices];
- nState = clGetDeviceIDs(nPlatformId, CL_DEVICE_TYPE_ALL, nDevices, pDevices, NULL);
- if(nState != CL_SUCCESS)
- return false;
-
- for(size_t i = 0; i < nDevices; ++i)
- {
- createDeviceInfo(pDevices[i], rPlatformInfo);
- }
-
- return true;
-}
-
-}
-
-const std::vector<OpenCLPlatformInfo>& fillOpenCLInfo()
-{
- static std::vector<OpenCLPlatformInfo> aPlatforms;
- if(!aPlatforms.empty())
- return aPlatforms;
-
- int status = clewInit(OPENCL_DLL_NAME);
- if (status < 0)
- return aPlatforms;
-
- cl_uint nPlatforms;
- cl_int nState = clGetPlatformIDs(0, NULL, &nPlatforms);
-
- if(nState != CL_SUCCESS)
- return aPlatforms;
-
- // memory leak that does not matter,
- // memory is stored in static instance aPlatforms
- cl_platform_id* pPlatforms = new cl_platform_id[nPlatforms];
- nState = clGetPlatformIDs(nPlatforms, pPlatforms, NULL);
-
- if(nState != CL_SUCCESS)
- return aPlatforms;
-
- for(size_t i = 0; i < nPlatforms; ++i)
- {
- OpenCLPlatformInfo aPlatformInfo;
- if(createPlatformInfo(pPlatforms[i], aPlatformInfo))
- aPlatforms.push_back(aPlatformInfo);
- }
-
- return aPlatforms;
-}
-
-namespace {
-
-cl_device_id findDeviceIdByDeviceString(const OUString& rString, const std::vector<OpenCLPlatformInfo>& rPlatforms)
-{
- std::vector<OpenCLPlatformInfo>::const_iterator it = rPlatforms.begin(), itEnd = rPlatforms.end();
- for(; it != itEnd; ++it)
- {
- std::vector<OpenCLDeviceInfo>::const_iterator itr = it->maDevices.begin(), itrEnd = it->maDevices.end();
- for(; itr != itrEnd; ++itr)
- {
- OUString aDeviceId = it->maVendor + " " + itr->maName;
- if(rString == aDeviceId)
- {
- return static_cast<cl_device_id>(itr->device);
- }
- }
- }
-
- return NULL;
-}
-
-void findDeviceInfoFromDeviceId(cl_device_id aDeviceId, size_t& rDeviceId, size_t& rPlatformId)
-{
- cl_platform_id platformId;
- cl_int nState = clGetDeviceInfo(aDeviceId, CL_DEVICE_PLATFORM,
- sizeof(platformId), &platformId, NULL);
-
- if(nState != CL_SUCCESS)
- return;
-
- const std::vector<OpenCLPlatformInfo>& rPlatforms = fillOpenCLInfo();
- for(size_t i = 0; i < rPlatforms.size(); ++i)
- {
- cl_platform_id platId = static_cast<cl_platform_id>(rPlatforms[i].platform);
- if(platId != platformId)
- continue;
-
- for(size_t j = 0; j < rPlatforms[i].maDevices.size(); ++j)
- {
- cl_device_id id = static_cast<cl_device_id>(rPlatforms[i].maDevices[j].device);
- if(id == aDeviceId)
- {
- rDeviceId = j;
- rPlatformId = i;
- return;
- }
- }
- }
-}
-
-}
-
-bool switchOpenCLDevice(const OUString* pDevice, bool bAutoSelect, bool bForceEvaluation)
-{
- if(fillOpenCLInfo().empty())
- return false;
-
- cl_device_id pDeviceId = NULL;
- if(pDevice)
- pDeviceId = findDeviceIdByDeviceString(*pDevice, fillOpenCLInfo());
-
- if(!pDeviceId || bAutoSelect)
- {
- int status = clewInit(OPENCL_DLL_NAME);
- if (status < 0)
- return false;
-
- OUString url("${$BRAND_BASE_DIR/" LIBO_ETC_FOLDER "/" SAL_CONFIGFILE("bootstrap") ":UserInstallation}/cache/");
- rtl::Bootstrap::expandMacros(url);
- OUString path;
- osl::FileBase::getSystemPathFromFileURL(url,path);
- OString dsFileName = rtl::OUStringToOString(path, RTL_TEXTENCODING_UTF8);
- ds_device pSelectedDevice = sc::OpenCLDevice::getDeviceSelection(dsFileName.getStr(), bForceEvaluation);
- pDeviceId = pSelectedDevice.oclDeviceID;
-
- }
-
- if(OpenCLDevice::gpuEnv.mpDevID == pDeviceId)
- {
- // we don't need to change anything
- // still the same device
- return pDeviceId != NULL;
- }
-
- cl_platform_id platformId;
- cl_int nState = clGetDeviceInfo(pDeviceId, CL_DEVICE_PLATFORM,
- sizeof(platformId), &platformId, NULL);
-
- cl_context_properties cps[3];
- cps[0] = CL_CONTEXT_PLATFORM;
- cps[1] = reinterpret_cast<cl_context_properties>(platformId);
- cps[2] = 0;
- cl_context context = clCreateContext( cps, 1, &pDeviceId, NULL, NULL, &nState );
-
- if(nState != CL_SUCCESS || context == NULL)
- {
- if(context != NULL)
- clReleaseContext(context);
-
- SAL_WARN("sc", "failed to set/switch opencl device");
- return false;
- }
-
- cl_command_queue command_queue = clCreateCommandQueue(
- context, pDeviceId, 0, &nState);
-
- if(command_queue == NULL || nState != CL_SUCCESS)
- {
- if(command_queue != NULL)
- clReleaseCommandQueue(command_queue);
-
- clReleaseContext(context);
- SAL_WARN("sc", "failed to set/switch opencl device");
- return false;
- }
-
- OpenCLDevice::releaseOpenCLEnv(&OpenCLDevice::gpuEnv);
- OpenCLEnv env;
- env.mpOclPlatformID = platformId;
- env.mpOclContext = context;
- env.mpOclDevsID = pDeviceId;
- env.mpOclCmdQueue = command_queue;
- OpenCLDevice::initOpenCLAttr(&env);
-
- // why do we need this at all?
-
- // (Assuming the above question refers to the mpArryDevsID
- // initialisation below.) Because otherwise the code crashes in
- // initOpenCLRunEnv(). Confused? You should be.
-
- OpenCLDevice::gpuEnv.mpArryDevsID = (cl_device_id*) malloc( sizeof(cl_device_id) );
- OpenCLDevice::gpuEnv.mpArryDevsID[0] = pDeviceId;
-
- return !OpenCLDevice::initOpenCLRunEnv(0);
-}
-
-void getOpenCLDeviceInfo(size_t& rDeviceId, size_t& rPlatformId)
-{
- int status = clewInit(OPENCL_DLL_NAME);
- if (status < 0)
- return;
-
- cl_device_id id = OpenCLDevice::gpuEnv.mpDevID;
- findDeviceInfoFromDeviceId(id, rDeviceId, rPlatformId);
-}
-
-}}
-
-/* vim:set shiftwidth=4 softtabstop=4 expandtab: */
diff --git a/sc/source/core/tool/formulagroup.cxx b/sc/source/core/tool/formulagroup.cxx
index 328e3c3b0083..cc4f80512376 100644
--- a/sc/source/core/tool/formulagroup.cxx
+++ b/sc/source/core/tool/formulagroup.cxx
@@ -35,7 +35,7 @@
#include <cstdio>
#if HAVE_FEATURE_OPENCL
-#include "openclwrapper.hxx"
+#include <opencl/openclwrapper.hxx>
#endif
namespace sc {
@@ -542,7 +542,7 @@ FormulaGroupInterpreter *FormulaGroupInterpreter::getStatic()
void FormulaGroupInterpreter::fillOpenCLInfo(std::vector<OpenCLPlatformInfo>& rPlatforms)
{
const std::vector<OpenCLPlatformInfo>& rPlatformsFromWrapper =
- sc::opencl::fillOpenCLInfo();
+ ::opencl::fillOpenCLInfo();
rPlatforms.assign(rPlatformsFromWrapper.begin(), rPlatformsFromWrapper.end());
}
@@ -564,7 +564,7 @@ bool FormulaGroupInterpreter::switchOpenCLDevice(const OUString& rDeviceId, bool
msInstance = new sc::FormulaGroupInterpreterSoftware();
return true;
}
- bool bSuccess = sc::opencl::switchOpenCLDevice(&rDeviceId, bAutoSelect, bForceEvaluation);
+ bool bSuccess = ::opencl::switchOpenCLDevice(&rDeviceId, bAutoSelect, bForceEvaluation);
if(!bSuccess)
return false;
@@ -591,7 +591,7 @@ void FormulaGroupInterpreter::getOpenCLDeviceInfo(sal_Int32& rDeviceId, sal_Int3
size_t aDeviceId = static_cast<size_t>(-1);
size_t aPlatformId = static_cast<size_t>(-1);
- sc::opencl::getOpenCLDeviceInfo(aDeviceId, aPlatformId);
+ ::opencl::getOpenCLDeviceInfo(aDeviceId, aPlatformId);
rDeviceId = aDeviceId;
rPlatformId = aPlatformId;
}