diff options
author | Jeroen Bakker <j.bakker@atmind.nl> | 2019-02-20 16:41:56 +0300 |
---|---|---|
committer | Jeroen Bakker <j.bakker@atmind.nl> | 2019-02-20 17:17:22 +0300 |
commit | 949ab753bb2e2d0f76921ed6d716f074ce863f21 (patch) | |
tree | cd39e16495b1c297557f9130273d0ccee45a4e16 /intern/cycles/kernel/kernels | |
parent | 667033e89e7fe5241592e72e088a19723ca906b5 (diff) |
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
Diffstat (limited to 'intern/cycles/kernel/kernels')
-rw-r--r-- | intern/cycles/kernel/kernels/opencl/kernel_base.cl (renamed from intern/cycles/kernel/kernels/opencl/kernel.cl) | 62 |
1 files changed, 1 insertions, 61 deletions
diff --git a/intern/cycles/kernel/kernels/opencl/kernel.cl b/intern/cycles/kernel/kernels/opencl/kernel_base.cl index aa837e2ae87..1c2d89e8a92 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_base.cl @@ -14,72 +14,14 @@ * limitations under the License. */ -/* OpenCL kernel entry points - unfinished */ +/* OpenCL base kernels entry points */ #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, @@ -144,5 +86,3 @@ __kernel void kernel_ocl_zero_buffer(ccl_global float4 *buffer, uint64_t size, u } } } - -#endif /* __COMPILE_ONLY_MEGAKERNEL__ */ |