You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by ma...@apache.org on 2022/11/15 23:03:05 UTC
[tvm] branch main updated: [OpenCL] Introduce OpenCL wrapper to TVM (#13362)
This is an automated email from the ASF dual-hosted git repository.
masahi pushed a commit to branch main
in repository https://gitbox.apache.org/repos/asf/tvm.git
The following commit(s) were added to refs/heads/main by this push:
new 2bb3382c89 [OpenCL] Introduce OpenCL wrapper to TVM (#13362)
2bb3382c89 is described below
commit 2bb3382c89419820dab51a5a0fe4f09bb7ce6ecd
Author: Egor Churaev <eg...@gmail.com>
AuthorDate: Wed Nov 16 02:02:59 2022 +0300
[OpenCL] Introduce OpenCL wrapper to TVM (#13362)
* [OpenCL] Introduce OpenCL wrapper to TVM
This wrapper helps dynamically loading OpenCL library. It allows us to
avoid of looking for and copying OpenCL library to host, looking for
OpenCL SDK.
* Update apps and documentation
* Apply comments
* Apply comments and fix Android build
Also, use OpenCL wrapper by default and fix Windows build
* Apply comments
* Update LICENSE file
---
.gitmodules | 3 +
3rdparty/OpenCL-Headers | 1 +
LICENSE | 1 +
apps/android_camera/app/src/main/jni/Android.mk | 1 +
.../android_camera/app/src/main/jni/make/config.mk | 2 +-
apps/android_camera/app/src/main/jni/tvm_runtime.h | 2 +
apps/android_deploy/README.md | 34 +-
apps/android_deploy/app/src/main/jni/Android.mk | 3 +-
.../android_deploy/app/src/main/jni/make/config.mk | 2 +-
apps/android_deploy/app/src/main/jni/tvm_runtime.h | 3 +
apps/android_rpc/README.md | 32 +-
apps/android_rpc/app/src/main/jni/Android.mk | 3 +-
apps/android_rpc/app/src/main/jni/make/config.mk | 2 +-
apps/android_rpc/app/src/main/jni/tvm_runtime.h | 1 +
apps/cpp_rpc/README.md | 10 +-
cmake/config.cmake | 3 +-
cmake/modules/OpenCL.cmake | 30 +-
cmake/utils/FindOpenCL.cmake | 2 +-
.../deploy_models/deploy_model_on_android.py | 5 +-
src/runtime/opencl/opencl_module.cc | 2 +-
src/runtime/opencl/opencl_wrapper/README.md | 25 +
.../opencl/opencl_wrapper/opencl_wrapper.cc | 574 +++++++++++++++++++++
tests/cpp-runtime/opencl/opencl_timer_test.cc | 4 +-
23 files changed, 663 insertions(+), 82 deletions(-)
diff --git a/.gitmodules b/.gitmodules
index e03336443d..66fd0390cf 100644
--- a/.gitmodules
+++ b/.gitmodules
@@ -16,3 +16,6 @@
[submodule "3rdparty/cutlass"]
path = 3rdparty/cutlass
url = https://github.com/NVIDIA/cutlass.git
+[submodule "3rdparty/OpenCL-Headers"]
+ path = 3rdparty/OpenCL-Headers
+ url = https://github.com/KhronosGroup/OpenCL-Headers.git
diff --git a/3rdparty/OpenCL-Headers b/3rdparty/OpenCL-Headers
new file mode 160000
index 0000000000..b590a6bfe0
--- /dev/null
+++ b/3rdparty/OpenCL-Headers
@@ -0,0 +1 @@
+Subproject commit b590a6bfe034ea3a418b7b523e3490956bcb367a
diff --git a/LICENSE b/LICENSE
index 345026985b..6524d530de 100644
--- a/LICENSE
+++ b/LICENSE
@@ -211,6 +211,7 @@ Apache Software Foundation License 2.0
3rdparty/dlpack
3rdparty/dmlc-core
+3rdparty/OpenCL-Headers
BSD 2-clause License
diff --git a/apps/android_camera/app/src/main/jni/Android.mk b/apps/android_camera/app/src/main/jni/Android.mk
index 513666a4ec..2201f66965 100644
--- a/apps/android_camera/app/src/main/jni/Android.mk
+++ b/apps/android_camera/app/src/main/jni/Android.mk
@@ -41,6 +41,7 @@ LOCAL_C_INCLUDES := $(ROOT_PATH)/include \
$(ROOT_PATH)/src/runtime/rpc \
$(ROOT_PATH)/3rdparty/dlpack/include \
$(ROOT_PATH)/3rdparty/dmlc-core/include \
+ $(ROOT_PATH)/3rdparty/OpenCL-Headers \
$(MY_PATH)
LOCAL_MODULE = tvm4j_runtime_packed
diff --git a/apps/android_camera/app/src/main/jni/make/config.mk b/apps/android_camera/app/src/main/jni/make/config.mk
index 49e332665a..1f601b9afb 100644
--- a/apps/android_camera/app/src/main/jni/make/config.mk
+++ b/apps/android_camera/app/src/main/jni/make/config.mk
@@ -34,7 +34,7 @@ APP_ABI = all
APP_PLATFORM = android-24
# whether enable OpenCL during compile
-USE_OPENCL = 0
+USE_OPENCL = 1
# whether to enable Vulkan during compile
USE_VULKAN = 0
diff --git a/apps/android_camera/app/src/main/jni/tvm_runtime.h b/apps/android_camera/app/src/main/jni/tvm_runtime.h
index 6585347801..0aac7f170a 100644
--- a/apps/android_camera/app/src/main/jni/tvm_runtime.h
+++ b/apps/android_camera/app/src/main/jni/tvm_runtime.h
@@ -62,6 +62,8 @@
#ifdef TVM_OPENCL_RUNTIME
#include "../src/runtime/opencl/opencl_device_api.cc"
#include "../src/runtime/opencl/opencl_module.cc"
+#include "../src/runtime/opencl/opencl_wrapper/opencl_wrapper.cc"
+#include "../src/runtime/opencl/texture_pool.cc"
#include "../src/runtime/source_utils.cc"
#endif
diff --git a/apps/android_deploy/README.md b/apps/android_deploy/README.md
index 32e601840f..4cfd9eb9da 100644
--- a/apps/android_deploy/README.md
+++ b/apps/android_deploy/README.md
@@ -21,7 +21,7 @@ This folder contains Android Demo app that allows us to show how to deploy model
You will need [JDK](http://www.oracle.com/technetwork/java/javase/downloads/jdk8-downloads-2133151.html), [Android SDK](https://developer.android.com/studio/index.html), [Android NDK](https://developer.android.com/ndk) and an Android device to use this. Make sure the `ANDROID_HOME` variable already points to your Android SDK folder or set it using `export ANDROID_HOME=[Path to your Android SDK, e.g., ~/Android/sdk]`. We use [Gradle](https://gradle.org) to build. Please follow [the instal [...]
-Alternatively, you may execute Docker image we provide which contains the required packages. Use the command below to build the image and enter interactive session. Note, that building with OpenCL was not tested from Docker.
+Alternatively, you may execute Docker image we provide which contains the required packages. Use the command below to build the image and enter interactive session.
```bash
./docker/build.sh demo_android -it bash
@@ -50,7 +50,7 @@ dependencies {
}
```
-Application default has CPU version TVM runtime flavor and follow below instruction to setup.
+Application default has CPU and GPU (OpenCL) versions TVM runtime flavor and follow below instruction to setup.
In `app/src/main/jni/make` you will find JNI Makefile config `config.mk` and copy it to `app/src/main/jni` and modify it.
```bash
@@ -64,9 +64,6 @@ Here's a piece of example for `config.mk`.
APP_ABI = arm64-v8a
APP_PLATFORM = android-17
-
-# whether enable OpenCL during compile
-USE_OPENCL = 0
```
Now use Gradle to compile JNI, resolve Java dependencies and build the Android application together with tvm4j. Run following script to generate the apk file.
@@ -82,28 +79,11 @@ Upload `tvmdemo-release.apk` to your Android device and install it.
### Build with OpenCL
-Application does not link with OpenCL library unless you configure it to. Modify JNI Makefile config `app/src/main/jni` with proper target OpenCL configuration.
-
-Here's a piece of example for `config.mk`.
-
-```makefile
-APP_ABI = arm64-v8a
-
-APP_PLATFORM = android-17
-
-# whether enable OpenCL during compile
-USE_OPENCL = 1
-
-# the additional include headers you want to add, e.g., SDK_PATH/adrenosdk/Development/Inc
-ADD_C_INCLUDES = /opt/adrenosdk-osx/Development/Inc
-
-# the additional link libs you want to add, e.g., ANDROID_LIB_PATH/libOpenCL.so
-ADD_LDLIBS = libOpenCL.so
-```
-
-Note that you should specify the correct GPU development headers for your android device. Run `adb shell dumpsys | grep GLES` to find out what GPU your android device uses. It is very likely the library (libOpenCL.so) is already present on the mobile device. For instance, I found it under `/system/vendor/lib64`. You can do `adb pull /system/vendor/lib64/libOpenCL.so ./` to get the file to your desktop.
-
-After you setup the `config.mk`, follow the instructions in [Build APK](#buildapk) to build the Android package with OpenCL flavor.
+Application is building with OpenCL support by default.
+[OpenCL-wrapper](../../src/runtime/opencl/opencl_wrapper) is used and will dynamically load OpenCL library on the device.
+If the device doesn't have OpenCL library on it, then you'll see in the runtime that OpenCL library cannot be opened.
+If you want to build this application without OpenCL then set `USE_OPENCL = 0`
+in [config.mk](./app/src/main/jni/make/config.mk)
## Cross Compile and Run on Android Devices
diff --git a/apps/android_deploy/app/src/main/jni/Android.mk b/apps/android_deploy/app/src/main/jni/Android.mk
index 1b06a6bdb8..ad9cee9bbd 100644
--- a/apps/android_deploy/app/src/main/jni/Android.mk
+++ b/apps/android_deploy/app/src/main/jni/Android.mk
@@ -38,7 +38,8 @@ LOCAL_LDFLAGS := -L$(SYSROOT)/usr/lib/ -llog
LOCAL_C_INCLUDES := $(ROOT_PATH)/include \
$(ROOT_PATH)/3rdparty/dlpack/include \
- $(ROOT_PATH)/3rdparty/dmlc-core/include
+ $(ROOT_PATH)/3rdparty/dmlc-core/include \
+ $(ROOT_PATH)/3rdparty/OpenCL-Headers
LOCAL_MODULE = tvm4j_runtime_packed
diff --git a/apps/android_deploy/app/src/main/jni/make/config.mk b/apps/android_deploy/app/src/main/jni/make/config.mk
index bcd56e3789..b06f42b264 100644
--- a/apps/android_deploy/app/src/main/jni/make/config.mk
+++ b/apps/android_deploy/app/src/main/jni/make/config.mk
@@ -34,7 +34,7 @@ APP_ABI = all
APP_PLATFORM = android-17
# whether enable OpenCL during compile
-USE_OPENCL = 0
+USE_OPENCL = 1
# the additional include headers you want to add, e.g., SDK_PATH/adrenosdk/Development/Inc
ADD_C_INCLUDES =
diff --git a/apps/android_deploy/app/src/main/jni/tvm_runtime.h b/apps/android_deploy/app/src/main/jni/tvm_runtime.h
index 725b5e1d3b..a2f10701d6 100644
--- a/apps/android_deploy/app/src/main/jni/tvm_runtime.h
+++ b/apps/android_deploy/app/src/main/jni/tvm_runtime.h
@@ -47,4 +47,7 @@
#ifdef TVM_OPENCL_RUNTIME
#include "../src/runtime/opencl/opencl_device_api.cc"
#include "../src/runtime/opencl/opencl_module.cc"
+#include "../src/runtime/opencl/opencl_wrapper/opencl_wrapper.cc"
+#include "../src/runtime/opencl/texture_pool.cc"
+#include "../src/runtime/source_utils.cc"
#endif
diff --git a/apps/android_rpc/README.md b/apps/android_rpc/README.md
index 2e301af6d9..d0a11b6121 100644
--- a/apps/android_rpc/README.md
+++ b/apps/android_rpc/README.md
@@ -74,33 +74,11 @@ $ANDROID_HOME/platform-tools/adb uninstall org.apache.tvm.tvmrpc
### Build with OpenCL
-This application does not link any OpenCL library unless you configure it to. In `app/src/main/jni/make` you will find JNI Makefile config `config.mk`. Copy it to `app/src/main/jni` and modify it.
-
-```bash
-cd apps/android_rpc/app/src/main/jni
-cp make/config.mk .
-```
-
-Here's a piece of example for `config.mk`.
-
-```makefile
-APP_ABI = arm64-v8a
-
-APP_PLATFORM = android-17
-
-# whether enable OpenCL during compile
-USE_OPENCL = 1
-
-# the additional include headers you want to add, e.g., SDK_PATH/adrenosdk/Development/Inc
-ADD_C_INCLUDES = /opt/adrenosdk-osx/Development/Inc
-
-# the additional link libs you want to add, e.g., ANDROID_LIB_PATH/libOpenCL.so
-ADD_LDLIBS = libOpenCL.so
-```
-
-Note that you should specify the correct GPU development headers for your android device. Run `adb shell dumpsys | grep GLES` to find out what GPU your android device uses. It is very likely the library (libOpenCL.so) is already present on the mobile device. For instance, I found it under `/system/vendor/lib64`. You can do `adb pull /system/vendor/lib64/libOpenCL.so ./` to get the file to your desktop.
-
-After you setup the `config.mk`, follow the instructions in [Build APK](#buildapk) to build the Android package.
+Application is building with OpenCL support by default.
+[OpenCL-wrapper](../../src/runtime/opencl/opencl_wrapper) is used and will dynamically load OpenCL library on the device.
+If the device doesn't have OpenCL library on it, then you'll see in the runtime that OpenCL library cannot be opened.
+If you want to build this application without OpenCL then set `USE_OPENCL = 0`
+in [config.mk](./app/src/main/jni/make/config.mk)
## Cross Compile and Run on Android Devices
diff --git a/apps/android_rpc/app/src/main/jni/Android.mk b/apps/android_rpc/app/src/main/jni/Android.mk
index 1b06a6bdb8..ad9cee9bbd 100644
--- a/apps/android_rpc/app/src/main/jni/Android.mk
+++ b/apps/android_rpc/app/src/main/jni/Android.mk
@@ -38,7 +38,8 @@ LOCAL_LDFLAGS := -L$(SYSROOT)/usr/lib/ -llog
LOCAL_C_INCLUDES := $(ROOT_PATH)/include \
$(ROOT_PATH)/3rdparty/dlpack/include \
- $(ROOT_PATH)/3rdparty/dmlc-core/include
+ $(ROOT_PATH)/3rdparty/dmlc-core/include \
+ $(ROOT_PATH)/3rdparty/OpenCL-Headers
LOCAL_MODULE = tvm4j_runtime_packed
diff --git a/apps/android_rpc/app/src/main/jni/make/config.mk b/apps/android_rpc/app/src/main/jni/make/config.mk
index 851430cd42..855a0af190 100644
--- a/apps/android_rpc/app/src/main/jni/make/config.mk
+++ b/apps/android_rpc/app/src/main/jni/make/config.mk
@@ -34,7 +34,7 @@ APP_ABI = all
APP_PLATFORM = android-24
# whether enable OpenCL during compile
-USE_OPENCL = 0
+USE_OPENCL = 1
# whether to enable Vulkan during compile
USE_VULKAN = 0
diff --git a/apps/android_rpc/app/src/main/jni/tvm_runtime.h b/apps/android_rpc/app/src/main/jni/tvm_runtime.h
index 543c9c8533..17a20bbaf9 100644
--- a/apps/android_rpc/app/src/main/jni/tvm_runtime.h
+++ b/apps/android_rpc/app/src/main/jni/tvm_runtime.h
@@ -64,6 +64,7 @@
#ifdef TVM_OPENCL_RUNTIME
#include "../src/runtime/opencl/opencl_device_api.cc"
#include "../src/runtime/opencl/opencl_module.cc"
+#include "../src/runtime/opencl/opencl_wrapper/opencl_wrapper.cc"
#include "../src/runtime/opencl/texture_pool.cc"
#include "../src/runtime/source_utils.cc"
#endif
diff --git a/apps/cpp_rpc/README.md b/apps/cpp_rpc/README.md
index d073fca819..58eb68055f 100644
--- a/apps/cpp_rpc/README.md
+++ b/apps/cpp_rpc/README.md
@@ -37,7 +37,15 @@ This folder contains a simple recipe to make RPC server in c++.
# Path to the desired C++ cross compiler
set(CMAKE_CXX_COMPILER /path/to/cross/compiler/executable)
```
-- If linking against a custom device OpenCL library is needed, in the config specify the path to the OpenCL SDK containing the include/CL headers and lib/ or lib64/libOpenCL.so:
+- If you need to build cpp_rpc with OpenCL support, specify variable `USE_OPENCL` in the config:
+ ```
+ set(USE_OPENCL ON)
+ ```
+ In this case [OpenCL-wrapper](../../src/runtime/opencl/opencl_wrapper) or OpenCL installed to your system will be used.
+ When OpenCL-wrapper is used, it will dynamically load OpenCL library on the device.
+ If the device doesn't have OpenCL library on it, then you'll see in the runtime that OpenCL library cannot be opened.
+
+ If linking against a custom device OpenCL library is needed, in the config specify the path to the OpenCL SDK containing the include/CL headers and lib/ or lib64/libOpenCL.so:
```
set(USE_OPENCL /path/to/opencl-sdk)
```
diff --git a/cmake/config.cmake b/cmake/config.cmake
index 22a548d298..679f5c459e 100644
--- a/cmake/config.cmake
+++ b/cmake/config.cmake
@@ -65,7 +65,8 @@ set(USE_AOCL OFF)
# Whether enable OpenCL runtime
#
# Possible values:
-# - ON: enable OpenCL with cmake's auto search
+# - ON: enable OpenCL with OpenCL wrapper to remove dependency during build
+# time and trigger dynamic search and loading of OpenCL in runtime
# - OFF: disable OpenCL
# - /path/to/opencl-sdk: use specific path to opencl-sdk
set(USE_OPENCL OFF)
diff --git a/cmake/modules/OpenCL.cmake b/cmake/modules/OpenCL.cmake
index 430af7e872..e738df7c56 100644
--- a/cmake/modules/OpenCL.cmake
+++ b/cmake/modules/OpenCL.cmake
@@ -15,15 +15,6 @@
# specific language governing permissions and limitations
# under the License.
-# OPENCL Module
-find_opencl(${USE_OPENCL})
-
-if(OpenCL_FOUND)
- # always set the includedir when cuda is available
- # avoid global retrigger of cmake
- include_directories(SYSTEM ${OpenCL_INCLUDE_DIRS})
-endif(OpenCL_FOUND)
-
if(USE_SDACCEL)
message(STATUS "Build with SDAccel support")
tvm_file_glob(GLOB RUNTIME_SDACCEL_SRCS src/runtime/opencl/sdaccel/*.cc)
@@ -49,12 +40,23 @@ else()
endif(USE_AOCL)
if(USE_OPENCL)
- if (NOT OpenCL_FOUND)
- find_package(OpenCL REQUIRED)
- endif()
- message(STATUS "Build with OpenCL support")
tvm_file_glob(GLOB RUNTIME_OPENCL_SRCS src/runtime/opencl/*.cc)
- list(APPEND TVM_RUNTIME_LINKER_LIBS ${OpenCL_LIBRARIES})
+
+ if(${USE_OPENCL} MATCHES ${IS_TRUE_PATTERN})
+ message(WARNING "Build with OpenCL wrapper")
+ file_glob_append(RUNTIME_OPENCL_SRCS
+ "src/runtime/opencl/opencl_wrapper/opencl_wrapper.cc"
+ )
+ include_directories(SYSTEM "3rdparty/OpenCL-Headers")
+ else()
+ find_opencl(${USE_OPENCL})
+ if(NOT OpenCL_FOUND)
+ message(FATAL_ERROR "Error! Cannot find specified OpenCL library")
+ endif()
+ message(STATUS "Build with OpenCL support")
+ include_directories(SYSTEM ${OpenCL_INCLUDE_DIRS})
+ list(APPEND TVM_RUNTIME_LINKER_LIBS ${OpenCL_LIBRARIES})
+ endif()
if(DEFINED USE_OPENCL_GTEST AND EXISTS ${USE_OPENCL_GTEST})
file_glob_append(RUNTIME_OPENCL_SRCS
diff --git a/cmake/utils/FindOpenCL.cmake b/cmake/utils/FindOpenCL.cmake
index f2931332fc..8eb35ab399 100644
--- a/cmake/utils/FindOpenCL.cmake
+++ b/cmake/utils/FindOpenCL.cmake
@@ -21,7 +21,7 @@
# Usage:
# find_opencl(${USE_OPENCL})
#
-# - When USE_OPENCL=ON, use auto search
+# - When USE_OPENCL=ON, use OpenCL wrapper for dynamic linking
# - When USE_OPENCL=/path/to/opencl-sdk-path, use the sdk.
# Can be useful when cross compiling and cannot rely on
# CMake to provide the correct library as part of the
diff --git a/gallery/how_to/deploy_models/deploy_model_on_android.py b/gallery/how_to/deploy_models/deploy_model_on_android.py
index 10e108239e..4bf86e2981 100644
--- a/gallery/how_to/deploy_models/deploy_model_on_android.py
+++ b/gallery/how_to/deploy_models/deploy_model_on_android.py
@@ -137,11 +137,10 @@ from tvm.contrib.download import download_testdata
#
# # the additional include headers you want to add, e.g., SDK_PATH/adrenosdk/Development/Inc
# ADD_C_INCLUDES += /work/adrenosdk-linux-5_0/Development/Inc
-# # downloaded from https://github.com/KhronosGroup/OpenCL-Headers
-# ADD_C_INCLUDES += /usr/local/OpenCL-Headers/
+# ADD_C_INCLUDES =
#
# # the additional link libs you want to add, e.g., ANDROID_LIB_PATH/libOpenCL.so
-# ADD_LDLIBS = /workspace/pull-from-android-device/libOpenCL.so
+# ADD_LDLIBS =
#
# .. note::
#
diff --git a/src/runtime/opencl/opencl_module.cc b/src/runtime/opencl/opencl_module.cc
index 9ae80d59d5..2fb157aac6 100644
--- a/src/runtime/opencl/opencl_module.cc
+++ b/src/runtime/opencl/opencl_module.cc
@@ -232,7 +232,7 @@ cl_kernel OpenCLModuleNode::InstallKernel(cl::OpenCLWorkspace* w, cl::OpenCLThre
cl_int err;
cl_device_id dev = w->devices[device_id];
programs_[func_name][device_id] =
- clCreateProgramWithBinary(w->context, 1, &dev, &len, &s, NULL, &err);
+ clCreateProgramWithBinary(w->context, 1, &dev, &len, &s, nullptr, &err);
OPENCL_CHECK_ERROR(err);
} else {
LOG(FATAL) << "Unknown OpenCL format " << fmt_;
diff --git a/src/runtime/opencl/opencl_wrapper/README.md b/src/runtime/opencl/opencl_wrapper/README.md
new file mode 100644
index 0000000000..7597a442c1
--- /dev/null
+++ b/src/runtime/opencl/opencl_wrapper/README.md
@@ -0,0 +1,25 @@
+<!--- 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 -->
+
+<!--- Unless required by applicable law or agreed to in writing, -->
+<!--- software distributed under the License is distributed on an -->
+<!--- "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY -->
+<!--- KIND, either express or implied. See the License for the -->
+<!--- specific language governing permissions and limitations -->
+<!--- under the License. -->
+
+# OpenCL Wrapper
+
+This wrapper helps dynamically loading OpenCL library. It allows us to avoid of
+looking for and copying library from phone to host, looking for OpenCL SDK.
+
+This can be done because OpenCL is a standard and number of functions are
+limited. We can safely wrap all required functions and their number will not
+grow.
diff --git a/src/runtime/opencl/opencl_wrapper/opencl_wrapper.cc b/src/runtime/opencl/opencl_wrapper/opencl_wrapper.cc
new file mode 100644
index 0000000000..c447ebcb53
--- /dev/null
+++ b/src/runtime/opencl/opencl_wrapper/opencl_wrapper.cc
@@ -0,0 +1,574 @@
+/*
+ * 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
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied. See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+/*!
+ * \file opencl_wrapper.cc
+ * \brief This wrapper is actual for OpenCL 1.2, but can be easily upgraded
+ * when TVM will use newer version of OpenCL
+ */
+
+#define CL_TARGET_OPENCL_VERSION 120
+#include <CL/cl.h>
+#include <CL/cl_gl.h>
+
+#if defined(_WIN32)
+#include <windows.h>
+#else
+#include <dlfcn.h>
+#endif
+
+#define DMLC_USE_LOGGING_LIBRARY <tvm/runtime/logging.h>
+#include <tvm/runtime/logging.h>
+
+#include <vector>
+
+namespace {
+#if defined(__APPLE__) || defined(__MACOSX)
+static const std::vector<const char*> default_so_paths = {
+ "libOpenCL.so", "/System/Library/Frameworks/OpenCL.framework/OpenCL"};
+#elif defined(__ANDROID__)
+static const std::vector<const char*> default_so_paths = {
+ "libOpenCL.so",
+ "/system/lib64/libOpenCL.so",
+ "/system/vendor/lib64/libOpenCL.so",
+ "/system/vendor/lib64/egl/libGLES_mali.so",
+ "/system/vendor/lib64/libPVROCL.so",
+ "/data/data/org.pocl.libs/files/lib64/libpocl.so",
+ "/system/lib/libOpenCL.so",
+ "/system/vendor/lib/libOpenCL.so",
+ "/system/vendor/lib/egl/libGLES_mali.so",
+ "/system/vendor/lib/libPVROCL.so",
+ "/data/data/org.pocl.libs/files/lib/libpocl.so"};
+#elif defined(_WIN32)
+static const std::vector<const TCHAR*> default_so_paths = {__TEXT("OpenCL.dll")};
+#elif defined(__linux__)
+static const std::vector<const char*> default_so_paths = {"libOpenCL.so",
+ "/usr/lib/libOpenCL.so",
+ "/usr/local/lib/libOpenCL.so",
+ "/usr/local/lib/libpocl.so",
+ "/usr/lib64/libOpenCL.so",
+ "/usr/lib32/libOpenCL.so"};
+#endif
+
+class LibOpenCLWrapper {
+ public:
+ static LibOpenCLWrapper& getInstance() {
+ static LibOpenCLWrapper instance;
+ return instance;
+ }
+ LibOpenCLWrapper(const LibOpenCLWrapper&) = delete;
+ LibOpenCLWrapper& operator=(const LibOpenCLWrapper&) = delete;
+ void* getOpenCLFunction(const char* funcName) {
+ if (m_libHandler == nullptr) openLibOpenCL();
+#if defined(_WIN32)
+ return GetProcAddress(m_libHandler, funcName);
+#else
+ return dlsym(m_libHandler, funcName);
+#endif
+ }
+
+ private:
+ LibOpenCLWrapper() {}
+ ~LibOpenCLWrapper() {
+#if defined(_WIN32)
+ if (m_libHandler) FreeLibrary(m_libHandler);
+#else
+ if (m_libHandler) dlclose(m_libHandler);
+#endif
+ }
+ void openLibOpenCL() {
+ for (const auto it : default_so_paths) {
+#if defined(_WIN32)
+ m_libHandler = LoadLibrary(it);
+#else
+ m_libHandler = dlopen(it, RTLD_LAZY);
+#endif
+ if (m_libHandler != nullptr) return;
+ }
+ ICHECK(m_libHandler != nullptr) << "Error! Cannot open libOpenCL!";
+ }
+
+ private:
+#if defined(_WIN32)
+ HMODULE m_libHandler = nullptr;
+#else
+ void* m_libHandler = nullptr;
+#endif
+};
+
+// Function pointers declaration
+using f_pfn_notify = void (*)(const char*, const void*, size_t, void*);
+using f_clGetPlatformIDs = cl_int (*)(cl_uint, cl_platform_id*, cl_uint*);
+using f_clGetPlatformInfo = cl_int (*)(cl_platform_id, cl_platform_info, size_t, void*, size_t*);
+using f_clGetDeviceIDs = cl_int (*)(cl_platform_id, cl_device_type, cl_uint, cl_device_id*,
+ cl_uint*);
+using f_clGetDeviceInfo = cl_int (*)(cl_device_id, cl_device_info, size_t, void*, size_t*);
+using f_clCreateContext = cl_context (*)(const cl_context_properties*, cl_uint, const cl_device_id*,
+ f_pfn_notify, void*, cl_int*);
+using f_clReleaseContext = cl_int (*)(cl_context);
+using f_clReleaseCommandQueue = cl_int (*)(cl_command_queue);
+using f_clGetCommandQueueInfo = cl_int (*)(cl_command_queue, cl_command_queue_info, size_t, void*,
+ size_t*);
+using f_clCreateBuffer = cl_mem (*)(cl_context, cl_mem_flags, size_t, void*, cl_int*);
+using f_clCreateImage = cl_mem (*)(cl_context, cl_mem_flags, const cl_image_format*,
+ const cl_image_desc*, void*, cl_int*);
+using f_clReleaseMemObject = cl_int (*)(cl_mem);
+using f_clCreateProgramWithSource = cl_program (*)(cl_context, cl_uint, const char**, const size_t*,
+ cl_int*);
+using f_clCreateProgramWithBinary = cl_program (*)(cl_context, cl_uint, const cl_device_id*,
+ const size_t*, const unsigned char**, cl_int*,
+ cl_int*);
+using f_clReleaseProgram = cl_int (*)(cl_program);
+using f_clBuildProgram = cl_int (*)(cl_program, cl_uint, const cl_device_id*, const char*,
+ void (*pfn_notify)(cl_program program, void* user_data), void*);
+using f_clGetProgramBuildInfo = cl_int (*)(cl_program, cl_device_id, cl_program_build_info, size_t,
+ void*, size_t*);
+using f_clCreateKernel = cl_kernel (*)(cl_program, const char*, cl_int*);
+using f_clReleaseKernel = cl_int (*)(cl_kernel);
+using f_clSetKernelArg = cl_int (*)(cl_kernel, cl_uint, size_t, const void*);
+using f_clWaitForEvents = cl_int (*)(cl_uint, const cl_event*);
+using f_clCreateUserEvent = cl_event (*)(cl_context, cl_int*);
+using f_clGetEventProfilingInfo = cl_int (*)(cl_event, cl_profiling_info, size_t, void*, size_t*);
+using f_clFlush = cl_int (*)(cl_command_queue);
+using f_clFinish = cl_int (*)(cl_command_queue);
+using f_clEnqueueReadBuffer = cl_int (*)(cl_command_queue, cl_mem, cl_bool, size_t, size_t, void*,
+ cl_uint, const cl_event*, cl_event*);
+using f_clEnqueueWriteBuffer = cl_int (*)(cl_command_queue, cl_mem, cl_bool, size_t, size_t,
+ const void*, cl_uint, const cl_event*, cl_event*);
+using f_clEnqueueCopyBuffer = cl_int (*)(cl_command_queue, cl_mem, cl_mem, size_t, size_t, size_t,
+ cl_uint, const cl_event*, cl_event*);
+using f_clEnqueueReadImage = cl_int (*)(cl_command_queue, cl_mem, cl_bool, const size_t*,
+ const size_t*, size_t, size_t, void*, cl_uint,
+ const cl_event*, cl_event*);
+using f_clEnqueueWriteImage = cl_int (*)(cl_command_queue, cl_mem, cl_bool, const size_t*,
+ const size_t*, size_t, size_t, const void*, cl_uint,
+ const cl_event*, cl_event*);
+using f_clEnqueueCopyImage = cl_int (*)(cl_command_queue, cl_mem, cl_mem, const size_t*,
+ const size_t*, const size_t*, cl_uint, const cl_event*,
+ cl_event*);
+using f_clEnqueueCopyImageToBuffer = cl_int (*)(cl_command_queue, cl_mem, cl_mem, const size_t*,
+ const size_t*, size_t, cl_uint, const cl_event*,
+ cl_event*);
+using f_clEnqueueCopyBufferToImage = cl_int (*)(cl_command_queue, cl_mem, cl_mem, size_t,
+ const size_t*, const size_t*, cl_uint,
+ const cl_event*, cl_event*);
+using f_clEnqueueNDRangeKernel = cl_int (*)(cl_command_queue, cl_kernel, cl_uint, const size_t*,
+ const size_t*, const size_t*, cl_uint, const cl_event*,
+ cl_event*);
+using f_clCreateCommandQueue = cl_command_queue (*)(cl_context, cl_device_id,
+ cl_command_queue_properties, cl_int*);
+} // namespace
+
+cl_int clGetPlatformIDs(cl_uint num_entries, cl_platform_id* platforms, cl_uint* num_platforms) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clGetPlatformIDs)lib.getOpenCLFunction("clGetPlatformIDs");
+ if (func) {
+ return func(num_entries, platforms, num_platforms);
+ } else {
+ return CL_INVALID_PLATFORM;
+ }
+}
+
+cl_int clGetPlatformInfo(cl_platform_id platform, cl_platform_info param_name,
+ size_t param_value_size, void* param_value, size_t* param_value_size_ret) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clGetPlatformInfo)lib.getOpenCLFunction("clGetPlatformInfo");
+ if (func) {
+ return func(platform, param_name, param_value_size, param_value, param_value_size_ret);
+ } else {
+ return CL_INVALID_PLATFORM;
+ }
+}
+
+cl_int clGetDeviceIDs(cl_platform_id platform, cl_device_type device_type, cl_uint num_entries,
+ cl_device_id* devices, cl_uint* num_devices) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clGetDeviceIDs)lib.getOpenCLFunction("clGetDeviceIDs");
+ if (func) {
+ return func(platform, device_type, num_entries, devices, num_devices);
+ } else {
+ return CL_INVALID_PLATFORM;
+ }
+}
+
+cl_int clGetDeviceInfo(cl_device_id device, cl_device_info param_name, size_t param_value_size,
+ void* param_value, size_t* param_value_size_ret) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clGetDeviceInfo)lib.getOpenCLFunction("clGetDeviceInfo");
+ if (func) {
+ return func(device, param_name, param_value_size, param_value, param_value_size_ret);
+ } else {
+ return CL_INVALID_PLATFORM;
+ }
+}
+
+cl_context clCreateContext(const cl_context_properties* properties, cl_uint num_devices,
+ const cl_device_id* devices,
+ void (*pfn_notify)(const char*, const void*, size_t, void*),
+ void* user_data, cl_int* errcode_ret) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clCreateContext)lib.getOpenCLFunction("clCreateContext");
+ if (func) {
+ return func(properties, num_devices, devices, pfn_notify, user_data, errcode_ret);
+ } else {
+ return nullptr;
+ }
+}
+
+cl_int clReleaseContext(cl_context context) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clReleaseContext)lib.getOpenCLFunction("clReleaseContext");
+
+ if (func) {
+ return func(context);
+ } else {
+ return CL_INVALID_PLATFORM;
+ }
+}
+
+cl_int clReleaseCommandQueue(cl_command_queue command_queue) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clReleaseCommandQueue)lib.getOpenCLFunction("clReleaseCommandQueue");
+ if (func) {
+ return func(command_queue);
+ } else {
+ return CL_INVALID_PLATFORM;
+ }
+}
+
+cl_int clGetCommandQueueInfo(cl_command_queue command_queue, cl_command_queue_info param_name,
+ size_t param_value_size, void* param_value,
+ size_t* param_value_size_ret) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clGetCommandQueueInfo)lib.getOpenCLFunction("clGetCommandQueueInfo");
+ if (func) {
+ return func(command_queue, param_name, param_value_size, param_value, param_value_size_ret);
+ } else {
+ return CL_INVALID_PLATFORM;
+ }
+}
+
+cl_mem clCreateBuffer(cl_context context, cl_mem_flags flags, size_t size, void* host_ptr,
+ cl_int* errcode_ret) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clCreateBuffer)lib.getOpenCLFunction("clCreateBuffer");
+ if (func) {
+ return func(context, flags, size, host_ptr, errcode_ret);
+ } else {
+ return nullptr;
+ }
+}
+
+cl_mem clCreateImage(cl_context context, cl_mem_flags flags, const cl_image_format* image_format,
+ const cl_image_desc* image_desc, void* host_ptr, cl_int* errcode_ret) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clCreateImage)lib.getOpenCLFunction("clCreateImage");
+ if (func) {
+ return func(context, flags, image_format, image_desc, host_ptr, errcode_ret);
+ } else {
+ return nullptr;
+ }
+}
+
+cl_int clReleaseMemObject(cl_mem memobj) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clReleaseMemObject)lib.getOpenCLFunction("clReleaseMemObject");
+ if (func) {
+ return func(memobj);
+ } else {
+ return CL_INVALID_PLATFORM;
+ }
+}
+
+cl_program clCreateProgramWithSource(cl_context context, cl_uint count, const char** strings,
+ const size_t* lengths, cl_int* errcode_ret) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clCreateProgramWithSource)lib.getOpenCLFunction("clCreateProgramWithSource");
+ if (func) {
+ return func(context, count, strings, lengths, errcode_ret);
+ } else {
+ return nullptr;
+ }
+}
+
+cl_program clCreateProgramWithBinary(cl_context context, cl_uint num_devices,
+ const cl_device_id* device_list, const size_t* lengths,
+ const unsigned char** binaries, cl_int* binary_status,
+ cl_int* errcode_ret) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clCreateProgramWithBinary)lib.getOpenCLFunction("clCreateProgramWithBinary");
+ if (func) {
+ return func(context, num_devices, device_list, lengths, binaries, binary_status, errcode_ret);
+ } else {
+ return nullptr;
+ }
+}
+
+cl_int clReleaseProgram(cl_program program) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clReleaseProgram)lib.getOpenCLFunction("clReleaseProgram");
+ if (func) {
+ return func(program);
+ } else {
+ return CL_INVALID_PLATFORM;
+ }
+}
+
+cl_int clBuildProgram(cl_program program, cl_uint num_devices, const cl_device_id* device_list,
+ const char* options, void (*pfn_notify)(cl_program program, void* user_data),
+ void* user_data) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clBuildProgram)lib.getOpenCLFunction("clBuildProgram");
+ if (func) {
+ return func(program, num_devices, device_list, options, pfn_notify, user_data);
+ } else {
+ return CL_INVALID_PLATFORM;
+ }
+}
+
+cl_int clGetProgramBuildInfo(cl_program program, cl_device_id device,
+ cl_program_build_info param_name, size_t param_value_size,
+ void* param_value, size_t* param_value_size_ret) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clGetProgramBuildInfo)lib.getOpenCLFunction("clGetProgramBuildInfo");
+ if (func) {
+ return func(program, device, param_name, param_value_size, param_value, param_value_size_ret);
+ } else {
+ return CL_INVALID_PLATFORM;
+ }
+}
+
+cl_kernel clCreateKernel(cl_program program, const char* kernel_name, cl_int* errcode_ret) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clCreateKernel)lib.getOpenCLFunction("clCreateKernel");
+ if (func) {
+ return func(program, kernel_name, errcode_ret);
+ } else {
+ return nullptr;
+ }
+}
+
+cl_int clReleaseKernel(cl_kernel kernel) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clReleaseKernel)lib.getOpenCLFunction("clReleaseKernel");
+ if (func) {
+ return func(kernel);
+ } else {
+ return CL_INVALID_PLATFORM;
+ }
+}
+
+cl_int clSetKernelArg(cl_kernel kernel, cl_uint arg_index, size_t arg_size, const void* arg_value) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clSetKernelArg)lib.getOpenCLFunction("clSetKernelArg");
+ if (func) {
+ return func(kernel, arg_index, arg_size, arg_value);
+ } else {
+ return CL_INVALID_PLATFORM;
+ }
+}
+
+cl_int clWaitForEvents(cl_uint num_events, const cl_event* event_list) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clWaitForEvents)lib.getOpenCLFunction("clWaitForEvents");
+ if (func) {
+ return func(num_events, event_list);
+ } else {
+ return CL_INVALID_PLATFORM;
+ }
+}
+
+cl_event clCreateUserEvent(cl_context context, cl_int* errcode_ret) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clCreateUserEvent)lib.getOpenCLFunction("clCreateUserEvent");
+ if (func) {
+ return func(context, errcode_ret);
+ } else {
+ return nullptr;
+ }
+}
+
+cl_int clGetEventProfilingInfo(cl_event event, cl_profiling_info param_name,
+ size_t param_value_size, void* param_value,
+ size_t* param_value_size_ret) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clGetEventProfilingInfo)lib.getOpenCLFunction("clGetEventProfilingInfo");
+ if (func) {
+ return func(event, param_name, param_value_size, param_value, param_value_size_ret);
+ } else {
+ return CL_INVALID_PLATFORM;
+ }
+}
+
+cl_int clFlush(cl_command_queue command_queue) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clFlush)lib.getOpenCLFunction("clFlush");
+ if (func) {
+ return func(command_queue);
+ } else {
+ return CL_INVALID_PLATFORM;
+ }
+}
+
+cl_int clFinish(cl_command_queue command_queue) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clFinish)lib.getOpenCLFunction("clFinish");
+ if (func) {
+ return func(command_queue);
+ } else {
+ return CL_INVALID_PLATFORM;
+ }
+}
+
+cl_int clEnqueueReadBuffer(cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_read,
+ size_t offset, size_t size, void* ptr, cl_uint num_events_in_wait_list,
+ const cl_event* event_wait_list, cl_event* event) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clEnqueueReadBuffer)lib.getOpenCLFunction("clEnqueueReadBuffer");
+ if (func) {
+ return func(command_queue, buffer, blocking_read, offset, size, ptr, num_events_in_wait_list,
+ event_wait_list, event);
+ } else {
+ return CL_INVALID_PLATFORM;
+ }
+}
+
+cl_int clEnqueueWriteBuffer(cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_write,
+ size_t offset, size_t size, const void* ptr,
+ cl_uint num_events_in_wait_list, const cl_event* event_wait_list,
+ cl_event* event) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clEnqueueWriteBuffer)lib.getOpenCLFunction("clEnqueueWriteBuffer");
+ if (func) {
+ return func(command_queue, buffer, blocking_write, offset, size, ptr, num_events_in_wait_list,
+ event_wait_list, event);
+ } else {
+ return CL_INVALID_PLATFORM;
+ }
+}
+
+cl_int clEnqueueCopyBuffer(cl_command_queue command_queue, cl_mem src_buffer, cl_mem dst_buffer,
+ size_t src_offset, size_t dst_offset, size_t size,
+ cl_uint num_events_in_wait_list, const cl_event* event_wait_list,
+ cl_event* event) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clEnqueueCopyBuffer)lib.getOpenCLFunction("clEnqueueCopyBuffer");
+ if (func) {
+ return func(command_queue, src_buffer, dst_buffer, src_offset, dst_offset, size,
+ num_events_in_wait_list, event_wait_list, event);
+ } else {
+ return CL_INVALID_PLATFORM;
+ }
+}
+
+cl_int clEnqueueReadImage(cl_command_queue command_queue, cl_mem image, cl_bool blocking_read,
+ const size_t* origin, const size_t* region, size_t row_pitch,
+ size_t slice_pitch, void* ptr, cl_uint num_events_in_wait_list,
+ const cl_event* event_wait_list, cl_event* event) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clEnqueueReadImage)lib.getOpenCLFunction("clEnqueueReadImage");
+ if (func) {
+ return func(command_queue, image, blocking_read, origin, region, row_pitch, slice_pitch, ptr,
+ num_events_in_wait_list, event_wait_list, event);
+ } else {
+ return CL_INVALID_PLATFORM;
+ }
+}
+
+cl_int clEnqueueWriteImage(cl_command_queue command_queue, cl_mem image, cl_bool blocking_write,
+ const size_t* origin, const size_t* region, size_t input_row_pitch,
+ size_t input_slice_pitch, const void* ptr,
+ cl_uint num_events_in_wait_list, const cl_event* event_wait_list,
+ cl_event* event) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clEnqueueWriteImage)lib.getOpenCLFunction("clEnqueueWriteImage");
+ if (func) {
+ return func(command_queue, image, blocking_write, origin, region, input_row_pitch,
+ input_slice_pitch, ptr, num_events_in_wait_list, event_wait_list, event);
+ } else {
+ return CL_INVALID_PLATFORM;
+ }
+}
+
+cl_int clEnqueueCopyImage(cl_command_queue command_queue, cl_mem src_image, cl_mem dst_image,
+ const size_t* src_origin, const size_t* dst_origin, const size_t* region,
+ cl_uint num_events_in_wait_list, const cl_event* event_wait_list,
+ cl_event* event) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clEnqueueCopyImage)lib.getOpenCLFunction("clEnqueueCopyImage");
+ if (func) {
+ return func(command_queue, src_image, dst_image, src_origin, dst_origin, region,
+ num_events_in_wait_list, event_wait_list, event);
+ } else {
+ return CL_INVALID_PLATFORM;
+ }
+}
+
+cl_int clEnqueueCopyImageToBuffer(cl_command_queue command_queue, cl_mem src_image,
+ cl_mem dst_buffer, const size_t* src_origin, const size_t* region,
+ size_t dst_offset, cl_uint num_events_in_wait_list,
+ const cl_event* event_wait_list, cl_event* event) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clEnqueueCopyImageToBuffer)lib.getOpenCLFunction("clEnqueueCopyImageToBuffer");
+ if (func) {
+ return func(command_queue, src_image, dst_buffer, src_origin, region, dst_offset,
+ num_events_in_wait_list, event_wait_list, event);
+ } else {
+ return CL_INVALID_PLATFORM;
+ }
+}
+
+cl_int clEnqueueCopyBufferToImage(cl_command_queue command_queue, cl_mem src_buffer,
+ cl_mem dst_image, size_t src_offset, const size_t* dst_origin,
+ const size_t* region, cl_uint num_events_in_wait_list,
+ const cl_event* event_wait_list, cl_event* event) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clEnqueueCopyBufferToImage)lib.getOpenCLFunction("clEnqueueCopyBufferToImage");
+ if (func) {
+ return func(command_queue, src_buffer, dst_image, src_offset, dst_origin, region,
+ num_events_in_wait_list, event_wait_list, event);
+ } else {
+ return CL_INVALID_PLATFORM;
+ }
+}
+
+cl_int clEnqueueNDRangeKernel(cl_command_queue command_queue, cl_kernel kernel, cl_uint work_dim,
+ const size_t* global_work_offset, const size_t* global_work_size,
+ const size_t* local_work_size, cl_uint num_events_in_wait_list,
+ const cl_event* event_wait_list, cl_event* event) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clEnqueueNDRangeKernel)lib.getOpenCLFunction("clEnqueueNDRangeKernel");
+ if (func) {
+ return func(command_queue, kernel, work_dim, global_work_offset, global_work_size,
+ local_work_size, num_events_in_wait_list, event_wait_list, event);
+ } else {
+ return CL_INVALID_PLATFORM;
+ }
+}
+
+cl_command_queue clCreateCommandQueue(cl_context context, cl_device_id device,
+ cl_command_queue_properties properties, cl_int* errcode_ret) {
+ auto& lib = LibOpenCLWrapper::getInstance();
+ auto func = (f_clCreateCommandQueue)lib.getOpenCLFunction("clCreateCommandQueue");
+ if (func) {
+ return func(context, device, properties, errcode_ret);
+ } else {
+ return nullptr;
+ }
+}
diff --git a/tests/cpp-runtime/opencl/opencl_timer_test.cc b/tests/cpp-runtime/opencl/opencl_timer_test.cc
index 6faf2f6a14..f6546c25ac 100644
--- a/tests/cpp-runtime/opencl/opencl_timer_test.cc
+++ b/tests/cpp-runtime/opencl/opencl_timer_test.cc
@@ -44,11 +44,11 @@ TEST(OpenCLTimerNode, nested_timers) {
cl_event ev = clCreateUserEvent(workspace->context, &err);
OPENCL_CHECK_ERROR(err);
cl_mem cl_buf = clCreateBuffer(workspace->context, CL_MEM_READ_ONLY, BUFF_SIZE * sizeof(cl_int),
- NULL, &err);
+ nullptr, &err);
OPENCL_CHECK_ERROR(err);
queue = workspace->GetQueue(thr->device);
OPENCL_CALL(clEnqueueWriteBuffer(queue, cl_buf, false, 0, BUFF_SIZE * sizeof(cl_int), tmp_buf,
- 0, NULL, &ev));
+ 0, nullptr, &ev));
OPENCL_CALL(clReleaseMemObject(cl_buf));
workspace->events[thr->device.device_id].push_back(ev);
nested_timer->Stop();