From 5021cb4d73f931f0d6456c2923429104f4bf56d9 Mon Sep 17 00:00:00 2001 From: Egor Churaev Date: Fri, 11 Nov 2022 12:13:55 +0300 Subject: [PATCH 1/6] [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. --- .gitmodules | 3 + 3rdparty/OpenCL-Headers | 1 + cmake/modules/OpenCL.cmake | 13 +- src/runtime/opencl/opencl_wrapper/README.md | 25 + .../opencl/opencl_wrapper/opencl_wrapper.cc | 551 ++++++++++++++++++ 5 files changed, 591 insertions(+), 2 deletions(-) create mode 160000 3rdparty/OpenCL-Headers create mode 100644 src/runtime/opencl/opencl_wrapper/README.md create mode 100644 src/runtime/opencl/opencl_wrapper/opencl_wrapper.cc 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/cmake/modules/OpenCL.cmake b/cmake/modules/OpenCL.cmake index 430af7e8722c..45560fb2a06f 100644 --- a/cmake/modules/OpenCL.cmake +++ b/cmake/modules/OpenCL.cmake @@ -50,9 +50,18 @@ endif(USE_AOCL) if(USE_OPENCL) if (NOT OpenCL_FOUND) - find_package(OpenCL REQUIRED) + find_package(OpenCL) + endif() + if (OpenCL_FOUND) + message(STATUS "Build with OpenCL support") + else() + message(WARNING "Build with OpenCL wrapper") + add_library(OpenCL STATIC src/runtime/opencl/opencl_wrapper/opencl_wrapper.cc) + set(OpenCL_FOUND true) + set(OpenCL_LIBRARIES OpenCL) + set(OpenCL_INCLUDE_DIRS "3rdparty/OpenCL-Headers") + include_directories(SYSTEM ${OpenCL_INCLUDE_DIRS}) 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}) 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..88458b243297 --- /dev/null +++ b/src/runtime/opencl/opencl_wrapper/opencl_wrapper.cc @@ -0,0 +1,551 @@ +/* + * 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 + */ + +#define CL_TARGET_OPENCL_VERSION 120 +#include +#include + +#define DMLC_USE_LOGGING_LIBRARY +#include +#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 = {"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(); + return dlsym(m_libHandler, funcName); + } + + private: + LibOpenCLWrapper() {} + ~LibOpenCLWrapper() { + if (m_libHandler) dlclose(m_libHandler); + } + void openLibOpenCL() { + for (const auto it : default_so_paths) { + m_libHandler = dlopen(it, RTLD_LAZY); + if (m_libHandler != NULL) return; + } + ICHECK(m_libHandler != NULL) << "Error! Cannot open libOpenCL!"; + } + + private: + void* m_libHandler = nullptr; +}; + +// 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 NULL; + } +} + +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 NULL; + } +} + +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 NULL; + } +} + +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 NULL; + } +} + +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 NULL; + } +} + +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 NULL; + } +} + +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 NULL; + } +} + +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 NULL; + } +} From fe80ad56d3abcca6ab73e0ef5719eb3b86a9c37e Mon Sep 17 00:00:00 2001 From: Egor Churaev Date: Mon, 14 Nov 2022 11:36:12 +0300 Subject: [PATCH 2/6] Update apps and documentation --- .../app/src/main/jni/Android.mk | 1 + .../app/src/main/jni/Application.mk | 6 ++-- .../app/src/main/jni/make/config.mk | 3 -- .../app/src/main/jni/tvm_runtime.h | 1 + apps/android_deploy/README.md | 32 +++---------------- .../app/src/main/jni/Android.mk | 3 +- .../app/src/main/jni/Application.mk | 6 ++-- .../app/src/main/jni/make/config.mk | 3 -- .../app/src/main/jni/tvm_runtime.h | 1 + apps/android_rpc/README.md | 28 +--------------- apps/android_rpc/app/src/main/jni/Android.mk | 3 +- .../app/src/main/jni/Application.mk | 6 ++-- .../app/src/main/jni/make/config.mk | 3 -- .../app/src/main/jni/tvm_runtime.h | 1 + apps/cpp_rpc/README.md | 8 ++++- .../deploy_models/deploy_model_on_android.py | 11 +++---- 16 files changed, 34 insertions(+), 82 deletions(-) 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/Application.mk b/apps/android_camera/app/src/main/jni/Application.mk index 83b7b4417b9e..047835183bad 100644 --- a/apps/android_camera/app/src/main/jni/Application.mk +++ b/apps/android_camera/app/src/main/jni/Application.mk @@ -32,9 +32,9 @@ APP_ABI ?= all APP_STL := c++_shared APP_CPPFLAGS += -DTVM_LOG_STACK_TRACE=0 -DTVM4J_ANDROID=1 -std=c++17 -Oz -frtti -ifeq ($(USE_OPENCL), 1) - APP_CPPFLAGS += -DTVM_OPENCL_RUNTIME=1 -endif + +# OpenCL support +APP_CPPFLAGS += -DTVM_OPENCL_RUNTIME=1 ifeq ($(USE_VULKAN), 1) APP_CPPFLAGS += -DTVM_VULKAN_RUNTIME=1 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..cf13e11c311c 100644 --- a/apps/android_camera/app/src/main/jni/make/config.mk +++ b/apps/android_camera/app/src/main/jni/make/config.mk @@ -33,9 +33,6 @@ APP_ABI = all APP_PLATFORM = android-24 -# whether enable OpenCL during compile -USE_OPENCL = 0 - # 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..6d09544fbb80 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,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/source_utils.cc" #endif diff --git a/apps/android_deploy/README.md b/apps/android_deploy/README.md index 32e601840f04..873c8ac5a630 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,9 @@ 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 open. ## 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/Application.mk b/apps/android_deploy/app/src/main/jni/Application.mk index 4a83907ff329..bc8562c3849f 100644 --- a/apps/android_deploy/app/src/main/jni/Application.mk +++ b/apps/android_deploy/app/src/main/jni/Application.mk @@ -28,6 +28,6 @@ include $(config) APP_STL := c++_static APP_CPPFLAGS += -DTVM_LOG_STACK_TRACE=0 -DTVM4J_ANDROID=1 -std=c++17 -Oz -frtti -ifeq ($(USE_OPENCL), 1) - APP_CPPFLAGS += -DTVM_OPENCL_RUNTIME=1 -endif + +# OpenCL support +APP_CPPFLAGS += -DTVM_OPENCL_RUNTIME=1 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..309aa581b6a8 100644 --- a/apps/android_deploy/app/src/main/jni/make/config.mk +++ b/apps/android_deploy/app/src/main/jni/make/config.mk @@ -33,9 +33,6 @@ APP_ABI = all APP_PLATFORM = android-17 -# whether enable OpenCL during compile -USE_OPENCL = 0 - # 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..f628f163cd34 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,5 @@ #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" #endif diff --git a/apps/android_rpc/README.md b/apps/android_rpc/README.md index 2e301af6d996..60ba5d2cf71f 100644 --- a/apps/android_rpc/README.md +++ b/apps/android_rpc/README.md @@ -74,33 +74,7 @@ $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 open. ## 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/Application.mk b/apps/android_rpc/app/src/main/jni/Application.mk index df560863f091..cebbcffcf6dd 100644 --- a/apps/android_rpc/app/src/main/jni/Application.mk +++ b/apps/android_rpc/app/src/main/jni/Application.mk @@ -32,9 +32,9 @@ APP_ABI ?= armeabi-v7a arm64-v8a x86 x86_64 mips APP_STL := c++_shared APP_CPPFLAGS += -DTVM_LOG_STACK_TRACE=0 -DTVM4J_ANDROID=1 -std=c++17 -Oz -frtti -ifeq ($(USE_OPENCL), 1) - APP_CPPFLAGS += -DTVM_OPENCL_RUNTIME=1 -endif + +# OpenCL support +APP_CPPFLAGS += -DTVM_OPENCL_RUNTIME=1 ifeq ($(USE_VULKAN), 1) APP_CPPFLAGS += -DTVM_VULKAN_RUNTIME=1 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..cd7a261120c1 100644 --- a/apps/android_rpc/app/src/main/jni/make/config.mk +++ b/apps/android_rpc/app/src/main/jni/make/config.mk @@ -33,9 +33,6 @@ APP_ABI = all APP_PLATFORM = android-24 -# whether enable OpenCL during compile -USE_OPENCL = 0 - # 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..f35523ad119e 100644 --- a/apps/cpp_rpc/README.md +++ b/apps/cpp_rpc/README.md @@ -37,7 +37,13 @@ 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 open. + + 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/gallery/how_to/deploy_models/deploy_model_on_android.py b/gallery/how_to/deploy_models/deploy_model_on_android.py index 10e108239ee7..6421d8477b74 100644 --- a/gallery/how_to/deploy_models/deploy_model_on_android.py +++ b/gallery/how_to/deploy_models/deploy_model_on_android.py @@ -115,7 +115,8 @@ # Follow this `readme page `_ to # install TVM RPC APK on the android device. # -# Here is an example of config.mk. I enabled OpenCL and Vulkan. +# Here is an example of config.mk. I enabled Vulkan. Dynamic OpenCL support is +# enabled by default. # # # .. code-block:: bash @@ -124,9 +125,6 @@ # # APP_PLATFORM = android-24 # -# # whether enable OpenCL during compile -# USE_OPENCL = 1 -# # # whether to enable Vulkan during compile # USE_VULKAN = 1 # @@ -137,11 +135,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:: # From fd5480b923c125f9c7b8d23cc206ada38db7b24a Mon Sep 17 00:00:00 2001 From: Egor Churaev Date: Mon, 14 Nov 2022 16:45:12 +0300 Subject: [PATCH 3/6] Apply comments --- .../app/src/main/jni/Application.mk | 6 +++--- .../app/src/main/jni/make/config.mk | 3 +++ apps/android_deploy/README.md | 4 +++- .../app/src/main/jni/Application.mk | 6 +++--- .../app/src/main/jni/make/config.mk | 3 +++ apps/android_rpc/README.md | 6 +++++- .../app/src/main/jni/Application.mk | 6 +++--- .../app/src/main/jni/make/config.mk | 3 +++ apps/cpp_rpc/README.md | 4 +++- cmake/modules/OpenCL.cmake | 16 ++++++--------- .../deploy_models/deploy_model_on_android.py | 6 ++++-- src/runtime/opencl/opencl_module.cc | 2 +- .../opencl/opencl_wrapper/opencl_wrapper.cc | 20 +++++++++---------- tests/cpp-runtime/opencl/opencl_timer_test.cc | 4 ++-- 14 files changed, 52 insertions(+), 37 deletions(-) diff --git a/apps/android_camera/app/src/main/jni/Application.mk b/apps/android_camera/app/src/main/jni/Application.mk index 047835183bad..83b7b4417b9e 100644 --- a/apps/android_camera/app/src/main/jni/Application.mk +++ b/apps/android_camera/app/src/main/jni/Application.mk @@ -32,9 +32,9 @@ APP_ABI ?= all APP_STL := c++_shared APP_CPPFLAGS += -DTVM_LOG_STACK_TRACE=0 -DTVM4J_ANDROID=1 -std=c++17 -Oz -frtti - -# OpenCL support -APP_CPPFLAGS += -DTVM_OPENCL_RUNTIME=1 +ifeq ($(USE_OPENCL), 1) + APP_CPPFLAGS += -DTVM_OPENCL_RUNTIME=1 +endif ifeq ($(USE_VULKAN), 1) APP_CPPFLAGS += -DTVM_VULKAN_RUNTIME=1 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 cf13e11c311c..1f601b9afb29 100644 --- a/apps/android_camera/app/src/main/jni/make/config.mk +++ b/apps/android_camera/app/src/main/jni/make/config.mk @@ -33,6 +33,9 @@ APP_ABI = all APP_PLATFORM = android-24 +# whether enable OpenCL during compile +USE_OPENCL = 1 + # whether to enable Vulkan during compile USE_VULKAN = 0 diff --git a/apps/android_deploy/README.md b/apps/android_deploy/README.md index 873c8ac5a630..4cfd9eb9daf2 100644 --- a/apps/android_deploy/README.md +++ b/apps/android_deploy/README.md @@ -81,7 +81,9 @@ Upload `tvmdemo-release.apk` to your Android device and install it. 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 open. +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/Application.mk b/apps/android_deploy/app/src/main/jni/Application.mk index bc8562c3849f..4a83907ff329 100644 --- a/apps/android_deploy/app/src/main/jni/Application.mk +++ b/apps/android_deploy/app/src/main/jni/Application.mk @@ -28,6 +28,6 @@ include $(config) APP_STL := c++_static APP_CPPFLAGS += -DTVM_LOG_STACK_TRACE=0 -DTVM4J_ANDROID=1 -std=c++17 -Oz -frtti - -# OpenCL support -APP_CPPFLAGS += -DTVM_OPENCL_RUNTIME=1 +ifeq ($(USE_OPENCL), 1) + APP_CPPFLAGS += -DTVM_OPENCL_RUNTIME=1 +endif 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 309aa581b6a8..b06f42b2647a 100644 --- a/apps/android_deploy/app/src/main/jni/make/config.mk +++ b/apps/android_deploy/app/src/main/jni/make/config.mk @@ -33,6 +33,9 @@ APP_ABI = all 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 = diff --git a/apps/android_rpc/README.md b/apps/android_rpc/README.md index 60ba5d2cf71f..d0a11b6121dc 100644 --- a/apps/android_rpc/README.md +++ b/apps/android_rpc/README.md @@ -74,7 +74,11 @@ $ANDROID_HOME/platform-tools/adb uninstall org.apache.tvm.tvmrpc ### Build with OpenCL -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 open. +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/Application.mk b/apps/android_rpc/app/src/main/jni/Application.mk index cebbcffcf6dd..df560863f091 100644 --- a/apps/android_rpc/app/src/main/jni/Application.mk +++ b/apps/android_rpc/app/src/main/jni/Application.mk @@ -32,9 +32,9 @@ APP_ABI ?= armeabi-v7a arm64-v8a x86 x86_64 mips APP_STL := c++_shared APP_CPPFLAGS += -DTVM_LOG_STACK_TRACE=0 -DTVM4J_ANDROID=1 -std=c++17 -Oz -frtti - -# OpenCL support -APP_CPPFLAGS += -DTVM_OPENCL_RUNTIME=1 +ifeq ($(USE_OPENCL), 1) + APP_CPPFLAGS += -DTVM_OPENCL_RUNTIME=1 +endif ifeq ($(USE_VULKAN), 1) APP_CPPFLAGS += -DTVM_VULKAN_RUNTIME=1 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 cd7a261120c1..855a0af19021 100644 --- a/apps/android_rpc/app/src/main/jni/make/config.mk +++ b/apps/android_rpc/app/src/main/jni/make/config.mk @@ -33,6 +33,9 @@ APP_ABI = all APP_PLATFORM = android-24 +# whether enable OpenCL during compile +USE_OPENCL = 1 + # whether to enable Vulkan during compile USE_VULKAN = 0 diff --git a/apps/cpp_rpc/README.md b/apps/cpp_rpc/README.md index f35523ad119e..58eb68055f4d 100644 --- a/apps/cpp_rpc/README.md +++ b/apps/cpp_rpc/README.md @@ -41,7 +41,9 @@ This folder contains a simple recipe to make RPC server in c++. ``` 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 open. + 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: ``` diff --git a/cmake/modules/OpenCL.cmake b/cmake/modules/OpenCL.cmake index 45560fb2a06f..831d99b44378 100644 --- a/cmake/modules/OpenCL.cmake +++ b/cmake/modules/OpenCL.cmake @@ -49,21 +49,17 @@ else() endif(USE_AOCL) if(USE_OPENCL) - if (NOT OpenCL_FOUND) - find_package(OpenCL) - endif() + tvm_file_glob(GLOB RUNTIME_OPENCL_SRCS src/runtime/opencl/*.cc) if (OpenCL_FOUND) message(STATUS "Build with OpenCL support") + list(APPEND TVM_RUNTIME_LINKER_LIBS ${OpenCL_LIBRARIES}) else() message(WARNING "Build with OpenCL wrapper") - add_library(OpenCL STATIC src/runtime/opencl/opencl_wrapper/opencl_wrapper.cc) - set(OpenCL_FOUND true) - set(OpenCL_LIBRARIES OpenCL) - set(OpenCL_INCLUDE_DIRS "3rdparty/OpenCL-Headers") - include_directories(SYSTEM ${OpenCL_INCLUDE_DIRS}) + file_glob_append(RUNTIME_OPENCL_SRCS + "src/runtime/opencl/opencl_wrapper/opencl_wrapper.cc" + ) + include_directories(SYSTEM "3rdparty/OpenCL-Headers") endif() - tvm_file_glob(GLOB RUNTIME_OPENCL_SRCS src/runtime/opencl/*.cc) - list(APPEND TVM_RUNTIME_LINKER_LIBS ${OpenCL_LIBRARIES}) if(DEFINED USE_OPENCL_GTEST AND EXISTS ${USE_OPENCL_GTEST}) file_glob_append(RUNTIME_OPENCL_SRCS 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 6421d8477b74..4bf86e2981a1 100644 --- a/gallery/how_to/deploy_models/deploy_model_on_android.py +++ b/gallery/how_to/deploy_models/deploy_model_on_android.py @@ -115,8 +115,7 @@ # Follow this `readme page `_ to # install TVM RPC APK on the android device. # -# Here is an example of config.mk. I enabled Vulkan. Dynamic OpenCL support is -# enabled by default. +# Here is an example of config.mk. I enabled OpenCL and Vulkan. # # # .. code-block:: bash @@ -125,6 +124,9 @@ # # APP_PLATFORM = android-24 # +# # whether enable OpenCL during compile +# USE_OPENCL = 1 +# # # whether to enable Vulkan during compile # USE_VULKAN = 1 # 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/opencl_wrapper.cc b/src/runtime/opencl/opencl_wrapper/opencl_wrapper.cc index 88458b243297..9d8db498ae6a 100644 --- a/src/runtime/opencl/opencl_wrapper/opencl_wrapper.cc +++ b/src/runtime/opencl/opencl_wrapper/opencl_wrapper.cc @@ -80,9 +80,9 @@ class LibOpenCLWrapper { void openLibOpenCL() { for (const auto it : default_so_paths) { m_libHandler = dlopen(it, RTLD_LAZY); - if (m_libHandler != NULL) return; + if (m_libHandler != nullptr) return; } - ICHECK(m_libHandler != NULL) << "Error! Cannot open libOpenCL!"; + ICHECK(m_libHandler != nullptr) << "Error! Cannot open libOpenCL!"; } private: @@ -204,7 +204,7 @@ cl_context clCreateContext(const cl_context_properties* properties, cl_uint num_ if (func) { return func(properties, num_devices, devices, pfn_notify, user_data, errcode_ret); } else { - return NULL; + return nullptr; } } @@ -248,7 +248,7 @@ cl_mem clCreateBuffer(cl_context context, cl_mem_flags flags, size_t size, void* if (func) { return func(context, flags, size, host_ptr, errcode_ret); } else { - return NULL; + return nullptr; } } @@ -259,7 +259,7 @@ cl_mem clCreateImage(cl_context context, cl_mem_flags flags, const cl_image_form if (func) { return func(context, flags, image_format, image_desc, host_ptr, errcode_ret); } else { - return NULL; + return nullptr; } } @@ -280,7 +280,7 @@ cl_program clCreateProgramWithSource(cl_context context, cl_uint count, const ch if (func) { return func(context, count, strings, lengths, errcode_ret); } else { - return NULL; + return nullptr; } } @@ -293,7 +293,7 @@ cl_program clCreateProgramWithBinary(cl_context context, cl_uint num_devices, if (func) { return func(context, num_devices, device_list, lengths, binaries, binary_status, errcode_ret); } else { - return NULL; + return nullptr; } } @@ -337,7 +337,7 @@ cl_kernel clCreateKernel(cl_program program, const char* kernel_name, cl_int* er if (func) { return func(program, kernel_name, errcode_ret); } else { - return NULL; + return nullptr; } } @@ -377,7 +377,7 @@ cl_event clCreateUserEvent(cl_context context, cl_int* errcode_ret) { if (func) { return func(context, errcode_ret); } else { - return NULL; + return nullptr; } } @@ -546,6 +546,6 @@ cl_command_queue clCreateCommandQueue(cl_context context, cl_device_id device, if (func) { return func(context, device, properties, errcode_ret); } else { - return NULL; + 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(); From 240ab5f27ba8873e71419eedb0dfca991e430363 Mon Sep 17 00:00:00 2001 From: Egor Churaev Date: Tue, 15 Nov 2022 11:23:43 +0300 Subject: [PATCH 4/6] Apply comments and fix Android build Also, use OpenCL wrapper by default and fix Windows build --- .../app/src/main/jni/tvm_runtime.h | 1 + .../app/src/main/jni/tvm_runtime.h | 2 ++ cmake/config.cmake | 2 +- cmake/modules/OpenCL.cmake | 23 +++++++--------- cmake/utils/FindOpenCL.cmake | 2 +- .../opencl/opencl_wrapper/opencl_wrapper.cc | 27 +++++++++++++++++-- 6 files changed, 40 insertions(+), 17 deletions(-) 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 6d09544fbb80..0aac7f170ab4 100644 --- a/apps/android_camera/app/src/main/jni/tvm_runtime.h +++ b/apps/android_camera/app/src/main/jni/tvm_runtime.h @@ -63,6 +63,7 @@ #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/app/src/main/jni/tvm_runtime.h b/apps/android_deploy/app/src/main/jni/tvm_runtime.h index f628f163cd34..a2f10701d6df 100644 --- a/apps/android_deploy/app/src/main/jni/tvm_runtime.h +++ b/apps/android_deploy/app/src/main/jni/tvm_runtime.h @@ -48,4 +48,6 @@ #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/cmake/config.cmake b/cmake/config.cmake index 22a548d29895..89bd2fc9547b 100644 --- a/cmake/config.cmake +++ b/cmake/config.cmake @@ -65,7 +65,7 @@ set(USE_AOCL OFF) # Whether enable OpenCL runtime # # Possible values: -# - ON: enable OpenCL with cmake's auto search +# - ON: enable OpenCL with OpenCL wrapper # - 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 831d99b44378..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) @@ -50,15 +41,21 @@ endif(USE_AOCL) if(USE_OPENCL) tvm_file_glob(GLOB RUNTIME_OPENCL_SRCS src/runtime/opencl/*.cc) - if (OpenCL_FOUND) - message(STATUS "Build with OpenCL support") - list(APPEND TVM_RUNTIME_LINKER_LIBS ${OpenCL_LIBRARIES}) - else() + + 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}) 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/src/runtime/opencl/opencl_wrapper/opencl_wrapper.cc b/src/runtime/opencl/opencl_wrapper/opencl_wrapper.cc index 9d8db498ae6a..c447ebcb5339 100644 --- a/src/runtime/opencl/opencl_wrapper/opencl_wrapper.cc +++ b/src/runtime/opencl/opencl_wrapper/opencl_wrapper.cc @@ -19,14 +19,21 @@ /*! * \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 -#define DMLC_USE_LOGGING_LIBRARY +#if defined(_WIN32) +#include +#else #include +#endif + +#define DMLC_USE_LOGGING_LIBRARY #include #include @@ -49,7 +56,7 @@ static const std::vector default_so_paths = { "/system/vendor/lib/libPVROCL.so", "/data/data/org.pocl.libs/files/lib/libpocl.so"}; #elif defined(_WIN32) -static const std::vector default_so_paths = {"OpenCL.dll"}; +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", @@ -69,24 +76,40 @@ class LibOpenCLWrapper { 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 From 04532862d774c1126799308db363cfa9dda3cdf0 Mon Sep 17 00:00:00 2001 From: Egor Churaev Date: Tue, 15 Nov 2022 20:04:44 +0300 Subject: [PATCH 5/6] Apply comments --- cmake/config.cmake | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/cmake/config.cmake b/cmake/config.cmake index 89bd2fc9547b..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 OpenCL wrapper +# - 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) From 46102e05453dcbd06d46061ecd0d60c72f6a1fe2 Mon Sep 17 00:00:00 2001 From: Egor Churaev Date: Tue, 15 Nov 2022 22:07:07 +0300 Subject: [PATCH 6/6] Update LICENSE file --- LICENSE | 1 + 1 file changed, 1 insertion(+) 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