diff options
Diffstat (limited to 'thirdparty/basis_universal/encoder/basisu_opencl.cpp')
| -rw-r--r-- | thirdparty/basis_universal/encoder/basisu_opencl.cpp | 1332 |
1 files changed, 1332 insertions, 0 deletions
diff --git a/thirdparty/basis_universal/encoder/basisu_opencl.cpp b/thirdparty/basis_universal/encoder/basisu_opencl.cpp new file mode 100644 index 0000000000..81e3090a26 --- /dev/null +++ b/thirdparty/basis_universal/encoder/basisu_opencl.cpp @@ -0,0 +1,1332 @@ +// basisu_opencl.cpp +// Copyright (C) 2019-2021 Binomial LLC. All Rights Reserved. +// +// Licensed 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. +#include "basisu_opencl.h" + +// If 1, the kernel source code will come from encoders/ocl_kernels.h. Otherwise, it will be read from the "ocl_kernels.cl" file in the current directory (for development). +#define BASISU_USE_OCL_KERNELS_HEADER (1) +#define BASISU_OCL_KERNELS_FILENAME "ocl_kernels.cl" + +#if BASISU_SUPPORT_OPENCL + +#include "basisu_enc.h" + +// We only use OpenCL v1.2 or less. +#define CL_TARGET_OPENCL_VERSION 120 + +#ifdef __APPLE__ +#include <OpenCL/opencl.h> +#else +#include <CL/cl.h> +#endif + +#define BASISU_OPENCL_ASSERT_ON_ANY_ERRORS (1) + +namespace basisu +{ +#if BASISU_USE_OCL_KERNELS_HEADER +#include "basisu_ocl_kernels.h" +#endif + + static void ocl_error_printf(const char* pFmt, ...) + { + va_list args; + va_start(args, pFmt); + error_vprintf(pFmt, args); + va_end(args); + +#if BASISU_OPENCL_ASSERT_ON_ANY_ERRORS + assert(0); +#endif + } + + class ocl + { + public: + ocl() + { + memset(&m_dev_fp_config, 0, sizeof(m_dev_fp_config)); + + m_ocl_mutex.lock(); + m_ocl_mutex.unlock(); + } + + ~ocl() + { + } + + bool is_initialized() const { return m_device_id != nullptr; } + + cl_device_id get_device_id() const { return m_device_id; } + cl_context get_context() const { return m_context; } + cl_command_queue get_command_queue() { return m_command_queue; } + cl_program get_program() const { return m_program; } + + bool init(bool force_serialization) + { + deinit(); + + interval_timer tm; + tm.start(); + + cl_uint num_platforms = 0; + cl_int ret = clGetPlatformIDs(0, NULL, &num_platforms); + if (ret != CL_SUCCESS) + { + ocl_error_printf("ocl::init: clGetPlatformIDs() failed with %i\n", ret); + return false; + } + + if ((!num_platforms) || (num_platforms > INT_MAX)) + { + ocl_error_printf("ocl::init: clGetPlatformIDs() returned an invalid number of num_platforms\n"); + return false; + } + + std::vector<cl_platform_id> platforms(num_platforms); + + ret = clGetPlatformIDs(num_platforms, platforms.data(), NULL); + if (ret != CL_SUCCESS) + { + ocl_error_printf("ocl::init: clGetPlatformIDs() failed\n"); + return false; + } + + cl_uint num_devices = 0; + ret = clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_GPU, 1, &m_device_id, &num_devices); + + if (ret == CL_DEVICE_NOT_FOUND) + { + ocl_error_printf("ocl::init: Couldn't get any GPU device ID's, trying CL_DEVICE_TYPE_CPU\n"); + + ret = clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_CPU, 1, &m_device_id, &num_devices); + } + + if (ret != CL_SUCCESS) + { + ocl_error_printf("ocl::init: Unable to get any device ID's\n"); + + m_device_id = nullptr; + return false; + } + + ret = clGetDeviceInfo(m_device_id, + CL_DEVICE_SINGLE_FP_CONFIG, + sizeof(m_dev_fp_config), + &m_dev_fp_config, + nullptr); + if (ret != CL_SUCCESS) + { + ocl_error_printf("ocl::init: clGetDeviceInfo() failed\n"); + return false; + } + + char plat_vers[256]; + size_t rv = 0; + ret = clGetPlatformInfo(platforms[0], CL_PLATFORM_VERSION, sizeof(plat_vers), plat_vers, &rv); + if (ret == CL_SUCCESS) + printf("OpenCL platform version: \"%s\"\n", plat_vers); + + // Serialize CL calls with the AMD driver to avoid lockups when multiple command queues per thread are used. This sucks, but what can we do? + m_use_mutex = (strstr(plat_vers, "AMD") != nullptr) || force_serialization; + + printf("Serializing OpenCL calls across threads: %u\n", (uint32_t)m_use_mutex); + + m_context = clCreateContext(nullptr, 1, &m_device_id, nullptr, nullptr, &ret); + if (ret != CL_SUCCESS) + { + ocl_error_printf("ocl::init: clCreateContext() failed\n"); + + m_device_id = nullptr; + m_context = nullptr; + return false; + } + + m_command_queue = clCreateCommandQueue(m_context, m_device_id, 0, &ret); + if (ret != CL_SUCCESS) + { + ocl_error_printf("ocl::init: clCreateCommandQueue() failed\n"); + + deinit(); + return false; + } + + printf("OpenCL init time: %3.3f secs\n", tm.get_elapsed_secs()); + + return true; + } + + bool deinit() + { + if (m_program) + { + clReleaseProgram(m_program); + m_program = nullptr; + } + + if (m_command_queue) + { + clReleaseCommandQueue(m_command_queue); + m_command_queue = nullptr; + } + + if (m_context) + { + clReleaseContext(m_context); + m_context = nullptr; + } + + m_device_id = nullptr; + + return true; + } + + cl_command_queue create_command_queue() + { + cl_serializer serializer(this); + + cl_int ret = 0; + cl_command_queue p = clCreateCommandQueue(m_context, m_device_id, 0, &ret); + if (ret != CL_SUCCESS) + return nullptr; + + return p; + } + + void destroy_command_queue(cl_command_queue p) + { + if (p) + { + cl_serializer serializer(this); + + clReleaseCommandQueue(p); + } + } + + bool init_program(const char* pSrc, size_t src_size) + { + cl_int ret; + + if (m_program != nullptr) + { + clReleaseProgram(m_program); + m_program = nullptr; + } + + m_program = clCreateProgramWithSource(m_context, 1, (const char**)&pSrc, (const size_t*)&src_size, &ret); + if (ret != CL_SUCCESS) + { + ocl_error_printf("ocl::init_program: clCreateProgramWithSource() failed!\n"); + return false; + } + + std::string options; + if (m_dev_fp_config & CL_FP_CORRECTLY_ROUNDED_DIVIDE_SQRT) + { + options += "-cl-fp32-correctly-rounded-divide-sqrt"; + } + + options += " -cl-std=CL1.2"; + //options += " -cl-opt-disable"; + //options += " -cl-mad-enable"; + //options += " -cl-fast-relaxed-math"; + + ret = clBuildProgram(m_program, 1, &m_device_id, + options.size() ? options.c_str() : nullptr, // options + nullptr, // notify + nullptr); // user_data + + if (ret != CL_SUCCESS) + { + const cl_int build_program_result = ret; + + size_t ret_val_size; + ret = clGetProgramBuildInfo(m_program, m_device_id, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size); + if (ret != CL_SUCCESS) + { + ocl_error_printf("ocl::init_program: clGetProgramBuildInfo() failed!\n"); + return false; + } + + std::vector<char> build_log(ret_val_size + 1); + + ret = clGetProgramBuildInfo(m_program, m_device_id, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log.data(), NULL); + + ocl_error_printf("\nclBuildProgram() failed with error %i:\n%s", build_program_result, build_log.data()); + + return false; + } + + return true; + } + + cl_kernel create_kernel(const char* pName) + { + if (!m_program) + return nullptr; + + cl_serializer serializer(this); + + cl_int ret; + cl_kernel kernel = clCreateKernel(m_program, pName, &ret); + if (ret != CL_SUCCESS) + { + ocl_error_printf("ocl::create_kernel: clCreateKernel() failed!\n"); + return nullptr; + } + + return kernel; + } + + bool destroy_kernel(cl_kernel k) + { + if (k) + { + cl_serializer serializer(this); + + cl_int ret = clReleaseKernel(k); + if (ret != CL_SUCCESS) + { + ocl_error_printf("ocl::destroy_kernel: clReleaseKernel() failed!\n"); + return false; + } + } + return true; + } + + cl_mem alloc_read_buffer(size_t size) + { + cl_serializer serializer(this); + + cl_int ret; + cl_mem obj = clCreateBuffer(m_context, CL_MEM_READ_ONLY, size, NULL, &ret); + if (ret != CL_SUCCESS) + { + ocl_error_printf("ocl::alloc_read_buffer: clCreateBuffer() failed!\n"); + return nullptr; + } + + return obj; + } + + cl_mem alloc_and_init_read_buffer(cl_command_queue command_queue, const void *pInit, size_t size) + { + cl_serializer serializer(this); + + cl_int ret; + cl_mem obj = clCreateBuffer(m_context, CL_MEM_READ_ONLY, size, NULL, &ret); + if (ret != CL_SUCCESS) + { + ocl_error_printf("ocl::alloc_and_init_read_buffer: clCreateBuffer() failed!\n"); + return nullptr; + } + +#if 0 + if (!write_to_buffer(command_queue, obj, pInit, size)) + { + destroy_buffer(obj); + return nullptr; + } +#else + ret = clEnqueueWriteBuffer(command_queue, obj, CL_TRUE, 0, size, pInit, 0, NULL, NULL); + if (ret != CL_SUCCESS) + { + ocl_error_printf("ocl::alloc_and_init_read_buffer: clEnqueueWriteBuffer() failed!\n"); + return nullptr; + } +#endif + + return obj; + } + + cl_mem alloc_write_buffer(size_t size) + { + cl_serializer serializer(this); + + cl_int ret; + cl_mem obj = clCreateBuffer(m_context, CL_MEM_WRITE_ONLY, size, NULL, &ret); + if (ret != CL_SUCCESS) + { + ocl_error_printf("ocl::alloc_write_buffer: clCreateBuffer() failed!\n"); + return nullptr; + } + + return obj; + } + + bool destroy_buffer(cl_mem buf) + { + if (buf) + { + cl_serializer serializer(this); + + cl_int ret = clReleaseMemObject(buf); + if (ret != CL_SUCCESS) + { + ocl_error_printf("ocl::destroy_buffer: clReleaseMemObject() failed!\n"); + return false; + } + } + + return true; + } + + bool write_to_buffer(cl_command_queue command_queue, cl_mem clmem, const void* d, const size_t m) + { + cl_serializer serializer(this); + + cl_int ret = clEnqueueWriteBuffer(command_queue, clmem, CL_TRUE, 0, m, d, 0, NULL, NULL); + if (ret != CL_SUCCESS) + { + ocl_error_printf("ocl::write_to_buffer: clEnqueueWriteBuffer() failed!\n"); + return false; + } + + return true; + } + + bool read_from_buffer(cl_command_queue command_queue, const cl_mem clmem, void* d, size_t m) + { + cl_serializer serializer(this); + + cl_int ret = clEnqueueReadBuffer(command_queue, clmem, CL_TRUE, 0, m, d, 0, NULL, NULL); + if (ret != CL_SUCCESS) + { + ocl_error_printf("ocl::read_from_buffer: clEnqueueReadBuffer() failed!\n"); + return false; + } + + return true; + } + + cl_mem create_read_image_u8(uint32_t width, uint32_t height, const void* pPixels, uint32_t bytes_per_pixel, bool normalized) + { + cl_image_format fmt = get_image_format(bytes_per_pixel, normalized); + + cl_image_desc desc; + memset(&desc, 0, sizeof(desc)); + desc.image_type = CL_MEM_OBJECT_IMAGE2D; + desc.image_width = width; + desc.image_height = height; + desc.image_row_pitch = width * bytes_per_pixel; + + cl_serializer serializer(this); + + cl_int ret; + cl_mem img = clCreateImage(m_context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, &fmt, &desc, (void*)pPixels, &ret); + if (ret != CL_SUCCESS) + { + ocl_error_printf("ocl::create_read_image_u8: clCreateImage() failed!\n"); + return nullptr; + } + + return img; + } + + cl_mem create_write_image_u8(uint32_t width, uint32_t height, uint32_t bytes_per_pixel, bool normalized) + { + cl_image_format fmt = get_image_format(bytes_per_pixel, normalized); + + cl_image_desc desc; + memset(&desc, 0, sizeof(desc)); + desc.image_type = CL_MEM_OBJECT_IMAGE2D; + desc.image_width = width; + desc.image_height = height; + + cl_serializer serializer(this); + + cl_int ret; + cl_mem img = clCreateImage(m_context, CL_MEM_WRITE_ONLY, &fmt, &desc, nullptr, &ret); + if (ret != CL_SUCCESS) + { + ocl_error_printf("ocl::create_write_image_u8: clCreateImage() failed!\n"); + return nullptr; + } + + return img; + } + + bool read_from_image(cl_command_queue command_queue, cl_mem img, void* pPixels, uint32_t ofs_x, uint32_t ofs_y, uint32_t width, uint32_t height) + { + cl_serializer serializer(this); + + size_t origin[3] = { ofs_x, ofs_y, 0 }, region[3] = { width, height, 1 }; + + cl_int err = clEnqueueReadImage(command_queue, img, CL_TRUE, origin, region, 0, 0, pPixels, 0, NULL, NULL); + if (err != CL_SUCCESS) + { + ocl_error_printf("ocl::read_from_image: clEnqueueReadImage() failed!\n"); + return false; + } + + return true; + } + + bool run_1D(cl_command_queue command_queue, const cl_kernel kernel, size_t num_items) + { + cl_serializer serializer(this); + + cl_int ret = clEnqueueNDRangeKernel(command_queue, kernel, + 1, // work_dim + nullptr, // global_work_offset + &num_items, // global_work_size + nullptr, // local_work_size + 0, // num_events_in_wait_list + nullptr, // event_wait_list + nullptr // event + ); + + if (ret != CL_SUCCESS) + { + ocl_error_printf("ocl::run_1D: clEnqueueNDRangeKernel() failed!\n"); + return false; + } + + return true; + } + + bool run_2D(cl_command_queue command_queue, const cl_kernel kernel, size_t width, size_t height) + { + cl_serializer serializer(this); + + size_t num_global_items[2] = { width, height }; + //size_t num_local_items[2] = { 1, 1 }; + + cl_int ret = clEnqueueNDRangeKernel(command_queue, kernel, + 2, // work_dim + nullptr, // global_work_offset + num_global_items, // global_work_size + nullptr, // local_work_size + 0, // num_events_in_wait_list + nullptr, // event_wait_list + nullptr // event + ); + + if (ret != CL_SUCCESS) + { + ocl_error_printf("ocl::run_2D: clEnqueueNDRangeKernel() failed!\n"); + return false; + } + + return true; + } + + bool run_2D(cl_command_queue command_queue, const cl_kernel kernel, size_t ofs_x, size_t ofs_y, size_t width, size_t height) + { + cl_serializer serializer(this); + + size_t global_ofs[2] = { ofs_x, ofs_y }; + size_t num_global_items[2] = { width, height }; + //size_t num_local_items[2] = { 1, 1 }; + + cl_int ret = clEnqueueNDRangeKernel(command_queue, kernel, + 2, // work_dim + global_ofs, // global_work_offset + num_global_items, // global_work_size + nullptr, // local_work_size + 0, // num_events_in_wait_list + nullptr, // event_wait_list + nullptr // event + ); + + if (ret != CL_SUCCESS) + { + ocl_error_printf("ocl::run_2D: clEnqueueNDRangeKernel() failed!\n"); + return false; + } + + return true; + } + + void flush(cl_command_queue command_queue) + { + cl_serializer serializer(this); + + clFlush(command_queue); + clFinish(command_queue); + } + + template<typename T> + bool set_kernel_arg(cl_kernel kernel, uint32_t index, const T& obj) + { + cl_serializer serializer(this); + + cl_int ret = clSetKernelArg(kernel, index, sizeof(T), (void*)&obj); + if (ret != CL_SUCCESS) + { + ocl_error_printf("ocl::set_kernel_arg: clSetKernelArg() failed!\n"); + return false; + } + return true; + } + + template<typename T> + bool set_kernel_args(cl_kernel kernel, const T& obj1) + { + cl_serializer serializer(this); + + cl_int ret = clSetKernelArg(kernel, 0, sizeof(T), (void*)&obj1); + if (ret != CL_SUCCESS) + { + ocl_error_printf("ocl::set_kernel_arg: clSetKernelArg() failed!\n"); + return false; + } + return true; + } + +#define BASISU_CHECK_ERR if (ret != CL_SUCCESS) { ocl_error_printf("ocl::set_kernel_args: clSetKernelArg() failed!\n"); return false; } + + template<typename T, typename U> + bool set_kernel_args(cl_kernel kernel, const T& obj1, const U& obj2) + { + cl_serializer serializer(this); + cl_int ret = clSetKernelArg(kernel, 0, sizeof(T), (void*)&obj1); BASISU_CHECK_ERR + ret = clSetKernelArg(kernel, 1, sizeof(U), (void*)&obj2); BASISU_CHECK_ERR + return true; + } + + template<typename T, typename U, typename V> + bool set_kernel_args(cl_kernel kernel, const T& obj1, const U& obj2, const V& obj3) + { + cl_serializer serializer(this); + cl_int ret = clSetKernelArg(kernel, 0, sizeof(T), (void*)&obj1); BASISU_CHECK_ERR + ret = clSetKernelArg(kernel, 1, sizeof(U), (void*)&obj2); BASISU_CHECK_ERR + ret = clSetKernelArg(kernel, 2, sizeof(V), (void*)&obj3); BASISU_CHECK_ERR + return true; + } + + template<typename T, typename U, typename V, typename W> + bool set_kernel_args(cl_kernel kernel, const T& obj1, const U& obj2, const V& obj3, const W& obj4) + { + cl_serializer serializer(this); + cl_int ret = clSetKernelArg(kernel, 0, sizeof(T), (void*)&obj1); BASISU_CHECK_ERR + ret = clSetKernelArg(kernel, 1, sizeof(U), (void*)&obj2); BASISU_CHECK_ERR + ret = clSetKernelArg(kernel, 2, sizeof(V), (void*)&obj3); BASISU_CHECK_ERR + ret = clSetKernelArg(kernel, 3, sizeof(W), (void*)&obj4); BASISU_CHECK_ERR + return true; + } + + template<typename T, typename U, typename V, typename W, typename X> + bool set_kernel_args(cl_kernel kernel, const T& obj1, const U& obj2, const V& obj3, const W& obj4, const X& obj5) + { + cl_serializer serializer(this); + cl_int ret = clSetKernelArg(kernel, 0, sizeof(T), (void*)&obj1); BASISU_CHECK_ERR + ret = clSetKernelArg(kernel, 1, sizeof(U), (void*)&obj2); BASISU_CHECK_ERR + ret = clSetKernelArg(kernel, 2, sizeof(V), (void*)&obj3); BASISU_CHECK_ERR + ret = clSetKernelArg(kernel, 3, sizeof(W), (void*)&obj4); BASISU_CHECK_ERR + ret = clSetKernelArg(kernel, 4, sizeof(X), (void*)&obj5); BASISU_CHECK_ERR + return true; + } + + template<typename T, typename U, typename V, typename W, typename X, typename Y> + bool set_kernel_args(cl_kernel kernel, const T& obj1, const U& obj2, const V& obj3, const W& obj4, const X& obj5, const Y& obj6) + { + cl_serializer serializer(this); + cl_int ret = clSetKernelArg(kernel, 0, sizeof(T), (void*)&obj1); BASISU_CHECK_ERR + ret = clSetKernelArg(kernel, 1, sizeof(U), (void*)&obj2); BASISU_CHECK_ERR + ret = clSetKernelArg(kernel, 2, sizeof(V), (void*)&obj3); BASISU_CHECK_ERR + ret = clSetKernelArg(kernel, 3, sizeof(W), (void*)&obj4); BASISU_CHECK_ERR + ret = clSetKernelArg(kernel, 4, sizeof(X), (void*)&obj5); BASISU_CHECK_ERR + ret = clSetKernelArg(kernel, 5, sizeof(Y), (void*)&obj6); BASISU_CHECK_ERR + return true; + } + + template<typename T, typename U, typename V, typename W, typename X, typename Y, typename Z> + bool set_kernel_args(cl_kernel kernel, const T& obj1, const U& obj2, const V& obj3, const W& obj4, const X& obj5, const Y& obj6, const Z& obj7) + { + cl_serializer serializer(this); + cl_int ret = clSetKernelArg(kernel, 0, sizeof(T), (void*)&obj1); BASISU_CHECK_ERR + ret = clSetKernelArg(kernel, 1, sizeof(U), (void*)&obj2); BASISU_CHECK_ERR + ret = clSetKernelArg(kernel, 2, sizeof(V), (void*)&obj3); BASISU_CHECK_ERR + ret = clSetKernelArg(kernel, 3, sizeof(W), (void*)&obj4); BASISU_CHECK_ERR + ret = clSetKernelArg(kernel, 4, sizeof(X), (void*)&obj5); BASISU_CHECK_ERR + ret = clSetKernelArg(kernel, 5, sizeof(Y), (void*)&obj6); BASISU_CHECK_ERR + ret = clSetKernelArg(kernel, 6, sizeof(Z), (void*)&obj7); BASISU_CHECK_ERR + return true; + } + + template<typename T, typename U, typename V, typename W, typename X, typename Y, typename Z, typename A> + bool set_kernel_args(cl_kernel kernel, const T& obj1, const U& obj2, const V& obj3, const W& obj4, const X& obj5, const Y& obj6, const Z& obj7, const A& obj8) + { + cl_serializer serializer(this); + cl_int ret = clSetKernelArg(kernel, 0, sizeof(T), (void*)&obj1); BASISU_CHECK_ERR + ret = clSetKernelArg(kernel, 1, sizeof(U), (void*)&obj2); BASISU_CHECK_ERR + ret = clSetKernelArg(kernel, 2, sizeof(V), (void*)&obj3); BASISU_CHECK_ERR + ret = clSetKernelArg(kernel, 3, sizeof(W), (void*)&obj4); BASISU_CHECK_ERR + ret = clSetKernelArg(kernel, 4, sizeof(X), (void*)&obj5); BASISU_CHECK_ERR + ret = clSetKernelArg(kernel, 5, sizeof(Y), (void*)&obj6); BASISU_CHECK_ERR + ret = clSetKernelArg(kernel, 6, sizeof(Z), (void*)&obj7); BASISU_CHECK_ERR + ret = clSetKernelArg(kernel, 7, sizeof(A), (void*)&obj8); BASISU_CHECK_ERR + return true; + } +#undef BASISU_CHECK_ERR + + private: + cl_device_id m_device_id = nullptr; + cl_context m_context = nullptr; + cl_command_queue m_command_queue = nullptr; + cl_program m_program = nullptr; + cl_device_fp_config m_dev_fp_config; + + bool m_use_mutex = false; + std::mutex m_ocl_mutex; + + // This helper object is used to optionally serialize all calls to the CL driver after initialization. + // Currently this is only used to work around race conditions in the Windows AMD driver. + struct cl_serializer + { + inline cl_serializer(const cl_serializer&); + cl_serializer& operator= (const cl_serializer&); + + inline cl_serializer(ocl *p) : m_p(p) + { + if (m_p->m_use_mutex) + m_p->m_ocl_mutex.lock(); + } + + inline ~cl_serializer() + { + if (m_p->m_use_mutex) + m_p->m_ocl_mutex.unlock(); + } + + private: + ocl* m_p; + }; + + cl_image_format get_image_format(uint32_t bytes_per_pixel, bool normalized) + { + cl_image_format fmt; + switch (bytes_per_pixel) + { + case 1: fmt.image_channel_order = CL_LUMINANCE; break; + case 2: fmt.image_channel_order = CL_RG; break; + case 3: fmt.image_channel_order = CL_RGB; break; + case 4: fmt.image_channel_order = CL_RGBA; break; + default: assert(0); fmt.image_channel_order = CL_LUMINANCE; break; + } + + fmt.image_channel_data_type = normalized ? CL_UNORM_INT8 : CL_UNSIGNED_INT8; + return fmt; + } + }; + + // Library blobal state + ocl g_ocl; + + bool opencl_init(bool force_serialization) + { + if (g_ocl.is_initialized()) + { + assert(0); + return false; + } + + if (!g_ocl.init(force_serialization)) + { + ocl_error_printf("opencl_init: Failed initializing OpenCL\n"); + return false; + } + + const char* pKernel_src = nullptr; + size_t kernel_src_size = 0; + uint8_vec kernel_src; + +#if BASISU_USE_OCL_KERNELS_HEADER + pKernel_src = reinterpret_cast<const char*>(ocl_kernels_cl); + kernel_src_size = ocl_kernels_cl_len; +#else + if (!read_file_to_vec(BASISU_OCL_KERNELS_FILENAME, kernel_src)) + { + ocl_error_printf("opencl_init: Cannot read OpenCL kernel source file \"%s\"\n", BASISU_OCL_KERNELS_FILENAME); + g_ocl.deinit(); + return false; + } + + pKernel_src = (char*)kernel_src.data(); + kernel_src_size = kernel_src.size(); +#endif + + if (!kernel_src_size) + { + ocl_error_printf("opencl_init: Invalid OpenCL kernel source file \"%s\"\n", BASISU_OCL_KERNELS_FILENAME); + g_ocl.deinit(); + return false; + } + + if (!g_ocl.init_program(pKernel_src, kernel_src_size)) + { + ocl_error_printf("opencl_init: Failed compiling OpenCL program\n"); + g_ocl.deinit(); + return false; + } + + printf("OpenCL support initialized successfully\n"); + + return true; + } + + void opencl_deinit() + { + g_ocl.deinit(); + } + + bool opencl_is_available() + { + return g_ocl.is_initialized(); + } + + struct opencl_context + { + uint32_t m_ocl_total_pixel_blocks; + cl_mem m_ocl_pixel_blocks; + + cl_command_queue m_command_queue; + + cl_kernel m_ocl_encode_etc1s_blocks_kernel; + cl_kernel m_ocl_refine_endpoint_clusterization_kernel; + cl_kernel m_ocl_encode_etc1s_from_pixel_cluster_kernel; + cl_kernel m_ocl_find_optimal_selector_clusters_for_each_block_kernel; + cl_kernel m_ocl_determine_selectors_kernel; + }; + + opencl_context_ptr opencl_create_context() + { + if (!opencl_is_available()) + { + ocl_error_printf("opencl_create_context: OpenCL not initialized\n"); + assert(0); + return nullptr; + } + + interval_timer tm; + tm.start(); + + opencl_context* pContext = static_cast<opencl_context * >(calloc(sizeof(opencl_context), 1)); + if (!pContext) + return nullptr; + + // To avoid driver bugs in some drivers - serialize this. Likely not necessary, we don't know. + // https://community.intel.com/t5/OpenCL-for-CPU/Bug-report-clCreateKernelsInProgram-is-not-thread-safe/td-p/1159771 + + pContext->m_command_queue = g_ocl.create_command_queue(); + if (!pContext->m_command_queue) + { + ocl_error_printf("opencl_create_context: Failed creating OpenCL command queue!\n"); + opencl_destroy_context(pContext); + return nullptr; + } + + pContext->m_ocl_encode_etc1s_blocks_kernel = g_ocl.create_kernel("encode_etc1s_blocks"); + if (!pContext->m_ocl_encode_etc1s_blocks_kernel) + { + ocl_error_printf("opencl_create_context: Failed creating OpenCL kernel encode_etc1s_block\n"); + opencl_destroy_context(pContext); + return nullptr; + } + + pContext->m_ocl_refine_endpoint_clusterization_kernel = g_ocl.create_kernel("refine_endpoint_clusterization"); + if (!pContext->m_ocl_refine_endpoint_clusterization_kernel) + { + ocl_error_printf("opencl_create_context: Failed creating OpenCL kernel refine_endpoint_clusterization\n"); + opencl_destroy_context(pContext); + return nullptr; + } + + pContext->m_ocl_encode_etc1s_from_pixel_cluster_kernel = g_ocl.create_kernel("encode_etc1s_from_pixel_cluster"); + if (!pContext->m_ocl_encode_etc1s_from_pixel_cluster_kernel) + { + ocl_error_printf("opencl_create_context: Failed creating OpenCL kernel encode_etc1s_from_pixel_cluster\n"); + opencl_destroy_context(pContext); + return nullptr; + } + + pContext->m_ocl_find_optimal_selector_clusters_for_each_block_kernel = g_ocl.create_kernel("find_optimal_selector_clusters_for_each_block"); + if (!pContext->m_ocl_find_optimal_selector_clusters_for_each_block_kernel) + { + ocl_error_printf("opencl_create_context: Failed creating OpenCL kernel find_optimal_selector_clusters_for_each_block\n"); + opencl_destroy_context(pContext); + return nullptr; + } + + pContext->m_ocl_determine_selectors_kernel = g_ocl.create_kernel("determine_selectors"); + if (!pContext->m_ocl_determine_selectors_kernel) + { + ocl_error_printf("opencl_create_context: Failed creating OpenCL kernel determine_selectors\n"); + opencl_destroy_context(pContext); + return nullptr; + } + + debug_printf("opencl_create_context: Elapsed time: %f secs\n", tm.get_elapsed_secs()); + + return pContext; + } + + void opencl_destroy_context(opencl_context_ptr pContext) + { + if (!pContext) + return; + + interval_timer tm; + tm.start(); + + g_ocl.destroy_buffer(pContext->m_ocl_pixel_blocks); + + g_ocl.destroy_kernel(pContext->m_ocl_determine_selectors_kernel); + g_ocl.destroy_kernel(pContext->m_ocl_find_optimal_selector_clusters_for_each_block_kernel); + g_ocl.destroy_kernel(pContext->m_ocl_encode_etc1s_from_pixel_cluster_kernel); + g_ocl.destroy_kernel(pContext->m_ocl_encode_etc1s_blocks_kernel); + g_ocl.destroy_kernel(pContext->m_ocl_refine_endpoint_clusterization_kernel); + + g_ocl.destroy_command_queue(pContext->m_command_queue); + + memset(pContext, 0, sizeof(opencl_context)); + + free(pContext); + + debug_printf("opencl_destroy_context: Elapsed time: %f secs\n", tm.get_elapsed_secs()); + } + +#pragma pack(push, 1) + struct cl_encode_etc1s_param_struct + { + int m_total_blocks; + int m_perceptual; + int m_total_perms; + }; +#pragma pack(pop) + + bool opencl_set_pixel_blocks(opencl_context_ptr pContext, uint32_t total_blocks, const cl_pixel_block* pPixel_blocks) + { + if (!opencl_is_available()) + return false; + + if (pContext->m_ocl_pixel_blocks) + { + g_ocl.destroy_buffer(pContext->m_ocl_pixel_blocks); + pContext->m_ocl_pixel_blocks = nullptr; + } + + pContext->m_ocl_pixel_blocks = g_ocl.alloc_and_init_read_buffer(pContext->m_command_queue, pPixel_blocks, sizeof(cl_pixel_block) * total_blocks); + if (!pContext->m_ocl_pixel_blocks) + return false; + + pContext->m_ocl_total_pixel_blocks = total_blocks; + + return true; + } + + bool opencl_encode_etc1s_blocks(opencl_context_ptr pContext, etc_block* pOutput_blocks, bool perceptual, uint32_t total_perms) + { + if (!opencl_is_available()) + return false; + + interval_timer tm; + tm.start(); + + assert(pContext->m_ocl_pixel_blocks); + if (!pContext->m_ocl_pixel_blocks) + return false; + + cl_encode_etc1s_param_struct ps; + ps.m_total_blocks = pContext->m_ocl_total_pixel_blocks; + ps.m_perceptual = perceptual; + ps.m_total_perms = total_perms; + + bool status = false; + + cl_mem vars = g_ocl.alloc_and_init_read_buffer(pContext->m_command_queue , &ps, sizeof(ps)); + cl_mem block_buf = g_ocl.alloc_write_buffer(sizeof(etc_block) * pContext->m_ocl_total_pixel_blocks); + + if (!vars || !block_buf) + goto exit; + + if (!g_ocl.set_kernel_args(pContext->m_ocl_encode_etc1s_blocks_kernel, vars, pContext->m_ocl_pixel_blocks, block_buf)) + goto exit; + + if (!g_ocl.run_2D(pContext->m_command_queue, pContext->m_ocl_encode_etc1s_blocks_kernel, pContext->m_ocl_total_pixel_blocks, 1)) + goto exit; + + if (!g_ocl.read_from_buffer(pContext->m_command_queue, block_buf, pOutput_blocks, pContext->m_ocl_total_pixel_blocks * sizeof(etc_block))) + goto exit; + + status = true; + + debug_printf("opencl_encode_etc1s_blocks: Elapsed time: %3.3f secs\n", tm.get_elapsed_secs()); + +exit: + g_ocl.destroy_buffer(block_buf); + g_ocl.destroy_buffer(vars); + + return status; + } + + bool opencl_encode_etc1s_pixel_clusters( + opencl_context_ptr pContext, + etc_block* pOutput_blocks, + uint32_t total_clusters, + const cl_pixel_cluster* pClusters, + uint64_t total_pixels, + const color_rgba* pPixels, const uint32_t* pPixel_weights, + bool perceptual, uint32_t total_perms) + { + if (!opencl_is_available()) + return false; + + interval_timer tm; + tm.start(); + + cl_encode_etc1s_param_struct ps; + ps.m_total_blocks = total_clusters; + ps.m_perceptual = perceptual; + ps.m_total_perms = total_perms; + + bool status = false; + + if (sizeof(size_t) == sizeof(uint32_t)) + { + if ( ((sizeof(cl_pixel_cluster) * total_clusters) > UINT32_MAX) || + ((sizeof(color_rgba) * total_pixels) > UINT32_MAX) || + ((sizeof(uint32_t) * total_pixels) > UINT32_MAX) ) + { + return false; + } + } + + cl_mem vars = g_ocl.alloc_and_init_read_buffer(pContext->m_command_queue , &ps, sizeof(ps)); + cl_mem input_clusters = g_ocl.alloc_and_init_read_buffer(pContext->m_command_queue, pClusters, (size_t)(sizeof(cl_pixel_cluster) * total_clusters)); + cl_mem input_pixels = g_ocl.alloc_and_init_read_buffer(pContext->m_command_queue, pPixels, (size_t)(sizeof(color_rgba) * total_pixels)); + cl_mem weights_buf = g_ocl.alloc_and_init_read_buffer(pContext->m_command_queue, pPixel_weights, (size_t)(sizeof(uint32_t) * total_pixels)); + cl_mem block_buf = g_ocl.alloc_write_buffer(sizeof(etc_block) * total_clusters); + + if (!vars || !input_clusters || !input_pixels || !weights_buf || !block_buf) + goto exit; + + if (!g_ocl.set_kernel_args(pContext->m_ocl_encode_etc1s_from_pixel_cluster_kernel, vars, input_clusters, input_pixels, weights_buf, block_buf)) + goto exit; + + if (!g_ocl.run_2D(pContext->m_command_queue, pContext->m_ocl_encode_etc1s_from_pixel_cluster_kernel, total_clusters, 1)) + goto exit; + + if (!g_ocl.read_from_buffer(pContext->m_command_queue, block_buf, pOutput_blocks, sizeof(etc_block) * total_clusters)) + goto exit; + + status = true; + + debug_printf("opencl_encode_etc1s_pixel_clusters: Elapsed time: %3.3f secs\n", tm.get_elapsed_secs()); + + exit: + g_ocl.destroy_buffer(block_buf); + g_ocl.destroy_buffer(weights_buf); + g_ocl.destroy_buffer(input_pixels); + g_ocl.destroy_buffer(input_clusters); + g_ocl.destroy_buffer(vars); + + return status; + } + +#pragma pack(push, 1) + struct cl_rec_param_struct + { + int m_total_blocks; + int m_perceptual; + }; +#pragma pack(pop) + + bool opencl_refine_endpoint_clusterization( + opencl_context_ptr pContext, + const cl_block_info_struct* pPixel_block_info, + uint32_t total_clusters, + const cl_endpoint_cluster_struct* pCluster_info, + const uint32_t* pSorted_block_indices, + uint32_t* pOutput_cluster_indices, + bool perceptual) + { + if (!opencl_is_available()) + return false; + + interval_timer tm; + tm.start(); + + assert(pContext->m_ocl_pixel_blocks); + if (!pContext->m_ocl_pixel_blocks) + return false; + + cl_rec_param_struct ps; + ps.m_total_blocks = pContext->m_ocl_total_pixel_blocks; + ps.m_perceptual = perceptual; + + bool status = false; + + cl_mem pixel_block_info = g_ocl.alloc_and_init_read_buffer(pContext->m_command_queue, pPixel_block_info, sizeof(cl_block_info_struct) * pContext->m_ocl_total_pixel_blocks); + cl_mem cluster_info = g_ocl.alloc_and_init_read_buffer(pContext->m_command_queue, pCluster_info, sizeof(cl_endpoint_cluster_struct) * total_clusters); + cl_mem sorted_block_indices = g_ocl.alloc_and_init_read_buffer(pContext->m_command_queue, pSorted_block_indices, sizeof(uint32_t) * pContext->m_ocl_total_pixel_blocks); + cl_mem output_buf = g_ocl.alloc_write_buffer(sizeof(uint32_t) * pContext->m_ocl_total_pixel_blocks); + + if (!pixel_block_info || !cluster_info || !sorted_block_indices || !output_buf) + goto exit; + + if (!g_ocl.set_kernel_args(pContext->m_ocl_refine_endpoint_clusterization_kernel, ps, pContext->m_ocl_pixel_blocks, pixel_block_info, cluster_info, sorted_block_indices, output_buf)) + goto exit; + + if (!g_ocl.run_2D(pContext->m_command_queue, pContext->m_ocl_refine_endpoint_clusterization_kernel, pContext->m_ocl_total_pixel_blocks, 1)) + goto exit; + + if (!g_ocl.read_from_buffer(pContext->m_command_queue, output_buf, pOutput_cluster_indices, pContext->m_ocl_total_pixel_blocks * sizeof(uint32_t))) + goto exit; + + debug_printf("opencl_refine_endpoint_clusterization: Elapsed time: %3.3f secs\n", tm.get_elapsed_secs()); + + status = true; + +exit: + g_ocl.destroy_buffer(pixel_block_info); + g_ocl.destroy_buffer(cluster_info); + g_ocl.destroy_buffer(sorted_block_indices); + g_ocl.destroy_buffer(output_buf); + + return status; + } + + bool opencl_find_optimal_selector_clusters_for_each_block( + opencl_context_ptr pContext, + const fosc_block_struct* pInput_block_info, // one per block + uint32_t total_input_selectors, + const fosc_selector_struct* pInput_selectors, + const uint32_t* pSelector_cluster_indices, + uint32_t* pOutput_selector_cluster_indices, // one per block + bool perceptual) + { + if (!opencl_is_available()) + return false; + + interval_timer tm; + tm.start(); + + assert(pContext->m_ocl_pixel_blocks); + if (!pContext->m_ocl_pixel_blocks) + return false; + + fosc_param_struct ps; + ps.m_total_blocks = pContext->m_ocl_total_pixel_blocks; + ps.m_perceptual = perceptual; + + bool status = false; + + cl_mem input_block_info = g_ocl.alloc_and_init_read_buffer(pContext->m_command_queue, pInput_block_info, sizeof(fosc_block_struct) * pContext->m_ocl_total_pixel_blocks); + cl_mem input_selectors = g_ocl.alloc_and_init_read_buffer(pContext->m_command_queue, pInput_selectors, sizeof(fosc_selector_struct) * total_input_selectors); + cl_mem selector_cluster_indices = g_ocl.alloc_and_init_read_buffer(pContext->m_command_queue, pSelector_cluster_indices, sizeof(uint32_t) * total_input_selectors); + cl_mem output_selector_cluster_indices = g_ocl.alloc_write_buffer(sizeof(uint32_t) * pContext->m_ocl_total_pixel_blocks); + + if (!input_block_info || !input_selectors || !selector_cluster_indices || !output_selector_cluster_indices) + goto exit; + + if (!g_ocl.set_kernel_args(pContext->m_ocl_find_optimal_selector_clusters_for_each_block_kernel, ps, pContext->m_ocl_pixel_blocks, input_block_info, input_selectors, selector_cluster_indices, output_selector_cluster_indices)) + goto exit; + + if (!g_ocl.run_2D(pContext->m_command_queue, pContext->m_ocl_find_optimal_selector_clusters_for_each_block_kernel, pContext->m_ocl_total_pixel_blocks, 1)) + goto exit; + + if (!g_ocl.read_from_buffer(pContext->m_command_queue, output_selector_cluster_indices, pOutput_selector_cluster_indices, pContext->m_ocl_total_pixel_blocks * sizeof(uint32_t))) + goto exit; + + debug_printf("opencl_find_optimal_selector_clusters_for_each_block: Elapsed time: %3.3f secs\n", tm.get_elapsed_secs()); + + status = true; + + exit: + g_ocl.destroy_buffer(input_block_info); + g_ocl.destroy_buffer(input_selectors); + g_ocl.destroy_buffer(selector_cluster_indices); + g_ocl.destroy_buffer(output_selector_cluster_indices); + + return status; + } + + bool opencl_determine_selectors( + opencl_context_ptr pContext, + const color_rgba* pInput_etc_color5_and_inten, + etc_block* pOutput_blocks, + bool perceptual) + { + if (!opencl_is_available()) + return false; + + interval_timer tm; + tm.start(); + + assert(pContext->m_ocl_pixel_blocks); + if (!pContext->m_ocl_pixel_blocks) + return false; + + ds_param_struct ps; + ps.m_total_blocks = pContext->m_ocl_total_pixel_blocks; + ps.m_perceptual = perceptual; + + bool status = false; + + cl_mem input_etc_color5_intens = g_ocl.alloc_and_init_read_buffer(pContext->m_command_queue, pInput_etc_color5_and_inten, sizeof(color_rgba) * pContext->m_ocl_total_pixel_blocks); + cl_mem output_blocks = g_ocl.alloc_write_buffer(sizeof(etc_block) * pContext->m_ocl_total_pixel_blocks); + + if (!input_etc_color5_intens || !output_blocks) + goto exit; + + if (!g_ocl.set_kernel_args(pContext->m_ocl_determine_selectors_kernel, ps, pContext->m_ocl_pixel_blocks, input_etc_color5_intens, output_blocks)) + goto exit; + + if (!g_ocl.run_2D(pContext->m_command_queue, pContext->m_ocl_determine_selectors_kernel, pContext->m_ocl_total_pixel_blocks, 1)) + goto exit; + + if (!g_ocl.read_from_buffer(pContext->m_command_queue, output_blocks, pOutput_blocks, pContext->m_ocl_total_pixel_blocks * sizeof(etc_block))) + goto exit; + + debug_printf("opencl_determine_selectors: Elapsed time: %3.3f secs\n", tm.get_elapsed_secs()); + + status = true; + + exit: + g_ocl.destroy_buffer(input_etc_color5_intens); + g_ocl.destroy_buffer(output_blocks); + + return status; + } + +#else +namespace basisu +{ + // No OpenCL support - all dummy functions that return false; + bool opencl_init(bool force_serialization) + { + BASISU_NOTE_UNUSED(force_serialization); + + return false; + } + + void opencl_deinit() + { + } + + bool opencl_is_available() + { + return false; + } + + opencl_context_ptr opencl_create_context() + { + return nullptr; + } + + void opencl_destroy_context(opencl_context_ptr context) + { + BASISU_NOTE_UNUSED(context); + } + + bool opencl_set_pixel_blocks(opencl_context_ptr pContext, uint32_t total_blocks, const cl_pixel_block* pPixel_blocks) + { + BASISU_NOTE_UNUSED(pContext); + BASISU_NOTE_UNUSED(total_blocks); + BASISU_NOTE_UNUSED(pPixel_blocks); + + return false; + } + + bool opencl_encode_etc1s_blocks(opencl_context_ptr pContext, etc_block* pOutput_blocks, bool perceptual, uint32_t total_perms) + { + BASISU_NOTE_UNUSED(pContext); + BASISU_NOTE_UNUSED(pOutput_blocks); + BASISU_NOTE_UNUSED(perceptual); + BASISU_NOTE_UNUSED(total_perms); + + return false; + } + + bool opencl_encode_etc1s_pixel_clusters( + opencl_context_ptr pContext, + etc_block* pOutput_blocks, + uint32_t total_clusters, + const cl_pixel_cluster* pClusters, + uint64_t total_pixels, + const color_rgba* pPixels, const uint32_t *pPixel_weights, + bool perceptual, uint32_t total_perms) + { + BASISU_NOTE_UNUSED(pContext); + BASISU_NOTE_UNUSED(pOutput_blocks); + BASISU_NOTE_UNUSED(total_clusters); + BASISU_NOTE_UNUSED(pClusters); + BASISU_NOTE_UNUSED(total_pixels); + BASISU_NOTE_UNUSED(pPixels); + BASISU_NOTE_UNUSED(pPixel_weights); + BASISU_NOTE_UNUSED(perceptual); + BASISU_NOTE_UNUSED(total_perms); + + return false; + } + + bool opencl_refine_endpoint_clusterization( + opencl_context_ptr pContext, + const cl_block_info_struct* pPixel_block_info, + uint32_t total_clusters, + const cl_endpoint_cluster_struct* pCluster_info, + const uint32_t* pSorted_block_indices, + uint32_t* pOutput_cluster_indices, + bool perceptual) + { + BASISU_NOTE_UNUSED(pContext); + BASISU_NOTE_UNUSED(pPixel_block_info); + BASISU_NOTE_UNUSED(total_clusters); + BASISU_NOTE_UNUSED(pCluster_info); + BASISU_NOTE_UNUSED(pSorted_block_indices); + BASISU_NOTE_UNUSED(pOutput_cluster_indices); + BASISU_NOTE_UNUSED(perceptual); + + return false; + } + + bool opencl_find_optimal_selector_clusters_for_each_block( + opencl_context_ptr pContext, + const fosc_block_struct* pInput_block_info, // one per block + uint32_t total_input_selectors, + const fosc_selector_struct* pInput_selectors, + const uint32_t* pSelector_cluster_indices, + uint32_t* pOutput_selector_cluster_indices, // one per block + bool perceptual) + { + BASISU_NOTE_UNUSED(pContext); + BASISU_NOTE_UNUSED(pInput_block_info); + BASISU_NOTE_UNUSED(total_input_selectors); + BASISU_NOTE_UNUSED(pInput_selectors); + BASISU_NOTE_UNUSED(pSelector_cluster_indices); + BASISU_NOTE_UNUSED(pOutput_selector_cluster_indices); + BASISU_NOTE_UNUSED(perceptual); + + return false; + } + + bool opencl_determine_selectors( + opencl_context_ptr pContext, + const color_rgba* pInput_etc_color5_and_inten, + etc_block* pOutput_blocks, + bool perceptual) + { + BASISU_NOTE_UNUSED(pContext); + BASISU_NOTE_UNUSED(pInput_etc_color5_and_inten); + BASISU_NOTE_UNUSED(pOutput_blocks); + BASISU_NOTE_UNUSED(perceptual); + + return false; + } + +#endif // BASISU_SUPPORT_OPENCL + +} // namespace basisu |