diff --git a/.gitmodules b/.gitmodules index e03336443d73..66fd0390cf35 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 000000000000..b590a6bfe034 --- /dev/null +++ b/3rdparty/OpenCL-Headers @@ -0,0 +1 @@ +Subproject commit b590a6bfe034ea3a418b7b523e3490956bcb367a diff --git a/LICENSE b/LICENSE index 345026985b07..6524d530deca 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 513666a4ecb4..2201f669653c 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 49e332665ad9..1f601b9afb29 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 658534780130..0aac7f170ab4 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 32e601840f04..4cfd9eb9daf2 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 installation instruction](https://gradle.org/install) for your operating system. -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 1b06a6bdb898..ad9cee9bbdb5 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 bcd56e37896d..b06f42b2647a 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 725b5e1d3b7a..a2f10701d6df 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 2e301af6d996..d0a11b6121dc 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 1b06a6bdb898..ad9cee9bbdb5 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 851430cd42a9..855a0af19021 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 543c9c85334e..17a20bbaf9a0 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 d073fca81921..58eb68055f4d 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 22a548d29895..679f5c459e87 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 430af7e8722c..e738df7c564c 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 f2931332fc90..8eb35ab3993e 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 10e108239ee7..4bf86e2981a1 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 @@ # # # 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 9ae80d59d565..2fb157aac6af 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 000000000000..7597a442c1a9 --- /dev/null +++ b/src/runtime/opencl/opencl_wrapper/README.md @@ -0,0 +1,25 @@ + + + + + + + + + + + + + + + + + +# 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 000000000000..c447ebcb5339 --- /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 +#include + +#if defined(_WIN32) +#include +#else +#include +#endif + +#define DMLC_USE_LOGGING_LIBRARY +#include + +#include + +namespace { +#if defined(__APPLE__) || defined(__MACOSX) +static const std::vector default_so_paths = { + "libOpenCL.so", "/System/Library/Frameworks/OpenCL.framework/OpenCL"}; +#elif defined(__ANDROID__) +static const std::vector 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 default_so_paths = {__TEXT("OpenCL.dll")}; +#elif defined(__linux__) +static const std::vector 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 6faf2f6a1482..f6546c25aca5 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();