From 949ab753bb2e2d0f76921ed6d716f074ce863f21 Mon Sep 17 00:00:00 2001 From: Jeroen Bakker Date: Wed, 20 Feb 2019 14:41:56 +0100 Subject: Cycles OpenCL: Remove OpenCL MegaKernel Using OpenCL MegaKernel has been slow and therefore not usefull. This patch will remove the mega kernel from the OpenCL codebase and the OpenCLDeviceBase class. T61736: removal of mega kernel T61703: baking does not work with mega kernel Tags: #cycles Differential Revision: https://developer.blender.org/D4383 --- intern/cycles/kernel/CMakeLists.txt | 2 +- intern/cycles/kernel/kernels/opencl/kernel.cl | 148 --------------------- intern/cycles/kernel/kernels/opencl/kernel_base.cl | 88 ++++++++++++ 3 files changed, 89 insertions(+), 149 deletions(-) delete mode 100644 intern/cycles/kernel/kernels/opencl/kernel.cl create mode 100644 intern/cycles/kernel/kernels/opencl/kernel_base.cl (limited to 'intern/cycles/kernel') diff --git a/intern/cycles/kernel/CMakeLists.txt b/intern/cycles/kernel/CMakeLists.txt index 0a2acd3f669..7332346a787 100644 --- a/intern/cycles/kernel/CMakeLists.txt +++ b/intern/cycles/kernel/CMakeLists.txt @@ -36,8 +36,8 @@ set(SRC_CUDA_KERNELS ) set(SRC_OPENCL_KERNELS - kernels/opencl/kernel.cl kernels/opencl/kernel_bake.cl + kernels/opencl/kernel_base.cl kernels/opencl/kernel_displace.cl kernels/opencl/kernel_background.cl kernels/opencl/kernel_state_buffer_size.cl diff --git a/intern/cycles/kernel/kernels/opencl/kernel.cl b/intern/cycles/kernel/kernels/opencl/kernel.cl deleted file mode 100644 index aa837e2ae87..00000000000 --- a/intern/cycles/kernel/kernels/opencl/kernel.cl +++ /dev/null @@ -1,148 +0,0 @@ -/* - * Copyright 2011-2013 Blender Foundation - * - * 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. - */ - -/* OpenCL kernel entry points - unfinished */ - -#include "kernel/kernel_compat_opencl.h" -#include "kernel/kernel_math.h" -#include "kernel/kernel_types.h" -#include "kernel/kernel_globals.h" -#include "kernel/kernel_color.h" -#include "kernel/kernels/opencl/kernel_opencl_image.h" - -#include "kernel/kernel_film.h" - -#if defined(__COMPILE_ONLY_MEGAKERNEL__) || !defined(__NO_BAKING__) -# include "kernel/kernel_path.h" -# include "kernel/kernel_path_branched.h" -#else /* __COMPILE_ONLY_MEGAKERNEL__ */ -/* Include only actually used headers for the case - * when path tracing kernels are not needed. - */ -# include "kernel/kernel_random.h" -# include "kernel/kernel_differential.h" -# include "kernel/kernel_montecarlo.h" -# include "kernel/kernel_projection.h" -# include "kernel/geom/geom.h" -# include "kernel/bvh/bvh.h" - -# include "kernel/kernel_accumulate.h" -# include "kernel/kernel_camera.h" -# include "kernel/kernel_shader.h" -#endif /* defined(__COMPILE_ONLY_MEGAKERNEL__) || !defined(__NO_BAKING__) */ - -#include "kernel/kernel_bake.h" - -#ifdef __COMPILE_ONLY_MEGAKERNEL__ - -__kernel void kernel_ocl_path_trace( - ccl_constant KernelData *data, - ccl_global float *buffer, - - KERNEL_BUFFER_PARAMS, - - int sample, - int sx, int sy, int sw, int sh, int offset, int stride) -{ - KernelGlobals kglobals, *kg = &kglobals; - - kg->data = data; - - kernel_set_buffer_pointers(kg, KERNEL_BUFFER_ARGS); - kernel_set_buffer_info(kg); - - int x = sx + ccl_global_id(0); - int y = sy + ccl_global_id(1); - bool thread_is_active = x < sx + sw && y < sy + sh; - if(thread_is_active) { - kernel_path_trace(kg, buffer, sample, x, y, offset, stride); - } - if(kernel_data.film.cryptomatte_passes) { - /* Make sure no thread is writing to the buffers. */ - ccl_barrier(CCL_LOCAL_MEM_FENCE); - if(thread_is_active) { - kernel_cryptomatte_post(kg, buffer, sample, x, y, offset, stride); - } - } -} - -#else /* __COMPILE_ONLY_MEGAKERNEL__ */ - -__kernel void kernel_ocl_convert_to_byte( - ccl_constant KernelData *data, - ccl_global uchar4 *rgba, - ccl_global float *buffer, - - KERNEL_BUFFER_PARAMS, - - float sample_scale, - int sx, int sy, int sw, int sh, int offset, int stride) -{ - KernelGlobals kglobals, *kg = &kglobals; - - kg->data = data; - - kernel_set_buffer_pointers(kg, KERNEL_BUFFER_ARGS); - kernel_set_buffer_info(kg); - - int x = sx + ccl_global_id(0); - int y = sy + ccl_global_id(1); - - if(x < sx + sw && y < sy + sh) - kernel_film_convert_to_byte(kg, rgba, buffer, sample_scale, x, y, offset, stride); -} - -__kernel void kernel_ocl_convert_to_half_float( - ccl_constant KernelData *data, - ccl_global uchar4 *rgba, - ccl_global float *buffer, - - KERNEL_BUFFER_PARAMS, - - float sample_scale, - int sx, int sy, int sw, int sh, int offset, int stride) -{ - KernelGlobals kglobals, *kg = &kglobals; - - kg->data = data; - - kernel_set_buffer_pointers(kg, KERNEL_BUFFER_ARGS); - kernel_set_buffer_info(kg); - - int x = sx + ccl_global_id(0); - int y = sy + ccl_global_id(1); - - if(x < sx + sw && y < sy + sh) - kernel_film_convert_to_half_float(kg, rgba, buffer, sample_scale, x, y, offset, stride); -} - -__kernel void kernel_ocl_zero_buffer(ccl_global float4 *buffer, uint64_t size, uint64_t offset) -{ - size_t i = ccl_global_id(0) + ccl_global_id(1) * ccl_global_size(0); - - if(i < size / sizeof(float4)) { - buffer[i+offset/sizeof(float4)] = make_float4(0.0f, 0.0f, 0.0f, 0.0f); - } - else if(i == size / sizeof(float4)) { - ccl_global uchar *b = (ccl_global uchar*)&buffer[i+offset/sizeof(float4)]; - - for(i = 0; i < size % sizeof(float4); i++) { - *(b++) = 0; - } - } -} - -#endif /* __COMPILE_ONLY_MEGAKERNEL__ */ diff --git a/intern/cycles/kernel/kernels/opencl/kernel_base.cl b/intern/cycles/kernel/kernels/opencl/kernel_base.cl new file mode 100644 index 00000000000..1c2d89e8a92 --- /dev/null +++ b/intern/cycles/kernel/kernels/opencl/kernel_base.cl @@ -0,0 +1,88 @@ +/* + * Copyright 2011-2013 Blender Foundation + * + * 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. + */ + +/* OpenCL base kernels entry points */ + +#include "kernel/kernel_compat_opencl.h" +#include "kernel/kernel_types.h" +#include "kernel/kernel_globals.h" + +#include "kernel/kernel_film.h" + + +__kernel void kernel_ocl_convert_to_byte( + ccl_constant KernelData *data, + ccl_global uchar4 *rgba, + ccl_global float *buffer, + + KERNEL_BUFFER_PARAMS, + + float sample_scale, + int sx, int sy, int sw, int sh, int offset, int stride) +{ + KernelGlobals kglobals, *kg = &kglobals; + + kg->data = data; + + kernel_set_buffer_pointers(kg, KERNEL_BUFFER_ARGS); + kernel_set_buffer_info(kg); + + int x = sx + ccl_global_id(0); + int y = sy + ccl_global_id(1); + + if(x < sx + sw && y < sy + sh) + kernel_film_convert_to_byte(kg, rgba, buffer, sample_scale, x, y, offset, stride); +} + +__kernel void kernel_ocl_convert_to_half_float( + ccl_constant KernelData *data, + ccl_global uchar4 *rgba, + ccl_global float *buffer, + + KERNEL_BUFFER_PARAMS, + + float sample_scale, + int sx, int sy, int sw, int sh, int offset, int stride) +{ + KernelGlobals kglobals, *kg = &kglobals; + + kg->data = data; + + kernel_set_buffer_pointers(kg, KERNEL_BUFFER_ARGS); + kernel_set_buffer_info(kg); + + int x = sx + ccl_global_id(0); + int y = sy + ccl_global_id(1); + + if(x < sx + sw && y < sy + sh) + kernel_film_convert_to_half_float(kg, rgba, buffer, sample_scale, x, y, offset, stride); +} + +__kernel void kernel_ocl_zero_buffer(ccl_global float4 *buffer, uint64_t size, uint64_t offset) +{ + size_t i = ccl_global_id(0) + ccl_global_id(1) * ccl_global_size(0); + + if(i < size / sizeof(float4)) { + buffer[i+offset/sizeof(float4)] = make_float4(0.0f, 0.0f, 0.0f, 0.0f); + } + else if(i == size / sizeof(float4)) { + ccl_global uchar *b = (ccl_global uchar*)&buffer[i+offset/sizeof(float4)]; + + for(i = 0; i < size % sizeof(float4); i++) { + *(b++) = 0; + } + } +} -- cgit v1.2.3