diff --git a/WORKSPACE b/WORKSPACE index 225cdd885cf16866c018ce018fe9732b24673e75..644692d297d79ad7eb4841c426f440a65bf89ef2 100644 --- a/WORKSPACE +++ b/WORKSPACE @@ -1,22 +1,22 @@ workspace(name = "mace") +# proto_library rules implicitly depend on @com_google_protobuf//:protoc, +# which is the proto-compiler. +# This statement defines the @com_google_protobuf repo. http_archive( - name = "org_tensorflow", - urls = ["http://v9.git.n.xiaomi.com/deep-learning/tensorflow/repository/archive.zip?ref=v1.3.0"], - strip_prefix = "tensorflow-v1.3.0-9e76bf324f6bac63137a02bb6e6ec9120703ea9b", - sha256 = "97049d3a59a77858e12c55422bd129261b14e869a91aebcdcc39439393c00dc7", + name = "com_google_protobuf", + urls = ["http://v9.git.n.xiaomi.com/deep-learning/protobuf/repository/archive.zip?ref=c7457ef65a7a8584b1e3bd396c401ccf8e275ffa"], + strip_prefix = "protobuf-c7457ef65a7a8584b1e3bd396c401ccf8e275ffa-c7457ef65a7a8584b1e3bd396c401ccf8e275ffa", + sha256 = "0a54cae83b77f4b54b7db4eaebadd81fbe91655e84a1ef3f6d29116d75f3a45f", ) -# TensorFlow depends on "io_bazel_rules_closure" so we need this here. -# Needs to be kept in sync with the same target in TensorFlow's WORKSPACE file. +# cc_proto_library rules implicitly depend on @com_google_protobuf_cc//:cc_toolchain, +# which is the C++ proto runtime (base classes and common utilities). http_archive( - name = "io_bazel_rules_closure", - sha256 = "60fc6977908f999b23ca65698c2bb70213403824a84f7904310b6000d78be9ce", - strip_prefix = "rules_closure-5ca1dab6df9ad02050f7ba4e816407f88690cf7d", - urls = [ - "http://bazel-mirror.storage.googleapis.com/github.com/bazelbuild/rules_closure/archive/5ca1dab6df9ad02050f7ba4e816407f88690cf7d.tar.gz", # 2017-02-03 - "https://github.com/bazelbuild/rules_closure/archive/5ca1dab6df9ad02050f7ba4e816407f88690cf7d.tar.gz", - ], + name = "com_google_protobuf_cc", + urls = ["http://v9.git.n.xiaomi.com/deep-learning/protobuf/repository/archive.zip?ref=c7457ef65a7a8584b1e3bd396c401ccf8e275ffa"], + strip_prefix = "protobuf-c7457ef65a7a8584b1e3bd396c401ccf8e275ffa-c7457ef65a7a8584b1e3bd396c401ccf8e275ffa", + sha256 = "0a54cae83b77f4b54b7db4eaebadd81fbe91655e84a1ef3f6d29116d75f3a45f", ) new_http_archive( @@ -27,16 +27,19 @@ new_http_archive( build_file = "mace/third_party/gtest.BUILD", ) -# Import all of the tensorflow dependencies. -load('@org_tensorflow//tensorflow:workspace.bzl', 'tf_workspace') -tf_workspace(tf_repo_name = "org_tensorflow") - new_http_archive( - name = "ncnn", - urls = ["http://v9.git.n.xiaomi.com/deep-learning/ncnn/repository/archive.zip?ref=bazel-fix"], - strip_prefix = "ncnn-bazel-fix-ce5e416164545e1ab37fe3544502624f605ca234/src", - sha256 = "e6d76356179bcdbb988279f0b42ab050c8af55970e1ad767787ad21d5b7aad51", - build_file = "mace/third_party/ncnn.BUILD", + name = "six_archive", + urls = [ + "http://mirror.bazel.build/pypi.python.org/packages/source/s/six/six-1.10.0.tar.gz", + "https://pypi.python.org/packages/source/s/six/six-1.10.0.tar.gz", + ], + sha256 = "105f8d68616f8248e24bf0e9372ef04d3cc10104f1980f54d57b2ce73a5ad56a", + strip_prefix = "six-1.10.0", + build_file = "mace/third_party/six.BUILD", +) +bind( + name = "six", + actual = "@six_archive//:six", ) new_http_archive( @@ -50,6 +53,6 @@ new_http_archive( # Set up Android NDK android_ndk_repository( name = "androidndk", - # Android 4.0 - api_level = 14 + # Android 5.0 + api_level = 21 ) diff --git a/mace/core/BUILD b/mace/core/BUILD index 73720cbac9e7dd7679f4f0b63b6ce5f73f0529aa..b6a087cad42aa8c290a5163d003d3d0770d978c0 100644 --- a/mace/core/BUILD +++ b/mace/core/BUILD @@ -12,11 +12,15 @@ load("//mace:mace.bzl", "if_android") cc_library( name = "opencl_runtime", srcs = glob([ - "platform/opencl/*.hpp", + "platform/opencl/cl.hpp", + "platform/opencl/cl2.hpp", + "platform/opencl/opencl_wrapper.h", + "platform/opencl/opencl_wrapper.cc", ]), copts = ["-std=c++11"], deps = [ - "@opencl_headers//:opencl20_headers", + "@opencl_headers//:opencl12_headers", + "core", ], ) diff --git a/mace/core/platform/opencl/opencl_wrapper.cc b/mace/core/platform/opencl/opencl_wrapper.cc new file mode 100644 index 0000000000000000000000000000000000000000..79b319b07649e3030139c74656cf9fc605a66fda --- /dev/null +++ b/mace/core/platform/opencl/opencl_wrapper.cc @@ -0,0 +1,615 @@ +// +// Copyright (c) 2017 XiaoMi All rights reserved. +// + +#include "CL/opencl.h" + +#include "mace/core/logging.h" +#include "mace/core/platform/opencl/opencl_wrapper.h" + +#include +#include + +/** + * Wrapper of OpenCL 1.2 + */ +namespace mace { +class OpenCLStub final { + public: + static OpenCLStub &Get(); + bool loaded() { return loaded_; } + + using clBuildProgramFunc = cl_int (*)(cl_program, + cl_uint, + const cl_device_id *, + const char *, + void (*pfn_notify)(cl_program, void *), + void *); + using clEnqueueNDRangeKernelFunc = 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 clSetKernelArgFunc = cl_int (*)(cl_kernel, + cl_uint, + size_t, + const void *); + using clRetainMemObjectFunc = cl_int (*)(cl_mem); + using clReleaseMemObjectFunc = cl_int (*)(cl_mem); + using clEnqueueUnmapMemObjectFunc = cl_int (*)( + cl_command_queue, cl_mem, void *, cl_uint, const cl_event *, cl_event *); + using clRetainCommandQueueFunc = cl_int (*)(cl_command_queue command_queue); + using clReleaseContextFunc = cl_int (*)(cl_context); + using clReleaseEventFunc = cl_int (*)(cl_event); + using clEnqueueWriteBufferFunc = cl_int (*)(cl_command_queue, + cl_mem, + cl_bool, + size_t, + size_t, + const void *, + cl_uint, + const cl_event *, + cl_event *); + using clEnqueueReadBufferFunc = cl_int (*)(cl_command_queue, + cl_mem, + cl_bool, + size_t, + size_t, + void *, + cl_uint, + const cl_event *, + cl_event *); + using clGetProgramBuildInfoFunc = cl_int (*)(cl_program, + cl_device_id, + cl_program_build_info, + size_t, + void *, + size_t *); + using clRetainProgramFunc = cl_int (*)(cl_program program); + using clEnqueueMapBufferFunc = void *(*)(cl_command_queue, + cl_mem, + cl_bool, + cl_map_flags, + size_t, + size_t, + cl_uint, + const cl_event *, + cl_event *, + cl_int *); + using clReleaseCommandQueueFunc = cl_int (*)(cl_command_queue); + using clCreateProgramWithBinaryFunc = cl_program (*)(cl_context, + cl_uint, + const cl_device_id *, + const size_t *, + const unsigned char **, + cl_int *, + cl_int *); + using clRetainContextFunc = cl_int (*)(cl_context context); + using clReleaseProgramFunc = cl_int (*)(cl_program program); + using clFlushFunc = cl_int (*)(cl_command_queue command_queue); + using clFinishFunc = cl_int (*)(cl_command_queue command_queue); + using clGetProgramInfoFunc = + cl_int (*)(cl_program, cl_program_info, size_t, void *, size_t *); + using clCreateKernelFunc = cl_kernel (*)(cl_program, const char *, cl_int *); + using clRetainKernelFunc = cl_int (*)(cl_kernel kernel); + using clCreateBufferFunc = + cl_mem (*)(cl_context, cl_mem_flags, size_t, void *, cl_int *); + using clCreateProgramWithSourceFunc = cl_program (*)( + cl_context, cl_uint, const char **, const size_t *, cl_int *); + using clReleaseKernelFunc = cl_int (*)(cl_kernel kernel); + using clGetDeviceInfoFunc = + cl_int (*)(cl_device_id, cl_device_info, size_t, void *, size_t *); + using clGetDeviceIDsFunc = cl_int (*)( + cl_platform_id, cl_device_type, cl_uint, cl_device_id *, cl_uint *); + using clRetainEventFunc = cl_int (*)(cl_event); + +#define DEFINE_FUNC_PTR(func) func##Func func = nullptr + + DEFINE_FUNC_PTR(clBuildProgram); + DEFINE_FUNC_PTR(clEnqueueNDRangeKernel); + DEFINE_FUNC_PTR(clSetKernelArg); + DEFINE_FUNC_PTR(clReleaseKernel); + DEFINE_FUNC_PTR(clCreateProgramWithSource); + DEFINE_FUNC_PTR(clCreateBuffer); + DEFINE_FUNC_PTR(clRetainKernel); + DEFINE_FUNC_PTR(clCreateKernel); + DEFINE_FUNC_PTR(clGetProgramInfo); + DEFINE_FUNC_PTR(clFlush); + DEFINE_FUNC_PTR(clFinish); + DEFINE_FUNC_PTR(clReleaseProgram); + DEFINE_FUNC_PTR(clRetainContext); + DEFINE_FUNC_PTR(clCreateProgramWithBinary); + DEFINE_FUNC_PTR(clReleaseCommandQueue); + DEFINE_FUNC_PTR(clEnqueueMapBuffer); + DEFINE_FUNC_PTR(clRetainProgram); + DEFINE_FUNC_PTR(clGetProgramBuildInfo); + DEFINE_FUNC_PTR(clEnqueueReadBuffer); + DEFINE_FUNC_PTR(clEnqueueWriteBuffer); + DEFINE_FUNC_PTR(clReleaseEvent); + DEFINE_FUNC_PTR(clReleaseContext); + DEFINE_FUNC_PTR(clRetainCommandQueue); + DEFINE_FUNC_PTR(clEnqueueUnmapMemObject); + DEFINE_FUNC_PTR(clRetainMemObject); + DEFINE_FUNC_PTR(clReleaseMemObject); + DEFINE_FUNC_PTR(clGetDeviceInfo); + DEFINE_FUNC_PTR(clGetDeviceIDs); + DEFINE_FUNC_PTR(clRetainEvent); + +#undef DEFINE_FUNC_PTR + + private: + bool TryLoadAll(); + bool Load(const std::string &library); + bool loaded_ = false; +}; + +OpenCLStub &OpenCLStub::Get() { + static std::once_flag load_once; + static OpenCLStub instance; + std::call_once(load_once, []() { instance.TryLoadAll(); }); + return instance; +} + +bool OpenCLStub::TryLoadAll() { + // TODO (heliangliang) Make this configurable + static const std::vector pathes = { +#if defined(__aarch64__) + // Qualcomm Adreno + "/system/vendor/lib64/libOpenCL.so", + "/system/lib64/libOpenCL.so", + // Mali + "/system/vendor/lib64/egl/libGLES_mali.so", + "/system/lib64/egl/libGLES_mali.so", +#else + // Qualcomm Adreno + "/system/vendor/lib/libOpenCL.so", + "/system/lib/libOpenCL.so", + // Mali + "/system/vendor/lib/egl/libGLES_mali.so", + "/system/lib/egl/libGLES_mali.so", +#endif + }; + + for (const auto &path : pathes) { + VLOG(2) << "Loading OpenCL from " << path; + if (Load(path)) { + return true; + } + } + + LOG(ERROR) << "Failed to load OpenCL library"; + return false; +} + +bool OpenCLStub::Load(const std::string &path) { + void *handle = dlopen(path.c_str(), RTLD_LAZY | RTLD_LOCAL); + + if (handle == nullptr) { + VLOG(2) << "Failed to load OpenCL library from path " << path + << " error code: " << dlerror(); + return false; + } + +#define ASSIGN_FROM_DLSYM(func) \ + do { \ + void *ptr = dlsym(handle, #func); \ + if (ptr == nullptr) { \ + LOG(ERROR) << "Failed to load " << #func << " from " << path; \ + loaded_ = false; \ + dlclose(handle); \ + return false; \ + } \ + func = reinterpret_cast(ptr); \ + VLOG(2) << "Loaded " << #func << " from " << path; \ + } while (false) + + ASSIGN_FROM_DLSYM(clBuildProgram); + ASSIGN_FROM_DLSYM(clEnqueueNDRangeKernel); + ASSIGN_FROM_DLSYM(clSetKernelArg); + ASSIGN_FROM_DLSYM(clReleaseKernel); + ASSIGN_FROM_DLSYM(clCreateProgramWithSource); + ASSIGN_FROM_DLSYM(clCreateBuffer); + ASSIGN_FROM_DLSYM(clRetainKernel); + ASSIGN_FROM_DLSYM(clCreateKernel); + ASSIGN_FROM_DLSYM(clGetProgramInfo); + ASSIGN_FROM_DLSYM(clFlush); + ASSIGN_FROM_DLSYM(clFinish); + ASSIGN_FROM_DLSYM(clReleaseProgram); + ASSIGN_FROM_DLSYM(clRetainContext); + ASSIGN_FROM_DLSYM(clCreateProgramWithBinary); + ASSIGN_FROM_DLSYM(clReleaseCommandQueue); + ASSIGN_FROM_DLSYM(clEnqueueMapBuffer); + ASSIGN_FROM_DLSYM(clRetainProgram); + ASSIGN_FROM_DLSYM(clGetProgramBuildInfo); + ASSIGN_FROM_DLSYM(clEnqueueReadBuffer); + ASSIGN_FROM_DLSYM(clEnqueueWriteBuffer); + ASSIGN_FROM_DLSYM(clReleaseEvent); + ASSIGN_FROM_DLSYM(clReleaseContext); + ASSIGN_FROM_DLSYM(clRetainCommandQueue); + ASSIGN_FROM_DLSYM(clEnqueueUnmapMemObject); + ASSIGN_FROM_DLSYM(clRetainMemObject); + ASSIGN_FROM_DLSYM(clReleaseMemObject); + ASSIGN_FROM_DLSYM(clGetDeviceInfo); + ASSIGN_FROM_DLSYM(clGetDeviceIDs); + ASSIGN_FROM_DLSYM(clRetainEvent); + +#undef ASSIGN_FROM_DLSYM + + loaded_ = true; + dlclose(handle); + return true; +} + +bool OpenCLSupported() { return OpenCLStub::Get().loaded(); } + +} // namespace mace + +cl_int clBuildProgram(cl_program program, + cl_uint num_devices, + const cl_device_id *device_list, + const char *options, + void(CL_CALLBACK *pfn_notify)(cl_program program, + void *user_data), + void *user_data) { + auto func = mace::OpenCLStub::Get().clBuildProgram; + if (func != nullptr) { + return func(program, num_devices, device_list, options, pfn_notify, + user_data); + } else { + return CL_OUT_OF_RESOURCES; + } +} + +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 func = mace::OpenCLStub::Get().clEnqueueNDRangeKernel; + if (func != nullptr) { + 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_OUT_OF_RESOURCES; + } +} + +cl_int clSetKernelArg(cl_kernel kernel, + cl_uint arg_index, + size_t arg_size, + const void *arg_value) { + auto func = mace::OpenCLStub::Get().clSetKernelArg; + if (func != nullptr) { + return func(kernel, arg_index, arg_size, arg_value); + } else { + return CL_OUT_OF_RESOURCES; + } +} + +cl_int clRetainMemObject(cl_mem memobj) { + auto func = mace::OpenCLStub::Get().clRetainMemObject; + if (func != nullptr) { + return func(memobj); + } else { + return CL_OUT_OF_RESOURCES; + } +} + +cl_int clReleaseMemObject(cl_mem memobj) { + auto func = mace::OpenCLStub::Get().clReleaseMemObject; + if (func != nullptr) { + return func(memobj); + } else { + return CL_OUT_OF_RESOURCES; + } +} + +cl_int clEnqueueUnmapMemObject(cl_command_queue command_queue, + cl_mem memobj, + void *mapped_ptr, + cl_uint num_events_in_wait_list, + const cl_event *event_wait_list, + cl_event *event) { + auto func = mace::OpenCLStub::Get().clEnqueueUnmapMemObject; + if (func != nullptr) { + return func(command_queue, memobj, mapped_ptr, num_events_in_wait_list, + event_wait_list, event); + } else { + return CL_OUT_OF_RESOURCES; + } +} + +cl_int clRetainCommandQueue(cl_command_queue command_queue) { + auto func = mace::OpenCLStub::Get().clRetainCommandQueue; + if (func != nullptr) { + return func(command_queue); + } else { + return CL_OUT_OF_RESOURCES; + } +} + +cl_int clReleaseContext(cl_context context) { + auto func = mace::OpenCLStub::Get().clReleaseContext; + if (func != nullptr) { + return func(context); + } else { + return CL_OUT_OF_RESOURCES; + } +} +cl_int clReleaseEvent(cl_event event) { + auto func = mace::OpenCLStub::Get().clReleaseEvent; + if (func != nullptr) { + return func(event); + } else { + return CL_OUT_OF_RESOURCES; + } +} + +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 func = mace::OpenCLStub::Get().clEnqueueWriteBuffer; + if (func != nullptr) { + return func(command_queue, buffer, blocking_write, offset, size, ptr, + num_events_in_wait_list, event_wait_list, event); + } else { + return CL_OUT_OF_RESOURCES; + } +} + +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 func = mace::OpenCLStub::Get().clEnqueueReadBuffer; + if (func != nullptr) { + return func(command_queue, buffer, blocking_read, offset, size, ptr, + num_events_in_wait_list, event_wait_list, event); + } else { + return CL_OUT_OF_RESOURCES; + } +} + +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 func = mace::OpenCLStub::Get().clGetProgramBuildInfo; + if (func != nullptr) { + return func(program, device, param_name, param_value_size, param_value, + param_value_size_ret); + } else { + return CL_OUT_OF_RESOURCES; + } +} + +cl_int clRetainProgram(cl_program program) { + auto func = mace::OpenCLStub::Get().clRetainProgram; + if (func != nullptr) { + return func(program); + } else { + return CL_OUT_OF_RESOURCES; + } +} + +void *clEnqueueMapBuffer(cl_command_queue command_queue, + cl_mem buffer, + cl_bool blocking_map, + cl_map_flags map_flags, + size_t offset, + size_t size, + cl_uint num_events_in_wait_list, + const cl_event *event_wait_list, + cl_event *event, + cl_int *errcode_ret) { + auto func = mace::OpenCLStub::Get().clEnqueueMapBuffer; + if (func != nullptr) { + return func(command_queue, buffer, blocking_map, map_flags, offset, size, + num_events_in_wait_list, event_wait_list, event, errcode_ret); + } else { + if (errcode_ret != nullptr) { + *errcode_ret = CL_OUT_OF_RESOURCES; + } + return nullptr; + } +} + +cl_int clReleaseCommandQueue(cl_command_queue command_queue) { + auto func = mace::OpenCLStub::Get().clReleaseCommandQueue; + if (func != nullptr) { + return func(command_queue); + } else { + return CL_OUT_OF_RESOURCES; + } +} + +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 func = mace::OpenCLStub::Get().clCreateProgramWithBinary; + if (func != nullptr) { + return func(context, num_devices, device_list, lengths, binaries, + binary_status, errcode_ret); + } else { + if (errcode_ret != nullptr) { + *errcode_ret = CL_OUT_OF_RESOURCES; + } + return nullptr; + } +} + +cl_int clRetainContext(cl_context context) { + auto func = mace::OpenCLStub::Get().clRetainContext; + if (func != nullptr) { + return func(context); + } else { + return CL_OUT_OF_RESOURCES; + } +} + +cl_int clReleaseProgram(cl_program program) { + auto func = mace::OpenCLStub::Get().clReleaseProgram; + if (func != nullptr) { + return func(program); + } else { + return CL_OUT_OF_RESOURCES; + } +} + +cl_int clFlush(cl_command_queue command_queue) { + auto func = mace::OpenCLStub::Get().clFlush; + if (func != nullptr) { + return func(command_queue); + } else { + return CL_OUT_OF_RESOURCES; + } +} + +cl_int clFinish(cl_command_queue command_queue) { + auto func = mace::OpenCLStub::Get().clFinish; + if (func != nullptr) { + return func(command_queue); + } else { + return CL_OUT_OF_RESOURCES; + } +} + +cl_int clGetProgramInfo(cl_program program, + cl_program_info param_name, + size_t param_value_size, + void *param_value, + size_t *param_value_size_ret) { + auto func = mace::OpenCLStub::Get().clGetProgramInfo; + if (func != nullptr) { + return func(program, param_name, param_value_size, param_value, + param_value_size_ret); + } else { + return CL_OUT_OF_RESOURCES; + } +} + +cl_kernel clCreateKernel(cl_program program, + const char *kernel_name, + cl_int *errcode_ret) { + auto func = mace::OpenCLStub::Get().clCreateKernel; + if (func != nullptr) { + return func(program, kernel_name, errcode_ret); + } else { + if (errcode_ret != nullptr) { + *errcode_ret = CL_OUT_OF_RESOURCES; + } + return nullptr; + } +} + +cl_int clRetainKernel(cl_kernel kernel) { + auto func = mace::OpenCLStub::Get().clRetainKernel; + if (func != nullptr) { + return func(kernel); + } else { + return CL_OUT_OF_RESOURCES; + } +} + +cl_mem clCreateBuffer(cl_context context, + cl_mem_flags flags, + size_t size, + void *host_ptr, + cl_int *errcode_ret) { + auto func = mace::OpenCLStub::Get().clCreateBuffer; + if (func != nullptr) { + return func(context, flags, size, host_ptr, errcode_ret); + } else { + if (errcode_ret != nullptr) { + *errcode_ret = CL_OUT_OF_RESOURCES; + } + return nullptr; + } +} + +cl_program clCreateProgramWithSource(cl_context context, + cl_uint count, + const char **strings, + const size_t *lengths, + cl_int *errcode_ret) { + auto func = mace::OpenCLStub::Get().clCreateProgramWithSource; + if (func != nullptr) { + return func(context, count, strings, lengths, errcode_ret); + } else { + if (errcode_ret != nullptr) { + *errcode_ret = CL_OUT_OF_RESOURCES; + } + return nullptr; + } +} + +cl_int clReleaseKernel(cl_kernel kernel) { + auto func = mace::OpenCLStub::Get().clReleaseKernel; + if (func != nullptr) { + return func(kernel); + } else { + return CL_OUT_OF_RESOURCES; + } +} + +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 func = mace::OpenCLStub::Get().clGetDeviceIDs; + if (func != nullptr) { + return func(platform, device_type, num_entries, devices, num_devices); + } else { + return CL_OUT_OF_RESOURCES; + } +} + +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 func = mace::OpenCLStub::Get().clGetDeviceInfo; + if (func != nullptr) { + return func(device, param_name, param_value_size, param_value, + param_value_size_ret); + } else { + return CL_OUT_OF_RESOURCES; + } +} + +cl_int clRetainEvent(cl_event event) { + auto func = mace::OpenCLStub::Get().clRetainEvent; + if (func != nullptr) { + return func(event); + } else { + return CL_OUT_OF_RESOURCES; + } +} diff --git a/mace/core/platform/opencl/opencl_wrapper.h b/mace/core/platform/opencl/opencl_wrapper.h new file mode 100644 index 0000000000000000000000000000000000000000..1a9f2e17cf7436eec073afcadec2e5985499b719 --- /dev/null +++ b/mace/core/platform/opencl/opencl_wrapper.h @@ -0,0 +1,14 @@ +// +// Copyright (c) 2017 XiaoMi All rights reserved. +// + +#include "CL/opencl.h" + +#include "mace/core/logging.h" + +#include +#include + +namespace mace { +bool OpenCLSupported(); +} // namespace mace diff --git a/mace/third_party/opencl-headers.BUILD b/mace/third_party/opencl-headers.BUILD index bdda258ceb525200bcdaa2296a393569856fa23a..117ef0f46b4eebef7ba95a15b9326e1f85f244ea 100644 --- a/mace/third_party/opencl-headers.BUILD +++ b/mace/third_party/opencl-headers.BUILD @@ -1,54 +1,54 @@ -cc_inc_library( +cc_library( name = "opencl10_headers", hdrs = glob([ "opencl10/CL/*.h", ]), - prefix = "opencl10", + strip_include_prefix = "opencl10", visibility = ["//visibility:public"], ) -cc_inc_library( +cc_library( name = "opencl11_headers", hdrs = glob([ "opencl11/CL/*.h", ]), - prefix = "opencli11", + strip_include_prefix = "opencl11", visibility = ["//visibility:public"], ) -cc_inc_library( +cc_library( name = "opencl12_headers", hdrs = glob([ "opencl12/CL/*.h", ]), - prefix = "opencl12", + strip_include_prefix = "opencl12", visibility = ["//visibility:public"], ) -cc_inc_library( +cc_library( name = "opencl20_headers", hdrs = glob([ "opencl20/CL/*.h", ]), - prefix = "opencl20", + strip_include_prefix = "opencl20", visibility = ["//visibility:public"], ) -cc_inc_library( +cc_library( name = "opencl21_headers", hdrs = glob([ "opencl21/CL/*.h", ]), - prefix = "opencl21", + strip_include_prefix = "opencl21", visibility = ["//visibility:public"], ) -cc_inc_library( +cc_library( name = "opencl22_headers", hdrs = glob([ "opencl22/CL/*.h", ]), - prefix = "opencl22", + strip_include_prefix = "opencl22", visibility = ["//visibility:public"], )