diff options
| author | Mai Lavelle <mai.lavelle@gmail.com> | 2017-05-12 02:23:49 +0300 |
|---|---|---|
| committer | Mai Lavelle <mai.lavelle@gmail.com> | 2017-06-10 11:08:49 +0300 |
| commit | eb293f59f2eb9847b8fd593ac2dde2781ac8ace1 (patch) | |
| tree | 7a939dbac4abcb34e10b9c5e1ac2c9c09dcc9b36 /intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl | |
| parent | 6238214159a4229ded91cd36d14c8e55ff427c28 (diff) | |
Cycles: Pass all buffers to each kernel call for OpenCL
Technically not passing all buffers used by a kernel is undefined
behavior. We haven't had any issues with this so far on AMD or
Nvidia, but it's known to be a problem with Intel and we received
a report from AMD that this is a problem on newer hardware, so we
need to make this change at some point.
Unfortunately there a cost to being correct, about 5% for the
benchmark scenes. For low sample counts it's even worse, I've
seen up to 50% slowdown. For the latter case I think adjusting
tile updating logic can help, but not sure what that would look
like yet (it would be just a few lines change however).
Diffstat (limited to 'intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl')
| -rw-r--r-- | intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl | 10 |
1 files changed, 4 insertions, 6 deletions
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl b/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl index 822d2287715..c314dc96c33 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl @@ -18,9 +18,7 @@ #include "kernel/split/kernel_split_common.h" #include "kernel/split/kernel_lamp_emission.h" -__kernel void kernel_ocl_path_trace_lamp_emission( - ccl_global char *kg, - ccl_constant KernelData *data) -{ - kernel_lamp_emission((KernelGlobals*)kg); -} +#define KERNEL_NAME lamp_emission +#include "kernel/kernels/opencl/kernel_split_function.h" +#undef KERNEL_NAME + |