[Libreoffice-commits] core.git: Branch 'feature/calc-group-interpreter' - 3 commits - config_host.mk.in configure.ac RepositoryExternal.mk sc/CppunitTest_sc_ucalc.mk sc/Library_sc.mk sc/source
Michael Meeks
michael.meeks at suse.com
Mon Jun 17 06:26:21 PDT 2013
RepositoryExternal.mk | 17
config_host.mk.in | 3
configure.ac | 34 +
sc/CppunitTest_sc_ucalc.mk | 3
sc/Library_sc.mk | 8
sc/source/core/opencl/oclkernels.hxx | 116 +++
sc/source/core/opencl/openclwrapper.cxx | 1031 ++++++++++++++++++++++++++++++++
sc/source/core/opencl/openclwrapper.hxx | 172 +++++
sc/source/core/tool/interpr1.cxx | 4
9 files changed, 1388 insertions(+)
New commits:
commit d967e607006c063c36e0cc85e6a527366bd19d81
Author: Michael Meeks <michael.meeks at suse.com>
Date: Sun Jun 16 17:55:11 2013 +0100
OpenCL code cleanups
Change-Id: Iab75b11f13856a6e631587e22b5d76977b8c7448
diff --git a/RepositoryExternal.mk b/RepositoryExternal.mk
index 479c08b..2b3166b 100644
--- a/RepositoryExternal.mk
+++ b/RepositoryExternal.mk
@@ -38,7 +38,7 @@
# External headers
-ifeq ($(ENABLE_OPENGCL),TRUE)
+ifeq ($(ENABLE_OPENCL),TRUE)
define gb_LinkTarget__use_opencl
$(call gb_LinkTarget_set_include,$(1),\
diff --git a/configure.ac b/configure.ac
index f81ab47..821233d 100644
--- a/configure.ac
+++ b/configure.ac
@@ -9688,8 +9688,8 @@ if test "z$with_opencl_sdk" = "z"; then
else
if test -d $with_opencl_sdk/include; then
ENABLE_OPENCL=TRUE
- OPENCL_CFLAGS="-I $with_opencl_sdk/include"
- OPENCL_LIBS="-L $with_opencl_sdk/lib/x86 -lOpenCL"
+ OPENCL_CFLAGS="-I$with_opencl_sdk/include"
+ OPENCL_LIBS="-L$with_opencl_sdk/lib/x86 -lOpenCL"
AC_MSG_RESULT([found at path $with_opencl_sdk])
else
AC_MSG_ERROR([no headers found found at $with_opencl_sdk/include ])
diff --git a/sc/CppunitTest_sc_ucalc.mk b/sc/CppunitTest_sc_ucalc.mk
index acf7bdb..7ed41df 100644
--- a/sc/CppunitTest_sc_ucalc.mk
+++ b/sc/CppunitTest_sc_ucalc.mk
@@ -20,6 +20,9 @@ $(eval $(call gb_CppunitTest_use_library_objects,sc_ucalc,sc))
ifeq ($(ENABLE_TELEPATHY),TRUE)
$(eval $(call gb_CppunitTest_use_libraries,sc_ucalc,tubes))
endif
+ifeq ($(ENABLE_OPENCL),TRUE)
+$(eval $(call gb_CppunitTest_use_externals,sc_ucalc,opencl))
+endif
$(eval $(call gb_CppunitTest_use_externals,sc_ucalc,\
boost_headers \
diff --git a/sc/source/core/opencl/openclwrapper.cxx b/sc/source/core/opencl/openclwrapper.cxx
index 05b169a..6c3935e 100644
--- a/sc/source/core/opencl/openclwrapper.cxx
+++ b/sc/source/core/opencl/openclwrapper.cxx
@@ -11,6 +11,7 @@
#include <stdlib.h>
#include <string.h>
+#include "random.hxx"
#include "openclwrapper.hxx"
#include "oclkernels.hxx"
@@ -1003,10 +1004,8 @@ double OclCalc::OclProcess(cl_kernel_function function, double *data,
double OclCalc::OclTest() {
double data[NUM];
- srand((unsigned int) time(NULL));
-
for (int i = 0; i < NUM; i++) {
- data[i] = rand() / (RAND_MAX + 1.0);
+ data[i] = sc::rng::uniform();
fprintf(stderr, "%f\t", data[i]);
}
OclProcess(&OclFormulax, data, AVG);
commit c16b1dce36d530a7d4a21fc86596ada17e392c0b
Author: Peng Gao <peng at multicorewareinc.com>
Date: Sun Jun 16 17:18:23 2013 +0100
Initial OpenCL pieces.
Change-Id: I3a52cb7085b2dd8b70863a346eca279444206be6
diff --git a/RepositoryExternal.mk b/RepositoryExternal.mk
index 2a58722..479c08b 100644
--- a/RepositoryExternal.mk
+++ b/RepositoryExternal.mk
@@ -38,6 +38,23 @@
# External headers
+ifeq ($(ENABLE_OPENGCL),TRUE)
+
+define gb_LinkTarget__use_opencl
+$(call gb_LinkTarget_set_include,$(1),\
+ $$(INCLUDE) \
+ $(OPENCL_CFLAGS) \
+)
+$(call gb_LinkTarget_add_libs,$(1),$(OPENCL_LIBS))
+
+endef
+
+else # ENABLE_OPENCL != TRUE
+
+gb_LinkTarget__use_opencl :=
+
+endif
+
ifeq ($(SYSTEM_MARIADB),YES)
define gb_LinkTarget__use_mariadb
diff --git a/sc/Library_sc.mk b/sc/Library_sc.mk
index f4ed57d..7117349 100644
--- a/sc/Library_sc.mk
+++ b/sc/Library_sc.mk
@@ -49,6 +49,14 @@ $(eval $(call gb_Library_add_exception_objects,sc,\
))
endif
+ifeq ($(ENABLE_OPENCL),TRUE)
+$(eval $(call gb_Library_use_externals,sc,opencl))
+
+$(eval $(call gb_Library_add_exception_objects,sc,\
+ sc/source/core/opencl/openclwrapper \
+))
+endif
+
$(eval $(call gb_Library_use_libraries,sc,\
avmedia \
basegfx \
diff --git a/sc/source/core/opencl/oclkernels.hxx b/sc/source/core/opencl/oclkernels.hxx
new file mode 100644
index 0000000..f9db447
--- /dev/null
+++ b/sc/source/core/opencl/oclkernels.hxx
@@ -0,0 +1,116 @@
+/* -*- 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 _OCL_KERNEL_H_
+#define _OCL_KERNEL_H_
+
+#ifndef USE_EXTERNAL_KERNEL
+#define KERNEL( ... )# __VA_ARGS__
+
+/////////////////////////////////////////////
+const char *kernel_src = KERNEL(
+//Add kernel here
+//use \n ... \n to define macro
+__kernel void hello(__global uint *buffer)
+
+{
+size_t idx = get_global_id(0);
+
+buffer[idx]=idx;
+
+}
+
+__kernel void oclformula(__global float *data,
+ const uint type)
+{
+ const unsigned int i = get_global_id(0);
+
+ switch (type)
+ {
+ case 0: //MAX
+ {
+ //printf("%i ?%f>%f\n",i,data[2*i],data[2*i+1]);
+ if(data[2*i]>data[2*i+1])
+ data[i] = data[2*i];
+ else
+ data[i] = data[2*i+1];
+ break;
+ }
+ case 1: //MIN
+ {
+ //printf("%d ?%d<%d\n",i,data[2*i],data[2*i+1]);
+ if(data[2*i]<data[2*i+1])
+ data[i] = data[2*i];
+ else
+ data[i] = data[2*i+1];
+ break;
+ }
+ case 2: //SUM
+ case 3: //AVG
+ {
+ //printf("%d %d+%d\n",i,data[2*i],data[2*i+1]);
+ data[i] = data[2*i] + data[2*i+1];
+ break;
+ }
+ default:
+ break;
+
+ }
+
+}
+
+__kernel void oclFormulaMin(__global float *data,
+ const uint type)
+{
+
+}
+
+__kernel void oclFormulaMax(__global float *data,
+ const uint type)
+{
+
+}
+
+__kernel void oclFormulaSum(__global float *data,
+ const uint type)
+{
+
+}
+
+__kernel void oclFormulaCount(__global float *data,
+ const uint type)
+{
+
+}
+
+__kernel void oclFormulaAverage(__global float *data,
+ const uint type)
+{
+
+}
+
+
+__kernel void oclFormulaSumproduct(__global float *data,
+ const uint type)
+{
+
+}
+
+__kernel void oclFormulaMinverse(__global float *data,
+ const uint type)
+{
+
+}
+
+);
+
+#endif // USE_EXTERNAL_KERNEL
+#endif //_OCL_KERNEL_H_
+
+/* vim:set shiftwidth=4 softtabstop=4 expandtab: */
diff --git a/sc/source/core/opencl/openclwrapper.cxx b/sc/source/core/opencl/openclwrapper.cxx
new file mode 100644
index 0000000..05b169a
--- /dev/null
+++ b/sc/source/core/opencl/openclwrapper.cxx
@@ -0,0 +1,1032 @@
+/* -*- 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 <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+
+#include "openclwrapper.hxx"
+#include "oclkernels.hxx"
+
+
+inline int OpenclDevice::add_kernel_cfg(int kCount, const char *kName) {
+ strcpy(gpu_env.kernel_names[kCount], kName);
+ gpu_env.kernel_count++;
+ return 0;
+}
+
+int OpenclDevice::regist_opencl_kernel() {
+ if (!gpu_env.isUserCreated) {
+ memset(&gpu_env, 0, sizeof(gpu_env));
+ }
+
+ gpu_env.file_count = 0; //argc;
+ gpu_env.kernel_count = 0UL;
+
+ add_kernel_cfg(0, (const char*) "hello");
+ add_kernel_cfg(1, (const char*) "oclformula");
+ add_kernel_cfg(2, (const char*) "oclFormulaMin");
+ add_kernel_cfg(3, (const char*) "oclFormulaMax");
+ add_kernel_cfg(4, (const char*) "oclFormulaSum");
+ add_kernel_cfg(5, (const char*) "oclFormulaCount");
+ add_kernel_cfg(6, (const char*) "oclFormulaAverage");
+ add_kernel_cfg(7, (const char*) "oclFormulaSumproduct");
+ add_kernel_cfg(8, (const char*) "oclFormulaMinverse");
+ return 0;
+}
+OpenclDevice::OpenclDevice() :
+ isInited(0) {
+
+}
+
+OpenclDevice::~OpenclDevice() {
+
+}
+#ifdef USE_KERNEL_FILE
+int OpenclDevice::convert_to_string(const char *filename, char **source) {
+ int file_size;
+ size_t result;
+ FILE *file = NULL;
+
+ file_size = 0;
+ result = 0;
+ file = fopen(filename, "rb+");
+ printf("open kernel file %s.\n", filename);
+
+ if (file != NULL) {
+ fseek(file, 0, SEEK_END);
+
+ file_size = ftell(file);
+ rewind(file);
+ *source = (char*) malloc(sizeof(char) * file_size + 1);
+ if (*source == (char*) NULL) {
+ return (0);
+ }
+ result = fread(*source, 1, file_size, file);
+ if (result != (size_t) file_size) {
+ free(*source);
+ return (0);
+ }
+ (*source)[file_size] = '\0';
+ fclose(file);
+
+ return (1);
+ }
+ printf("open kernel file failed.\n");
+ return (0);
+}
+#endif
+int OpenclDevice::binary_generated(cl_context context,
+ const char * cl_file_name, FILE ** fhandle) {
+ unsigned int i = 0;
+ cl_int status;
+
+ size_t numDevices;
+
+ cl_device_id *devices;
+
+ char *str = NULL;
+
+ FILE *fd = NULL;
+
+ status = clGetContextInfo(context, CL_CONTEXT_NUM_DEVICES,
+ sizeof(numDevices), &numDevices, NULL);
+
+ CHECK_OPENCL(status)
+
+ devices = (cl_device_id*) malloc(sizeof(cl_device_id) * numDevices);
+
+ if (devices == NULL) {
+ return 0;
+ }
+
+ /* grab the handles to all of the devices in the context. */
+ status = clGetContextInfo(context, CL_CONTEXT_DEVICES,
+ sizeof(cl_device_id) * numDevices, devices, NULL);
+
+ status = 0;
+ /* dump out each binary into its own separate file. */
+ for (i = 0; i < numDevices; i++) {
+ char fileName[256] = { 0 }, cl_name[128] = { 0 };
+
+ if (devices[i] != 0) {
+ char deviceName[1024];
+ status = clGetDeviceInfo(devices[i], CL_DEVICE_NAME,
+ sizeof(deviceName), deviceName, NULL);
+ CHECK_OPENCL(status)
+ str = (char*) strstr(cl_file_name, (char*) ".cl");
+ memcpy(cl_name, cl_file_name, str - cl_file_name);
+ cl_name[str - cl_file_name] = '\0';
+ sprintf(fileName, "./%s-%s.bin", cl_name, deviceName);
+ fd = fopen(fileName, "rb");
+ status = (fd != NULL) ? 1 : 0;
+ }
+
+ }
+
+ if (devices != NULL) {
+ free(devices);
+ devices = NULL;
+ }
+
+ if (fd != NULL) {
+ *fhandle = fd;
+ }
+
+ return status;
+}
+
+int OpenclDevice::write_binary_to_file(const char* fileName, const char* birary,
+ size_t numBytes) {
+ FILE *output = NULL;
+ output = fopen(fileName, "wb");
+ if (output == NULL) {
+ return 0;
+ }
+
+ fwrite(birary, sizeof(char), numBytes, output);
+ fclose(output);
+
+ return 1;
+}
+
+int OpenclDevice::generat_bin_from_kernel_source(cl_program program,
+ const char * cl_file_name) {
+ unsigned int i = 0;
+ cl_int status;
+ size_t *binarySizes, numDevices;
+ cl_device_id *devices;
+ char **binaries, *str = NULL;
+
+ status = clGetProgramInfo(program, CL_PROGRAM_NUM_DEVICES,
+ sizeof(numDevices), &numDevices, NULL);
+ CHECK_OPENCL(status)
+
+ devices = (cl_device_id*) malloc(sizeof(cl_device_id) * numDevices);
+ if (devices == NULL) {
+ return 0;
+ }
+ /* grab the handles to all of the devices in the program. */
+ status = clGetProgramInfo(program, CL_PROGRAM_DEVICES,
+ sizeof(cl_device_id) * numDevices, devices, NULL);
+ CHECK_OPENCL(status)
+
+ /* figure out the sizes of each of the binaries. */
+ binarySizes = (size_t*) malloc(sizeof(size_t) * numDevices);
+
+ status = clGetProgramInfo(program, CL_PROGRAM_BINARY_SIZES,
+ sizeof(size_t) * numDevices, binarySizes, NULL);
+ CHECK_OPENCL(status)
+
+ /* copy over all of the generated binaries. */
+ binaries = (char**) malloc(sizeof(char *) * numDevices);
+ if (binaries == NULL) {
+ return 0;
+ }
+
+ for (i = 0; i < numDevices; i++) {
+ if (binarySizes[i] != 0) {
+ binaries[i] = (char*) malloc(sizeof(char) * binarySizes[i]);
+ if (binaries[i] == NULL) {
+ return 0;
+ }
+ } else {
+ binaries[i] = NULL;
+ }
+ }
+
+ status = clGetProgramInfo(program, CL_PROGRAM_BINARIES,
+ sizeof(char *) * numDevices, binaries, NULL);
+ CHECK_OPENCL(status)
+
+ /* dump out each binary into its own separate file. */
+ for (i = 0; i < numDevices; i++) {
+ char fileName[256] = { 0 }, cl_name[128] = { 0 };
+
+ if (binarySizes[i] != 0) {
+ char deviceName[1024];
+ status = clGetDeviceInfo(devices[i], CL_DEVICE_NAME,
+ sizeof(deviceName), deviceName, NULL);
+ CHECK_OPENCL(status)
+
+ str = (char*) strstr(cl_file_name, (char*) ".cl");
+ memcpy(cl_name, cl_file_name, str - cl_file_name);
+ cl_name[str - cl_file_name] = '\0';
+ sprintf(fileName, "./%s-%s.bin", cl_name, deviceName);
+
+ if (!write_binary_to_file(fileName, binaries[i], binarySizes[i])) {
+ printf("opencl-wrapper: write binary[%s] failds\n", fileName);
+ return 0;
+ } //else
+ printf("opencl-wrapper: write binary[%s] succesfully\n", fileName);
+ }
+ }
+
+ // Release all resouces and memory
+ for (i = 0; i < numDevices; i++) {
+ if (binaries[i] != NULL) {
+ free(binaries[i]);
+ binaries[i] = NULL;
+ }
+ }
+
+ if (binaries != NULL) {
+ free(binaries);
+ binaries = NULL;
+ }
+
+ if (binarySizes != NULL) {
+ free(binarySizes);
+ binarySizes = NULL;
+ }
+
+ if (devices != NULL) {
+ free(devices);
+ devices = NULL;
+ }
+ return 1;
+}
+
+int OpenclDevice::init_opencl_attr(OpenCLEnv * env) {
+ if (gpu_env.isUserCreated) {
+ return 1;
+ }
+
+ gpu_env.context = env->context;
+ gpu_env.platform = env->platform;
+ gpu_env.dev = env->devices;
+ gpu_env.command_queue = env->command_queue;
+
+ gpu_env.isUserCreated = 1;
+
+ return 0;
+}
+
+int OpenclDevice::create_kernel(char * kernelname, KernelEnv * env) {
+ int status;
+
+ env->kernel = clCreateKernel(gpu_env.programs[0], kernelname, &status);
+ env->context = gpu_env.context;
+ env->command_queue = gpu_env.command_queue;
+ return status != CL_SUCCESS ? 1 : 0;
+}
+
+int OpenclDevice::release_kernel(KernelEnv * env) {
+ int status = clReleaseKernel(env->kernel);
+ return status != CL_SUCCESS ? 1 : 0;
+}
+
+int OpenclDevice::init_opencl_env(GPUEnv *gpu_info) {
+ size_t length;
+ cl_int status;
+ cl_uint numPlatforms, numDevices;
+ cl_platform_id *platforms;
+ cl_context_properties cps[3];
+ char platformName[100];
+ unsigned int i;
+
+ /*
+ * Have a look at the available platforms.
+ */
+ if (!gpu_info->isUserCreated) {
+ status = clGetPlatformIDs(0, NULL, &numPlatforms);
+ if (status != CL_SUCCESS) {
+ return (1);
+ }
+ gpu_info->platform = NULL;
+ ;
+ if (0 < numPlatforms) {
+ platforms = (cl_platform_id*) malloc(
+ numPlatforms * sizeof(cl_platform_id));
+ if (platforms == (cl_platform_id*) NULL) {
+ return (1);
+ }
+ status = clGetPlatformIDs(numPlatforms, platforms, NULL);
+
+ if (status != CL_SUCCESS) {
+ return (1);
+ }
+
+ for (i = 0; i < numPlatforms; i++) {
+ status = clGetPlatformInfo(platforms[i], CL_PLATFORM_VENDOR,
+ sizeof(platformName), platformName, NULL);
+
+ if (status != CL_SUCCESS) {
+ return (1);
+ }
+ gpu_info->platform = platforms[i];
+
+ //if (!strcmp(platformName, "Intel(R) Coporation"))
+ //if( !strcmp( platformName, "Advanced Micro Devices, Inc." ))
+ {
+ gpu_info->platform = platforms[i];
+
+ status = clGetDeviceIDs(gpu_info->platform /* platform */,
+ CL_DEVICE_TYPE_GPU /* device_type */,
+ 0 /* num_entries */, NULL /* devices */,
+ &numDevices);
+
+ if (status != CL_SUCCESS) {
+ return (1);
+ }
+
+ if (numDevices) {
+ break;
+ }
+ }
+ }
+ free(platforms);
+ }
+ if (NULL == gpu_info->platform) {
+ return (1);
+ }
+
+ /*
+ * Use available platform.
+ */
+ cps[0] = CL_CONTEXT_PLATFORM;
+ cps[1] = (cl_context_properties) gpu_info->platform;
+ cps[2] = 0;
+ /* Check for GPU. */
+ gpu_info->dType = CL_DEVICE_TYPE_GPU;
+ gpu_info->context = clCreateContextFromType(cps, gpu_info->dType, NULL,
+ NULL, &status);
+
+ if ((gpu_info->context == (cl_context) NULL)
+ || (status != CL_SUCCESS)) {
+ gpu_info->dType = CL_DEVICE_TYPE_CPU;
+ gpu_info->context = clCreateContextFromType(cps, gpu_info->dType,
+ NULL, NULL, &status);
+ }
+ if ((gpu_info->context == (cl_context) NULL)
+ || (status != CL_SUCCESS)) {
+ gpu_info->dType = CL_DEVICE_TYPE_DEFAULT;
+ gpu_info->context = clCreateContextFromType(cps, gpu_info->dType,
+ NULL, NULL, &status);
+ }
+ if ((gpu_info->context == (cl_context) NULL)
+ || (status != CL_SUCCESS)) {
+ return (1);
+ }
+ /* Detect OpenCL devices. */
+ /* First, get the size of device list data */
+ status = clGetContextInfo(gpu_info->context, CL_CONTEXT_DEVICES, 0,
+ NULL, &length);
+
+ if ((status != CL_SUCCESS) || (length == 0)) {
+ return (1);
+ }
+ /* Now allocate memory for device list based on the size we got earlier */
+ gpu_info->devices = (cl_device_id*) malloc(length);
+ if (gpu_info->devices == (cl_device_id*) NULL) {
+ return (1);
+ }
+ /* Now, get the device list data */
+ status = clGetContextInfo(gpu_info->context, CL_CONTEXT_DEVICES, length,
+ gpu_info->devices, NULL);
+
+ if (status != CL_SUCCESS) {
+ return (1);
+ }
+
+ /* Create OpenCL command queue. */
+ gpu_info->command_queue = clCreateCommandQueue(gpu_info->context,
+ gpu_info->devices[0], 0, &status);
+
+ if (status != CL_SUCCESS) {
+ return (1);
+ }
+ }
+
+ status = clGetCommandQueueInfo(gpu_info->command_queue,
+ CL_QUEUE_THREAD_HANDLE_AMD, 0, NULL, NULL);
+
+ return 0;
+}
+
+int OpenclDevice::release_opencl_env(GPUEnv *gpu_info) {
+ int i = 0;
+ int status = 0;
+
+ if (!isInited) {
+ return 1;
+ }
+
+ for (i = 0; i < gpu_env.file_count; i++) {
+ if (gpu_env.programs[i]) {
+ status = clReleaseProgram(gpu_env.programs[i]);
+ CHECK_OPENCL(status)
+ gpu_env.programs[i] = NULL;
+ }
+ }
+ if (gpu_env.command_queue) {
+ clReleaseCommandQueue(gpu_env.command_queue);
+ gpu_env.command_queue = NULL;
+ }
+ if (gpu_env.context) {
+ clReleaseContext(gpu_env.context);
+ gpu_env.context = NULL;
+ }
+ isInited = 0;
+ gpu_info->isUserCreated = 0;
+ free(gpu_info->devices);
+ return 1;
+}
+
+int OpenclDevice::run_kernel_wrapper(cl_kernel_function function,
+ char * kernel_name, void **usrdata) {
+ printf("oclwrapper:run_kernel_wrapper...\n");
+ if (register_kernel_wrapper(kernel_name, function) != 1) {
+ fprintf(stderr,
+ "Error:run_kernel_wrapper:register_kernel_wrapper fail!\n");
+ return -1;
+ }
+ return (run_kernel(kernel_name, usrdata));
+}
+
+int OpenclDevice::register_kernel_wrapper(const char *kernel_name,
+ cl_kernel_function function) {
+ int i;
+ printf("oclwrapper:register_kernel_wrapper...%d\n", gpu_env.kernel_count);
+ for (i = 0; i < gpu_env.kernel_count; i++) {
+ //printf("oclwrapper:register_kernel_wrapper kname...%s\n", kernel_name);
+ //printf("oclwrapper:register_kernel_wrapper kname...%s\n", gpu_env.kernel_names[i]);
+ if (strcasecmp(kernel_name, gpu_env.kernel_names[i]) == 0) {
+ //printf("oclwrapper:register_kernel_wrapper if()...\n");
+ gpu_env.kernel_functions[i] = function;
+ return (1);
+ }
+ }
+ return (0);
+}
+
+int OpenclDevice::cached_of_kerner_prg(const GPUEnv *gpu_env_cached,
+ const char * cl_file_name) {
+ int i;
+ for (i = 0; i < gpu_env_cached->file_count; i++) {
+ if (strcasecmp(gpu_env_cached->kernelSrcFile[i], cl_file_name) == 0) {
+ if (gpu_env_cached->programs[i] != NULL) {
+ return (1);
+ }
+ }
+ }
+
+ return (0);
+}
+
+int OpenclDevice::compile_kernel_file(GPUEnv *gpu_info, const char *build_option) {
+ cl_int status;
+
+ size_t length;
+
+ char *buildLog = NULL, *binary;
+
+ const char *source;
+ size_t source_size[1];
+
+ int b_error, binary_status, binaryExisted, idx;
+
+ size_t numDevices;
+
+ cl_device_id *devices;
+
+ FILE *fd, *fd1;
+ const char* filename = "kernel.cl";
+ if (cached_of_kerner_prg(gpu_info, filename) == 1) {
+ return (1);
+ }
+
+ idx = gpu_info->file_count;
+
+ source = kernel_src;
+
+ source_size[0] = strlen(source);
+
+ binaryExisted = 0;
+ if ((binaryExisted = binary_generated(gpu_info->context, filename, &fd)) == 1) {
+ status = clGetContextInfo(gpu_info->context, CL_CONTEXT_NUM_DEVICES,
+ sizeof(numDevices), &numDevices, NULL);
+ CHECK_OPENCL(status)
+
+ devices = (cl_device_id*) malloc(sizeof(cl_device_id) * numDevices);
+ if (devices == NULL) {
+ return 0;
+ }
+
+ b_error = 0;
+ length = 0;
+ b_error |= fseek(fd, 0, SEEK_END) < 0;
+ b_error |= (length = ftell(fd)) <= 0;
+ b_error |= fseek(fd, 0, SEEK_SET) < 0;
+ if (b_error) {
+ return 0;
+ }
+
+ binary = (char*) malloc(length + 2);
+ if (!binary) {
+ return 0;
+ }
+
+ memset(binary, 0, length + 2);
+ b_error |= fread(binary, 1, length, fd) != length;
+ if (binary[length - 1] != '\n') {
+ binary[length++] = '\n';
+ }
+
+ fclose(fd);
+ fd = NULL;
+ // grab the handles to all of the devices in the context.
+ status = clGetContextInfo(gpu_info->context, CL_CONTEXT_DEVICES,
+ sizeof(cl_device_id) * numDevices, devices, NULL);
+ CHECK_OPENCL(status)
+
+ gpu_info->programs[idx] = clCreateProgramWithBinary(gpu_info->context,
+ numDevices, devices, &length, (const unsigned char**) &binary,
+ &binary_status, &status);
+ CHECK_OPENCL(status)
+
+ free(binary);
+ free(devices);
+ devices = NULL;
+ } else {
+
+ // create a CL program using the kernel source
+ gpu_info->programs[idx] = clCreateProgramWithSource(gpu_info->context,
+ 1, &source, source_size, &status);
+ CHECK_OPENCL(status)
+
+ printf("clCreateProgramWithSource.\n");
+ }
+
+ if (gpu_info->programs[idx] == (cl_program) NULL) {
+ return (0);
+ }
+
+ //char options[512];
+ // create a cl program executable for all the devices specified
+ if (!gpu_info->isUserCreated) {
+ status = clBuildProgram(gpu_info->programs[idx], 1, gpu_info->devices,
+ build_option, NULL, NULL);
+ CHECK_OPENCL(status)
+ } else {
+ status = clBuildProgram(gpu_info->programs[idx], 1, &(gpu_info->dev),
+ build_option, NULL, NULL);
+ CHECK_OPENCL(status)
+ }
+ printf("BuildProgram.\n");
+
+ if (status != CL_SUCCESS) {
+ if (!gpu_info->isUserCreated) {
+ status = clGetProgramBuildInfo(gpu_info->programs[idx],
+ gpu_info->devices[0], CL_PROGRAM_BUILD_LOG, 0, NULL,
+ &length);
+ } else {
+ status = clGetProgramBuildInfo(gpu_info->programs[idx],
+ gpu_info->dev, CL_PROGRAM_BUILD_LOG, 0, NULL, &length);
+ }
+ if (status != CL_SUCCESS) {
+ printf("opencl create build log fail\n");
+ return (0);
+ }
+ buildLog = (char*) malloc(length);
+ if (buildLog == (char*) NULL) {
+ return (0);
+ }
+ if (!gpu_info->isUserCreated) {
+ status = clGetProgramBuildInfo(gpu_info->programs[idx],
+ gpu_info->devices[0], CL_PROGRAM_BUILD_LOG, length,
+ buildLog, &length);
+ } else {
+ status = clGetProgramBuildInfo(gpu_info->programs[idx],
+ gpu_info->dev, CL_PROGRAM_BUILD_LOG, length, buildLog,
+ &length);
+ }
+
+ fd1 = fopen("kernel-build.log", "w+");
+ if (fd1 != NULL) {
+ fwrite(buildLog, sizeof(char), length, fd1);
+ fclose(fd1);
+ }
+
+ free(buildLog);
+ return (0);
+ }
+
+ strcpy(gpu_env.kernelSrcFile[idx], filename);
+
+ if (binaryExisted == 0)
+ generat_bin_from_kernel_source(gpu_env.programs[idx], filename);
+
+ gpu_info->file_count += 1;
+
+ return (1);
+}
+
+int OpenclDevice::compile_kernel_file(const char *filename, GPUEnv *gpu_info,
+ const char *build_option) {
+ cl_int status;
+
+ size_t length;
+
+#ifdef USE_KERNEL_FILE
+ char
+ *source_str;
+#endif
+ char *buildLog = NULL, *binary;
+
+ const char *source;
+ size_t source_size[1];
+
+ int b_error, binary_status, binaryExisted, idx;
+
+ size_t numDevices;
+
+ cl_device_id *devices;
+
+ FILE *fd, *fd1;
+
+ if (cached_of_kerner_prg(gpu_info, filename) == 1) {
+ return (1);
+ }
+
+ idx = gpu_info->file_count;
+#ifdef USE_KERNEL_FILE
+ status = convert_to_string( filename, &source_str, gpu_info, idx );
+
+ if( status == 0 )
+ {
+ printf("convert_to_string failed.\n");
+ return(0);
+ }
+ source = source_str;
+#else
+
+ source = kernel_src;
+#endif
+ source_size[0] = strlen(source);
+
+ binaryExisted = 0;
+ if ((binaryExisted = binary_generated(gpu_info->context, filename, &fd))
+ == 1) {
+ status = clGetContextInfo(gpu_info->context, CL_CONTEXT_NUM_DEVICES,
+ sizeof(numDevices), &numDevices, NULL);
+ CHECK_OPENCL(status)
+
+ devices = (cl_device_id*) malloc(sizeof(cl_device_id) * numDevices);
+ if (devices == NULL) {
+ return 0;
+ }
+
+ b_error = 0;
+ length = 0;
+ b_error |= fseek(fd, 0, SEEK_END) < 0;
+ b_error |= (length = ftell(fd)) <= 0;
+ b_error |= fseek(fd, 0, SEEK_SET) < 0;
+ if (b_error) {
+ return 0;
+ }
+
+ binary = (char*) malloc(length + 2);
+ if (!binary) {
+ return 0;
+ }
+
+ memset(binary, 0, length + 2);
+ b_error |= fread(binary, 1, length, fd) != length;
+ if (binary[length - 1] != '\n') {
+ binary[length++] = '\n';
+ }
+
+ fclose(fd);
+ fd = NULL;
+ /* grab the handles to all of the devices in the context. */
+ status = clGetContextInfo(gpu_info->context, CL_CONTEXT_DEVICES,
+ sizeof(cl_device_id) * numDevices, devices, NULL);
+ CHECK_OPENCL(status)
+
+ gpu_info->programs[idx] = clCreateProgramWithBinary(gpu_info->context,
+ numDevices, devices, &length, (const unsigned char**) &binary,
+ &binary_status, &status);
+ CHECK_OPENCL(status)
+
+ free(binary);
+ free(devices);
+ devices = NULL;
+ } else {
+
+ // create a CL program using the kernel source
+ gpu_info->programs[idx] = clCreateProgramWithSource(gpu_info->context,
+ 1, &source, source_size, &status);
+ CHECK_OPENCL(status)
+#ifdef USE_KERNEL_FILE
+ free((char*)source);
+#endif
+ printf("clCreateProgramWithSource.\n");
+ }
+
+ if (gpu_info->programs[idx] == (cl_program) NULL) {
+ return (0);
+ }
+
+ //char options[512];
+ // create a cl program executable for all the devices specified
+ if (!gpu_info->isUserCreated) {
+ status = clBuildProgram(gpu_info->programs[idx], 1, gpu_info->devices,
+ build_option, NULL, NULL);
+ CHECK_OPENCL(status)
+ } else {
+ status = clBuildProgram(gpu_info->programs[idx], 1, &(gpu_info->dev),
+ build_option, NULL, NULL);
+ CHECK_OPENCL(status)
+ }
+ printf("BuildProgram.\n");
+
+ if (status != CL_SUCCESS) {
+ if (!gpu_info->isUserCreated) {
+ status = clGetProgramBuildInfo(gpu_info->programs[idx],
+ gpu_info->devices[0], CL_PROGRAM_BUILD_LOG, 0, NULL,
+ &length);
+ } else {
+ status = clGetProgramBuildInfo(gpu_info->programs[idx],
+ gpu_info->dev, CL_PROGRAM_BUILD_LOG, 0, NULL, &length);
+ }
+ if (status != CL_SUCCESS) {
+ printf("opencl create build log fail\n");
+ return (0);
+ }
+ buildLog = (char*) malloc(length);
+ if (buildLog == (char*) NULL) {
+ return (0);
+ }
+ if (!gpu_info->isUserCreated) {
+ status = clGetProgramBuildInfo(gpu_info->programs[idx],
+ gpu_info->devices[0], CL_PROGRAM_BUILD_LOG, length,
+ buildLog, &length);
+ } else {
+ status = clGetProgramBuildInfo(gpu_info->programs[idx],
+ gpu_info->dev, CL_PROGRAM_BUILD_LOG, length, buildLog,
+ &length);
+ }
+
+ fd1 = fopen("kernel-build.log", "w+");
+ if (fd1 != NULL) {
+ fwrite(buildLog, sizeof(char), length, fd1);
+ fclose(fd1);
+ }
+
+ free(buildLog);
+ return (0);
+ }
+
+ strcpy(gpu_env.kernelSrcFile[idx], filename);
+
+ if (binaryExisted == 0)
+ generat_bin_from_kernel_source(gpu_env.programs[idx], filename);
+
+ gpu_info->file_count += 1;
+
+ return (1);
+}
+
+int OpenclDevice::get_kernel_env_and_func(const char *kernel_name,
+ KernelEnv *env, cl_kernel_function *function) {
+ int i; //,program_idx ;
+ for (i = 0; i < gpu_env.kernel_count; i++) {
+ if (strcasecmp(kernel_name, gpu_env.kernel_names[i]) == 0) {
+ env->context = gpu_env.context;
+ env->command_queue = gpu_env.command_queue;
+ env->program = gpu_env.programs[0];
+ env->kernel = gpu_env.kernels[i];
+ *function = gpu_env.kernel_functions[i];
+ return (1);
+ }
+ }
+ return (0);
+}
+
+int OpenclDevice::run_kernel(const char *kernel_name, void **userdata) {
+ KernelEnv env;
+
+ cl_kernel_function function;
+
+ int status;
+
+ memset(&env, 0, sizeof(KernelEnv));
+ status = get_kernel_env_and_func(kernel_name, &env, &function);
+ strcpy(env.kernel_name, kernel_name);
+ if (status == 1) {
+ if (&env == (KernelEnv *) NULL
+ || &function == (cl_kernel_function *) NULL) {
+ return (0);
+ }
+ return (function(userdata, &env));
+ }
+ return (0);
+}
+
+int OpenclDevice::init_opencl_run_env(int argc, const char *build_option_kernelfiles)
+{
+ int status = 0;
+
+ if (MAX_CLKERNEL_NUM <= 0) {
+ return 1;
+ }
+ if ((argc > MAX_CLFILE_NUM) || (argc < 0)) {
+ return 1;
+ }
+
+ if (!isInited) {
+ printf("regist_opencl_kernel start.\n");
+ regist_opencl_kernel();
+ //initialize devices, context, comand_queue
+ status = init_opencl_env(&gpu_env);
+ if (status) {
+ printf("init_opencl_env failed.\n");
+ return (1);
+ }
+ printf("init_opencl_env successed.\n");
+ //initialize program, kernel_name, kernel_count
+ status = compile_kernel_file( &gpu_env, build_option_kernelfiles);
+ if (status == 0 || gpu_env.kernel_count == 0) {
+ printf("compile_kernel_file failed.\n");
+ return (1);
+ }
+ printf("compile_kernel_file successed.\n");
+ isInited = 1;
+ }
+
+ return (0);
+}
+
+int OpenclDevice::init_opencl_run_env(int argc, const char *argv_kernelfiles[],
+ const char *build_option_kernelfiles) {
+ int status = 0;
+
+ if (MAX_CLKERNEL_NUM <= 0) {
+ return 1;
+ }
+ if ((argc > MAX_CLFILE_NUM) || (argc < 0)) {
+ return 1;
+ }
+
+ if (!isInited) {
+ printf("regist_opencl_kernel start.\n");
+ regist_opencl_kernel();
+ //initialize devices, context, comand_queue
+ status = init_opencl_env(&gpu_env);
+ if (status) {
+ printf("init_opencl_env failed.\n");
+ return (1);
+ }
+ printf("init_opencl_env successed.\n");
+ //initialize program, kernel_name, kernel_count
+ status = compile_kernel_file(argv_kernelfiles[0], &gpu_env,
+ build_option_kernelfiles);
+ if (status == 0 || gpu_env.kernel_count == 0) {
+ printf("compile_kernel_file failed.\n");
+ return (1);
+ }
+ printf("compile_kernel_file successed.\n");
+ isInited = 1;
+ }
+
+ return (0);
+}
+
+int OpenclDevice::release_opencl_run_env() {
+ return release_opencl_env(&gpu_env);
+}
+
+void OpenclDevice::setOpenclState(int state) {
+ isInited = state;
+}
+
+int OpenclDevice::getOpenclState() {
+ return isInited;
+}
+//ocldbg
+int OclFormulaMin(void ** usrdata, KernelEnv *env) { return 0; }
+int OclFormulaMax(void ** usrdata, KernelEnv *env) { return 0; }
+int OclFormulaSum(void ** usrdata, KernelEnv *env) { return 0; }
+int OclFormulaCount(void ** usrdata, KernelEnv *env) { return 0; }
+int OclFormulaAverage(void ** usrdata, KernelEnv *env) { return 0; }
+int OclFormulaSumproduct(void ** usrdata, KernelEnv *env) { return 0; }
+int OclFormulaMinverse(void ** usrdata, KernelEnv *env) { return 0; }
+
+int OclFormulax(void ** usrdata, KernelEnv *env) {
+ fprintf(stderr, "In OpenclDevice,...Formula_proc\n");
+ cl_int clStatus;
+ int status;
+ size_t global_work_size[1];
+ float tdata[NUM];
+
+ double *data = (double *) usrdata[0];
+ const formulax type = *((const formulax *) usrdata[1]);
+ double ret = 0.0;
+
+ for (int i = 0; i < NUM; i++) {
+ tdata[i] = (float) data[i];
+ }
+
+ env->kernel = clCreateKernel(env->program, "oclformula", &clStatus);
+ //printf("ScInterpreter::IterateParameters...after clCreateKernel.\n");
+ //fprintf(stderr, "\nIn OpenclDevice,...after clCreateKernel\n");
+ int size = NUM;
+
+ cl_mem formula_data = clCreateBuffer(env->context,
+ (cl_mem_flags) (CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR),
+ size * sizeof(float), (void *) tdata, &clStatus);
+ //fprintf(stderr, "\nIn OpenclDevice,...after clCreateBuffer\n");
+
+ status = clSetKernelArg(env->kernel, 0, sizeof(cl_mem),
+ (void *) &formula_data);
+ CHECK_OPENCL(status)
+ status = clSetKernelArg(env->kernel, 1, sizeof(unsigned int),
+ (void *) &type);
+ CHECK_OPENCL(status)
+
+ global_work_size[0] = size;
+ //fprintf(stderr, "\nIn OpenclDevice,...after global_work_size\n");
+ //PPAStartCpuEvent(ppa_proc);
+
+ while (global_work_size[0] != 1) {
+ global_work_size[0] = global_work_size[0] / 2;
+ status = clEnqueueNDRangeKernel(env->command_queue, env->kernel, 1,
+ NULL, global_work_size, NULL, 0, NULL, NULL);
+ CHECK_OPENCL(status)
+
+ }
+ //fprintf(stderr, "\nIn OpenclDevice,...before clEnqueueReadBuffer\n");
+ status = clEnqueueReadBuffer(env->command_queue, formula_data, CL_FALSE, 0,
+ sizeof(float), (void *) &tdata, 0, NULL, NULL);
+ CHECK_OPENCL(status)
+ status = clFinish(env->command_queue);
+ CHECK_OPENCL(status)
+
+ //PPAStopCpuEvent(ppa_proc);
+ //fprintf(stderr, "\nIn OpenclDevice,...before clReleaseKernel\n");
+ status = clReleaseKernel(env->kernel);
+ CHECK_OPENCL(status)
+ status = clReleaseMemObject(formula_data);
+ CHECK_OPENCL(status)
+
+ if (type == AVG)
+ ret = (double) tdata[0] / NUM;
+ else
+ ret = (double) tdata[0];
+
+ printf("size = %d ret = %f.\n\n", NUM, ret);
+
+ return 0;
+}
+double OclCalc::OclProcess(cl_kernel_function function, double *data,
+ formulax type) {
+ fprintf(stderr, "\In OpenclDevice, proc...begin\n");
+ double ret = 0;
+
+ void *usrdata[2];
+
+ usrdata[0] = (void *) data;
+ usrdata[1] = (void *) &type;
+
+ run_kernel_wrapper(function, "oclformula", usrdata);
+ //fprintf(stderr, "\In OpenclDevice, proc...after run_kernel_wrapper\n");
+ return ret;
+}
+
+double OclCalc::OclTest() {
+ double data[NUM];
+
+ srand((unsigned int) time(NULL));
+
+ for (int i = 0; i < NUM; i++) {
+ data[i] = rand() / (RAND_MAX + 1.0);
+ fprintf(stderr, "%f\t", data[i]);
+ }
+ OclProcess(&OclFormulax, data, AVG);
+ //fprintf(stderr, "\nIn OpenclDevice,OclTest() after proc,data0...%f\n", data[0]);
+
+ return 0.0;
+}
+
+OclCalc::OclCalc()
+{
+ OpenclDevice::init_opencl_run_env(0, NULL);
+ OpenclDevice::setOpenclState(1);
+ fprintf(stderr,"OclCalc:: init opencl.\n");
+}
+
+OclCalc::~OclCalc()
+{
+ OpenclDevice::release_opencl_run_env();
+ OpenclDevice::setOpenclState(0);
+ fprintf(stderr,"OclCalc:: opencl end.\n");
+}
+
+/* vim:set shiftwidth=4 softtabstop=4 expandtab: */
diff --git a/sc/source/core/opencl/openclwrapper.hxx b/sc/source/core/opencl/openclwrapper.hxx
new file mode 100644
index 0000000..4646954
--- /dev/null
+++ b/sc/source/core/opencl/openclwrapper.hxx
@@ -0,0 +1,172 @@
+/* -*- 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/.
+ *
+ * This file incorporates work covered by the following license notice:
+ *
+ * Licensed to the Apache Software Foundation (ASF) under one or more
+ * contributor license agreements. See the NOTICE file distributed
+ * with this work for additional information regarding copyright
+ * ownership. The ASF licenses this file to you under the Apache
+ * License, Version 2.0 (the "License"); you may not use this file
+ * except in compliance with the License. You may obtain a copy of
+ * the License at http://www.apache.org/licenses/LICENSE-2.0 .
+ */
+
+#ifndef _OPENCL_WRAPPER_H_
+#define _OPENCL_WRAPPER_H_
+
+#include <CL/cl.h>
+
+#define MaxTextExtent 4096
+//support AMD opencl
+#define CL_QUEUE_THREAD_HANDLE_AMD 0x403E
+#define CL_MAP_WRITE_INVALIDATE_REGION (1 << 2)
+
+#if defined(_MSC_VER)
+#ifndef strcasecmp
+#define strcasecmp strcmp
+#endif
+#endif
+
+typedef struct _KernelEnv {
+ cl_context context;
+ cl_command_queue command_queue;
+ cl_program program;
+ cl_kernel kernel;
+ char kernel_name[150];
+} KernelEnv;
+
+typedef struct _OpenCLEnv {
+ cl_platform_id platform;
+ cl_context context;
+ cl_device_id devices;
+ cl_command_queue command_queue;
+} OpenCLEnv;
+
+#if defined __cplusplus
+extern "C" {
+#endif
+
+//user defined, this is function wrapper which is used to set the input parameters ,
+//luanch kernel and copy data from GPU to CPU or CPU to GPU.
+typedef int (*cl_kernel_function)(void **userdata, KernelEnv *kenv);
+
+#if defined __cplusplus
+
+}
+#endif
+
+#define CHECK_OPENCL(status) \
+if(status != CL_SUCCESS) \
+{ \
+ printf ("error code is %d.",status); \
+ return (0); \
+}
+#endif
+
+#define MAX_KERNEL_STRING_LEN 64
+#define MAX_CLFILE_NUM 50
+#define MAX_CLKERNEL_NUM 200
+#define MAX_KERNEL_NAME_LEN 64
+
+typedef struct _GPUEnv {
+ //share vb in all modules in hb library
+ cl_platform_id platform;
+
+ cl_device_type dType;
+
+ cl_context context;
+
+ cl_device_id *devices;
+
+ cl_device_id dev;
+
+ cl_command_queue command_queue;
+
+ cl_kernel kernels[MAX_CLFILE_NUM];
+
+ cl_program programs[MAX_CLFILE_NUM]; //one program object maps one kernel source file
+
+ char kernelSrcFile[MAX_CLFILE_NUM][256], //the max len of kernel file name is 256
+ kernel_names[MAX_CLKERNEL_NUM][MAX_KERNEL_STRING_LEN + 1];
+
+ cl_kernel_function kernel_functions[MAX_CLKERNEL_NUM];
+
+ int kernel_count, file_count, // only one kernel file
+ isUserCreated; // 1: created , 0:no create and needed to create by opencl wrapper
+
+} GPUEnv;
+
+typedef struct {
+ char kernelName[MAX_KERNEL_NAME_LEN + 1];
+ char *kernelStr;
+} kernel_node;
+
+class OpenclDevice {
+private:
+ GPUEnv gpu_env;
+ int isInited;
+
+public:
+ OpenclDevice();
+ ~OpenclDevice();
+ int regist_opencl_kernel();
+ int convert_to_string(const char *filename, char **source);
+ int binary_generated(cl_context context, const char * cl_file_name,
+ FILE ** fhandle);
+ int write_binary_to_file(const char* fileName, const char* birary,
+ size_t numBytes);
+ int generat_bin_from_kernel_source(cl_program program,
+ const char * cl_file_name);
+ int init_opencl_attr(OpenCLEnv * env);
+ int create_kernel(char * kernelname, KernelEnv * env);
+ int release_kernel(KernelEnv * env);
+ int init_opencl_env(GPUEnv *gpu_info);
+ int release_opencl_env(GPUEnv *gpu_info);
+ int run_kernel_wrapper(cl_kernel_function function, char * kernel_name,
+ void **usrdata);
+ int register_kernel_wrapper(const char *kernel_name,
+ cl_kernel_function function);
+ int cached_of_kerner_prg(const GPUEnv *gpu_env_cached,
+ const char * cl_file_name);
+ int compile_kernel_file(GPUEnv *gpu_info, const char *build_option);
+ int compile_kernel_file(const char *filename, GPUEnv *gpu_info,
+ const char *build_option);
+ int get_kernel_env_and_func(const char *kernel_name, KernelEnv *env,
+ cl_kernel_function *function);
+ int run_kernel(const char *kernel_name, void **userdata);
+ int init_opencl_run_env(int argc, const char *build_option_kernelfiles);
+ int init_opencl_run_env(int argc, const char *argv_kernelfiles[],
+ const char *build_option_kernelfiles);
+ int release_opencl_run_env();
+ void setOpenclState(int state);
+ int getOpenclState();
+ inline int add_kernel_cfg(int kCount, const char *kName);
+
+};
+
+#define NUM 4//(16*16*16)
+typedef enum _formulax_ {
+ MIN, MAX, SUM, AVG, COUNT, SUMPRODUCT, MINVERSE
+} formulax;
+class OclCalc: public OpenclDevice {
+public:
+ OclCalc();
+ ~OclCalc();
+ double OclProcess(cl_kernel_function function, double *data, formulax type);
+ double OclTest();
+ double OclMin();
+ double OclMax();
+ double OclSum();
+ double OclCount();
+ double OclAverage();
+ double OclSumproduct();
+ double OclMinverse();
+
+};
+
diff --git a/sc/source/core/tool/interpr1.cxx b/sc/source/core/tool/interpr1.cxx
index ff9f7cb..bf56fde 100644
--- a/sc/source/core/tool/interpr1.cxx
+++ b/sc/source/core/tool/interpr1.cxx
@@ -4328,6 +4328,10 @@ void ScInterpreter::ScProduct()
void ScInterpreter::ScAverage( bool bTextAsZero )
{
RTL_LOGFILE_CONTEXT_AUTHOR( aLogger, "sc", "er", "ScInterpreter::ScAverage" );
+#ifdef ENABLE_OPENCL
+ static OclCalc ocl_calc;
+ ocl_calc.OclTest();
+#endif
PushDouble( IterateParameters( ifAVERAGE, bTextAsZero ) );
}
commit 45f2ecfc93d9074453838438e57ff48996c95d8b
Author: Michael Meeks <michael.meeks at suse.com>
Date: Thu Jun 13 21:36:53 2013 +0100
initial opencl configure pieces.
Change-Id: I5592fdab8f6029e371b3d70f3334b737a6262eac
diff --git a/config_host.mk.in b/config_host.mk.in
index 6e5e0cd..c8b3a25 100644
--- a/config_host.mk.in
+++ b/config_host.mk.in
@@ -147,6 +147,7 @@ export ENABLE_NPAPI_FROM_BROWSER=@ENABLE_NPAPI_FROM_BROWSER@
export ENABLE_NPAPI_INTO_BROWSER=@ENABLE_NPAPI_INTO_BROWSER@
export ENABLE_ONLINE_UPDATE=@ENABLE_ONLINE_UPDATE@
export ENABLE_OPENGL=@ENABLE_OPENGL@
+export ENABLE_OPENCL=@ENABLE_OPENCL@
export ENABLE_PACKAGEKIT=@ENABLE_PACKAGEKIT@
export ENABLE_PCH=@ENABLE_PCH@
export ENABLE_PDFIMPORT=@ENABLE_PDFIMPORT@
@@ -393,6 +394,8 @@ export OOOP_SAMPLES_PACK=@OOOP_SAMPLES_PACK@
export OOOP_TEMPLATES_PACK=@OOOP_TEMPLATES_PACK@
export OOO_JUNIT_JAR=@OOO_JUNIT_JAR@
export OOO_VENDOR=@OOO_VENDOR@
+export OPENCL_CFLAGS=$(gb_SPACE)@OPENCL_CFLAGS@
+export OPENCL_LIBS=$(gb_SPACE)@OPENCL_LIBS@
export OPENSSL_CFLAGS=$(gb_SPACE)@OPENSSL_CFLAGS@
export OPENSSL_LIBS=$(gb_SPACE)@OPENSSL_LIBS@
export ORCUS_CFLAGS=$(gb_SPACE)@ORCUS_CFLAGS@
diff --git a/configure.ac b/configure.ac
index 53c6d9d..f81ab47 100644
--- a/configure.ac
+++ b/configure.ac
@@ -1752,6 +1752,16 @@ AC_ARG_WITH(windows-sdk,
],
,)
+AC_ARG_WITH(opencl-sdk,
+ AS_HELP_STRING([--with-opencl-sdk],
+ [If you have the AMD HSA / opencl SDK installed (cf.
+ http://developer.amd.com/tools-and-sdks/heterogeneous-computing/amd-accelerated-parallel-processing-app-sdk/downloads/ )
+ Pass the absolute path to where that SDK is unpacked.])
+ [
+ Usage: --with-opencl-sdk=/opt/AMDAPP
+ ],
+,)
+
AC_ARG_WITH(lang,
AS_HELP_STRING([--with-lang],
[Use this option to build LibreOffice with additional language support.
@@ -9665,6 +9675,30 @@ fi
AC_SUBST(SYSTEM_MESA_HEADERS)
AC_SUBST(ENABLE_OPENGL)
+dnl =================================================
+dnl Check whether the OpenCL libraries are available
+dnl =================================================
+
+OPENCL_LIBS=
+OPENCL_CFLAGS=
+ENABLE_OPENCL=
+AC_MSG_CHECKING([opencl sdk])
+if test "z$with_opencl_sdk" = "z"; then
+ AC_MSG_RESULT([no])
+else
+ if test -d $with_opencl_sdk/include; then
+ ENABLE_OPENCL=TRUE
+ OPENCL_CFLAGS="-I $with_opencl_sdk/include"
+ OPENCL_LIBS="-L $with_opencl_sdk/lib/x86 -lOpenCL"
+ AC_MSG_RESULT([found at path $with_opencl_sdk])
+ else
+ AC_MSG_ERROR([no headers found found at $with_opencl_sdk/include ])
+ fi
+fi
+AC_SUBST(OPENCL_CFLAGS)
+AC_SUBST(OPENCL_LIBS)
+AC_SUBST(ENABLE_OPENCL)
+
# presenter minimizer extension?
AC_MSG_CHECKING([whether to build the Presentation Minimizer extension])
if test "x$enable_ext_presenter_minimizer" != "xno" -a "x$enable_extension_integration" != "xno"; then
More information about the Libreoffice-commits
mailing list