diff options
author | Jeroen Bakker <j.bakker@atmind.nl> | 2019-02-19 18:31:31 +0300 |
---|---|---|
committer | Jeroen Bakker <j.bakker@atmind.nl> | 2019-02-19 18:34:55 +0300 |
commit | 667033e89e7fe5241592e72e088a19723ca906b5 (patch) | |
tree | f6993fa977e76f8f81ed688a77dec32b44727927 /intern/cycles/kernel | |
parent | e6f5632eb11b37a2398f80841a77674656243dcf (diff) |
T61463: Separate Baking kernels
Cycles OpenCL: Split baking kernels in own program
Fix T61463. Before this patch baking was part of the base kernels. There
are 3 baking kernels that and all 3 uses shader evaluation. Only for one
of these kernels the functionality was wrapped in the __NO_BAKING__
compile directive.
When you start baking this leads to long compile times. By separating
in individual programs will reduce the compile times.
Also wrapped all baking kernels with __NO_BAKING__ to reduce the
compilation times.
Impact on compilation time
job | scene_name | previous | new | percentage
--------+-----------------+----------+-------+------------
T61463 | empty | 10.63 | 7.27 | 32%
T61463 | bmw | 17.91 | 14.24 | 20%
T61463 | fishycat | 19.57 | 15.08 | 23%
T61463 | barbershop | 54.10 | 48.18 | 11%
T61463 | classroom | 17.55 | 14.42 | 18%
T61463 | koro | 18.92 | 17.15 | 9%
T61463 | pavillion | 17.43 | 14.23 | 18%
T61463 | splash279 | 16.48 | 15.33 | 7%
T61463 | volume_emission | 36.22 | 34.19 | 6%
Impact on render time
job | scene_name | previous | new | percentage
--------+-----------------+----------+---------+------------
T61463 | empty | 21.06 | 20.54 | 2%
T61463 | bmw | 198.44 | 189.59 | 4%
T61463 | fishycat | 394.20 | 388.50 | 1%
T61463 | barbershop | 1188.16 | 1185.49 | 0%
T61463 | classroom | 341.08 | 339.27 | 1%
T61463 | koro | 472.43 | 360.70 | 24%
T61463 | pavillion | 905.77 | 902.14 | 0%
T61463 | splash279 | 55.26 | 54.92 | 1%
T61463 | volume_emission | 62.59 | 39.09 | 38%
I don't have a grounded explanation why koro and volume_emission is this much
faster; I have done several tests though...
Maniphest Tasks: T61463
Differential Revision: https://developer.blender.org/D4376
Diffstat (limited to 'intern/cycles/kernel')
-rw-r--r-- | intern/cycles/kernel/CMakeLists.txt | 3 | ||||
-rw-r--r-- | intern/cycles/kernel/kernels/opencl/kernel.cl | 72 | ||||
-rw-r--r-- | intern/cycles/kernel/kernels/opencl/kernel_background.cl | 39 | ||||
-rw-r--r-- | intern/cycles/kernel/kernels/opencl/kernel_bake.cl | 38 | ||||
-rw-r--r-- | intern/cycles/kernel/kernels/opencl/kernel_displace.cl | 40 |
5 files changed, 120 insertions, 72 deletions
diff --git a/intern/cycles/kernel/CMakeLists.txt b/intern/cycles/kernel/CMakeLists.txt index f7041ee2783..0a2acd3f669 100644 --- a/intern/cycles/kernel/CMakeLists.txt +++ b/intern/cycles/kernel/CMakeLists.txt @@ -37,6 +37,9 @@ set(SRC_CUDA_KERNELS set(SRC_OPENCL_KERNELS kernels/opencl/kernel.cl + kernels/opencl/kernel_bake.cl + kernels/opencl/kernel_displace.cl + kernels/opencl/kernel_background.cl kernels/opencl/kernel_state_buffer_size.cl kernels/opencl/kernel_split.cl kernels/opencl/kernel_split_bundle.cl diff --git a/intern/cycles/kernel/kernels/opencl/kernel.cl b/intern/cycles/kernel/kernels/opencl/kernel.cl index de1f5088629..aa837e2ae87 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel.cl @@ -81,78 +81,6 @@ __kernel void kernel_ocl_path_trace( #else /* __COMPILE_ONLY_MEGAKERNEL__ */ -__kernel void kernel_ocl_displace( - ccl_constant KernelData *data, - ccl_global uint4 *input, - ccl_global float4 *output, - - KERNEL_BUFFER_PARAMS, - - int type, int sx, int sw, int offset, int sample) -{ - 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); - - if(x < sx + sw) { - kernel_displace_evaluate(kg, input, output, x); - } -} -__kernel void kernel_ocl_background( - ccl_constant KernelData *data, - ccl_global uint4 *input, - ccl_global float4 *output, - - KERNEL_BUFFER_PARAMS, - - int type, int sx, int sw, int offset, int sample) -{ - 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); - - if(x < sx + sw) { - kernel_background_evaluate(kg, input, output, x); - } -} - -__kernel void kernel_ocl_bake( - ccl_constant KernelData *data, - ccl_global uint4 *input, - ccl_global float4 *output, - - KERNEL_BUFFER_PARAMS, - - int type, int filter, int sx, int sw, int offset, int sample) -{ - 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); - - if(x < sx + sw) { -#ifdef __NO_BAKING__ - output[x] = make_float4(0.0f, 0.0f, 0.0f, 0.0f); -#else - kernel_bake_evaluate(kg, input, output, (ShaderEvalType)type, filter, x, offset, sample); -#endif - } -} - __kernel void kernel_ocl_convert_to_byte( ccl_constant KernelData *data, ccl_global uchar4 *rgba, diff --git a/intern/cycles/kernel/kernels/opencl/kernel_background.cl b/intern/cycles/kernel/kernels/opencl/kernel_background.cl new file mode 100644 index 00000000000..c7c709c0ad7 --- /dev/null +++ b/intern/cycles/kernel/kernels/opencl/kernel_background.cl @@ -0,0 +1,39 @@ + +#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_path.h" +#include "kernel/kernel_path_branched.h" + +#include "kernel/kernel_bake.h" + +__kernel void kernel_ocl_background( + ccl_constant KernelData *data, + ccl_global uint4 *input, + ccl_global float4 *output, + + KERNEL_BUFFER_PARAMS, + + int type, int sx, int sw, int offset, int sample) +{ + 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); + + if(x < sx + sw) { +#ifdef __NO_BAKING__ + output[x] = make_float4(0.0f, 0.0f, 0.0f, 0.0f); +#else + kernel_background_evaluate(kg, input, output, x); +#endif + } +} diff --git a/intern/cycles/kernel/kernels/opencl/kernel_bake.cl b/intern/cycles/kernel/kernels/opencl/kernel_bake.cl new file mode 100644 index 00000000000..041312b53cb --- /dev/null +++ b/intern/cycles/kernel/kernels/opencl/kernel_bake.cl @@ -0,0 +1,38 @@ +#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_path.h" +#include "kernel/kernel_path_branched.h" + +#include "kernel/kernel_bake.h" + +__kernel void kernel_ocl_bake( + ccl_constant KernelData *data, + ccl_global uint4 *input, + ccl_global float4 *output, + + KERNEL_BUFFER_PARAMS, + + int type, int filter, int sx, int sw, int offset, int sample) +{ + 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); + + if(x < sx + sw) { +#ifdef __NO_BAKING__ + output[x] = make_float4(0.0f, 0.0f, 0.0f, 0.0f); +#else + kernel_bake_evaluate(kg, input, output, (ShaderEvalType)type, filter, x, offset, sample); +#endif + } +} diff --git a/intern/cycles/kernel/kernels/opencl/kernel_displace.cl b/intern/cycles/kernel/kernels/opencl/kernel_displace.cl new file mode 100644 index 00000000000..288bfd5eadc --- /dev/null +++ b/intern/cycles/kernel/kernels/opencl/kernel_displace.cl @@ -0,0 +1,40 @@ + +#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_path.h" +#include "kernel/kernel_path_branched.h" + +#include "kernel/kernel_bake.h + +__kernel void kernel_ocl_displace( + ccl_constant KernelData *data, + ccl_global uint4 *input, + ccl_global float4 *output, + + KERNEL_BUFFER_PARAMS, + + int type, int sx, int sw, int offset, int sample) +{ + 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); + + if(x < sx + sw) { +#ifdef __NO_BAKING__ + output[x] = make_float4(0.0f, 0.0f, 0.0f, 0.0f); +#else + kernel_displace_evaluate(kg, input, output, x); +#endif + } +} + |