| // 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 |