Cycles: Make all #include statements relative to cycles source directory

The idea is to make include statements more explicit and obvious where the file is coming from, additionally reducing chance of wrong header being picked up. For example, it was not obvious whether bvh.h was refferring to builder or traversal, whenter node.h is a generic graph node or a shader node and cases like that. Surely this might look obvious for the active developers, but after some time of not touching the code it becomes less obvious where file is coming from. This was briefly mentioned in T50824 and seems @brecht is fine with such explicitness, but need to agree with all active developers before committing this. Please note that this patch is lacking changes related on GPU/OpenCL support. This will be solved if/when we all agree this is a good idea to move forward. Reviewers: brecht, lukasstockner97, maiself, nirved, dingto, juicyfruit, swerner Reviewed By: lukasstockner97, maiself, nirved, dingto Subscribers: brecht Differential Revision: https://developer.blender.org/D2586
2017-03-28 20:39:14 +02:00
parent 61db9ee27a
commit 0579eaae1f
233 changed files with 1421 additions and 1464 deletions
--- a/intern/cycles/kernel/kernels/cpu/kernel.cpp
+++ b/intern/cycles/kernel/kernels/cpu/kernel.cpp
@@ -56,9 +56,9 @@
    /* do nothing */
 #endif

-#include "kernel.h"
+#include "kernel/kernel.h"
 #define KERNEL_ARCH cpu
-#include "kernel_cpu_impl.h"
+#include "kernel/kernels/cpu/kernel_cpu_impl.h"

 CCL_NAMESPACE_BEGIN

@@ -90,7 +90,7 @@ void kernel_tex_copy(KernelGlobals *kg,
 		kg->tname.width = width; \
 	}
 #define KERNEL_IMAGE_TEX(type, ttype, tname)
-#include "kernel_textures.h"
+#include "kernel/kernel_textures.h"

 	else if(strstr(name, "__tex_image_float4")) {
 		texture_image_float4 *tex = NULL;
--- a/intern/cycles/kernel/kernels/cpu/kernel_avx.cpp
+++ b/intern/cycles/kernel/kernels/cpu/kernel_avx.cpp
@@ -28,10 +28,10 @@
 #  define __KERNEL_AVX__
 #endif

-#include "util_optimization.h"
+#include "util/util_optimization.h"

 #ifdef WITH_CYCLES_OPTIMIZED_KERNEL_AVX
-#  include "kernel.h"
+#  include "kernel/kernel.h"
 #  define KERNEL_ARCH cpu_avx
-#  include "kernel_cpu_impl.h"
+#  include "kernel/kernels/cpu/kernel_cpu_impl.h"
 #endif  /* WITH_CYCLES_OPTIMIZED_KERNEL_AVX */
--- a/intern/cycles/kernel/kernels/cpu/kernel_avx2.cpp
+++ b/intern/cycles/kernel/kernels/cpu/kernel_avx2.cpp
@@ -29,10 +29,10 @@
 #  define __KERNEL_AVX2__
 #endif

-#include "util_optimization.h"
+#include "util/util_optimization.h"

 #ifdef WITH_CYCLES_OPTIMIZED_KERNEL_AVX2
-#  include "kernel.h"
+#  include "kernel/kernel.h"
 #  define KERNEL_ARCH cpu_avx2
-#  include "kernel_cpu_impl.h"
+#  include "kernel/kernels/cpu/kernel_cpu_impl.h"
 #endif  /* WITH_CYCLES_OPTIMIZED_KERNEL_AVX2 */
--- a/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h
+++ b/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h
@@ -20,39 +20,39 @@
 * simply includes this file without worry of copying actual implementation over.
 */

-#include "kernel_compat_cpu.h"
+#include "kernel/kernel_compat_cpu.h"

 #ifndef __SPLIT_KERNEL__
-#  include "kernel_math.h"
-#  include "kernel_types.h"
+#  include "kernel/kernel_math.h"
+#  include "kernel/kernel_types.h"

-#  include "split/kernel_split_data.h"
-#  include "kernel_globals.h"
+#  include "kernel/split/kernel_split_data.h"
+#  include "kernel/kernel_globals.h"

-#  include "kernel_cpu_image.h"
-#  include "kernel_film.h"
-#  include "kernel_path.h"
-#  include "kernel_path_branched.h"
-#  include "kernel_bake.h"
+#  include "kernel/kernels/cpu/kernel_cpu_image.h"
+#  include "kernel/kernel_film.h"
+#  include "kernel/kernel_path.h"
+#  include "kernel/kernel_path_branched.h"
+#  include "kernel/kernel_bake.h"
 #else
-#  include "split/kernel_split_common.h"
+#  include "kernel/split/kernel_split_common.h"

-#  include "split/kernel_data_init.h"
-#  include "split/kernel_path_init.h"
-#  include "split/kernel_scene_intersect.h"
-#  include "split/kernel_lamp_emission.h"
-#  include "split/kernel_do_volume.h"
-#  include "split/kernel_queue_enqueue.h"
-#  include "split/kernel_indirect_background.h"
-#  include "split/kernel_shader_eval.h"
-#  include "split/kernel_holdout_emission_blurring_pathtermination_ao.h"
-#  include "split/kernel_subsurface_scatter.h"
-#  include "split/kernel_direct_lighting.h"
-#  include "split/kernel_shadow_blocked_ao.h"
-#  include "split/kernel_shadow_blocked_dl.h"
-#  include "split/kernel_next_iteration_setup.h"
-#  include "split/kernel_indirect_subsurface.h"
-#  include "split/kernel_buffer_update.h"
+#  include "kernel/split/kernel_data_init.h"
+#  include "kernel/split/kernel_path_init.h"
+#  include "kernel/split/kernel_scene_intersect.h"
+#  include "kernel/split/kernel_lamp_emission.h"
+#  include "kernel/split/kernel_do_volume.h"
+#  include "kernel/split/kernel_queue_enqueue.h"
+#  include "kernel/split/kernel_indirect_background.h"
+#  include "kernel/split/kernel_shader_eval.h"
+#  include "kernel/split/kernel_holdout_emission_blurring_pathtermination_ao.h"
+#  include "kernel/split/kernel_subsurface_scatter.h"
+#  include "kernel/split/kernel_direct_lighting.h"
+#  include "kernel/split/kernel_shadow_blocked_ao.h"
+#  include "kernel/split/kernel_shadow_blocked_dl.h"
+#  include "kernel/split/kernel_next_iteration_setup.h"
+#  include "kernel/split/kernel_indirect_subsurface.h"
+#  include "kernel/split/kernel_buffer_update.h"
 #endif

 CCL_NAMESPACE_BEGIN
--- a/intern/cycles/kernel/kernels/cpu/kernel_split.cpp
+++ b/intern/cycles/kernel/kernels/cpu/kernel_split.cpp
@@ -57,7 +57,7 @@
    /* do nothing */
 #endif

-#include "kernel.h"
+#include "kernel/kernel.h"
 #define KERNEL_ARCH cpu
-#include "kernel_cpu_impl.h"
+#include "kernel/kernels/cpu/kernel_cpu_impl.h"

--- a/intern/cycles/kernel/kernels/cpu/kernel_split_avx.cpp
+++ b/intern/cycles/kernel/kernels/cpu/kernel_split_avx.cpp
@@ -29,10 +29,10 @@

 #define __SPLIT_KERNEL__

-#include "util_optimization.h"
+#include "util/util_optimization.h"

 #ifdef WITH_CYCLES_OPTIMIZED_KERNEL_AVX
-#  include "kernel.h"
+#  include "kernel/kernel.h"
 #  define KERNEL_ARCH cpu_avx
-#  include "kernel_cpu_impl.h"
+#  include "kernel/kernels/cpu/kernel_cpu_impl.h"
 #endif  /* WITH_CYCLES_OPTIMIZED_KERNEL_AVX */
--- a/intern/cycles/kernel/kernels/cpu/kernel_split_avx2.cpp
+++ b/intern/cycles/kernel/kernels/cpu/kernel_split_avx2.cpp
@@ -31,10 +31,10 @@

 #define __SPLIT_KERNEL__

-#include "util_optimization.h"
+#include "util/util_optimization.h"

 #ifdef WITH_CYCLES_OPTIMIZED_KERNEL_AVX2
-#  include "kernel.h"
+#  include "kernel/kernel.h"
 #  define KERNEL_ARCH cpu_avx2
-#  include "kernel_cpu_impl.h"
+#  include "kernel/kernels/cpu/kernel_cpu_impl.h"
 #endif  /* WITH_CYCLES_OPTIMIZED_KERNEL_AVX2 */
--- a/intern/cycles/kernel/kernels/cpu/kernel_split_sse2.cpp
+++ b/intern/cycles/kernel/kernels/cpu/kernel_split_sse2.cpp
@@ -25,10 +25,10 @@

 #define __SPLIT_KERNEL__

-#include "util_optimization.h"
+#include "util/util_optimization.h"

 #ifdef WITH_CYCLES_OPTIMIZED_KERNEL_SSE2
-#  include "kernel.h"
+#  include "kernel/kernel.h"
 #  define KERNEL_ARCH cpu_sse2
-#  include "kernel_cpu_impl.h"
+#  include "kernel/kernels/cpu/kernel_cpu_impl.h"
 #endif  /* WITH_CYCLES_OPTIMIZED_KERNEL_SSE2 */
--- a/intern/cycles/kernel/kernels/cpu/kernel_split_sse3.cpp
+++ b/intern/cycles/kernel/kernels/cpu/kernel_split_sse3.cpp
@@ -27,10 +27,10 @@

 #define __SPLIT_KERNEL__

-#include "util_optimization.h"
+#include "util/util_optimization.h"

 #ifdef WITH_CYCLES_OPTIMIZED_KERNEL_SSE3
-#  include "kernel.h"
+#  include "kernel/kernel.h"
 #  define KERNEL_ARCH cpu_sse3
-#  include "kernel_cpu_impl.h"
+#  include "kernel/kernels/cpu/kernel_cpu_impl.h"
 #endif  /* WITH_CYCLES_OPTIMIZED_KERNEL_SSE3 */
--- a/intern/cycles/kernel/kernels/cpu/kernel_split_sse41.cpp
+++ b/intern/cycles/kernel/kernels/cpu/kernel_split_sse41.cpp
@@ -28,10 +28,10 @@

 #define __SPLIT_KERNEL__

-#include "util_optimization.h"
+#include "util/util_optimization.h"

 #ifdef WITH_CYCLES_OPTIMIZED_KERNEL_SSE41
-#  include "kernel.h"
+#  include "kernel/kernel.h"
 #  define KERNEL_ARCH cpu_sse41
-#  include "kernel_cpu_impl.h"
+#  include "kernel/kernels/cpu/kernel_cpu_impl.h"
 #endif  /* WITH_CYCLES_OPTIMIZED_KERNEL_SSE41 */
--- a/intern/cycles/kernel/kernels/cpu/kernel_sse2.cpp
+++ b/intern/cycles/kernel/kernels/cpu/kernel_sse2.cpp
@@ -23,10 +23,10 @@
 #  define __KERNEL_SSE2__
 #endif

-#include "util_optimization.h"
+#include "util/util_optimization.h"

 #ifdef WITH_CYCLES_OPTIMIZED_KERNEL_SSE2
-#  include "kernel.h"
+#  include "kernel/kernel.h"
 #  define KERNEL_ARCH cpu_sse2
-#  include "kernel_cpu_impl.h"
+#  include "kernel/kernels/cpu/kernel_cpu_impl.h"
 #endif  /* WITH_CYCLES_OPTIMIZED_KERNEL_SSE2 */
--- a/intern/cycles/kernel/kernels/cpu/kernel_sse3.cpp
+++ b/intern/cycles/kernel/kernels/cpu/kernel_sse3.cpp
@@ -25,10 +25,10 @@
 #  define __KERNEL_SSSE3__
 #endif

-#include "util_optimization.h"
+#include "util/util_optimization.h"

 #ifdef WITH_CYCLES_OPTIMIZED_KERNEL_SSE3
-#  include "kernel.h"
+#  include "kernel/kernel.h"
 #  define KERNEL_ARCH cpu_sse3
-#  include "kernel_cpu_impl.h"
+#  include "kernel/kernels/cpu/kernel_cpu_impl.h"
 #endif  /* WITH_CYCLES_OPTIMIZED_KERNEL_SSE3 */
--- a/intern/cycles/kernel/kernels/cpu/kernel_sse41.cpp
+++ b/intern/cycles/kernel/kernels/cpu/kernel_sse41.cpp
@@ -26,10 +26,10 @@
 #  define __KERNEL_SSE41__
 #endif

-#include "util_optimization.h"
+#include "util/util_optimization.h"

 #ifdef WITH_CYCLES_OPTIMIZED_KERNEL_SSE41
-#  include "kernel.h"
+#  include "kernel/kernel.h"
 #  define KERNEL_ARCH cpu_sse41
-#  include "kernel_cpu_impl.h"
+#  include "kernel/kernels/cpu//kernel_cpu_impl.h"
 #endif  /* WITH_CYCLES_OPTIMIZED_KERNEL_SSE41 */
--- a/intern/cycles/kernel/kernels/cuda/kernel.cu
+++ b/intern/cycles/kernel/kernels/cuda/kernel.cu
@@ -18,15 +18,15 @@

 #ifdef __CUDA_ARCH__

-#include "../../kernel_compat_cuda.h"
+#include "kernel/kernel_compat_cuda.h"
 #include "kernel_config.h"
-#include "../../kernel_math.h"
-#include "../../kernel_types.h"
-#include "../../kernel_globals.h"
-#include "../../kernel_film.h"
-#include "../../kernel_path.h"
-#include "../../kernel_path_branched.h"
-#include "../../kernel_bake.h"
+#include "kernel/kernel_math.h"
+#include "kernel/kernel_types.h"
+#include "kernel/kernel_globals.h"
+#include "kernel/kernel_film.h"
+#include "kernel/kernel_path.h"
+#include "kernel/kernel_path_branched.h"
+#include "kernel/kernel_bake.h"

 /* kernels */
 extern "C" __global__ void
--- a/intern/cycles/kernel/kernels/cuda/kernel_split.cu
+++ b/intern/cycles/kernel/kernels/cuda/kernel_split.cu
@@ -20,28 +20,28 @@

 #define __SPLIT_KERNEL__

-#include "../../kernel_compat_cuda.h"
+#include "kernel/kernel_compat_cuda.h"
 #include "kernel_config.h"

-#include "../../split/kernel_split_common.h"
-#include "../../split/kernel_data_init.h"
-#include "../../split/kernel_path_init.h"
-#include "../../split/kernel_scene_intersect.h"
-#include "../../split/kernel_lamp_emission.h"
-#include "../../split/kernel_do_volume.h"
-#include "../../split/kernel_queue_enqueue.h"
-#include "../../split/kernel_indirect_background.h"
-#include "../../split/kernel_shader_eval.h"
-#include "../../split/kernel_holdout_emission_blurring_pathtermination_ao.h"
-#include "../../split/kernel_subsurface_scatter.h"
-#include "../../split/kernel_direct_lighting.h"
-#include "../../split/kernel_shadow_blocked_ao.h"
-#include "../../split/kernel_shadow_blocked_dl.h"
-#include "../../split/kernel_next_iteration_setup.h"
-#include "../../split/kernel_indirect_subsurface.h"
-#include "../../split/kernel_buffer_update.h"
+#include "kernel/split/kernel_split_common.h"
+#include "kernel/split/kernel_data_init.h"
+#include "kernel/split/kernel_path_init.h"
+#include "kernel/split/kernel_scene_intersect.h"
+#include "kernel/split/kernel_lamp_emission.h"
+#include "kernel/split/kernel_do_volume.h"
+#include "kernel/split/kernel_queue_enqueue.h"
+#include "kernel/split/kernel_indirect_background.h"
+#include "kernel/split/kernel_shader_eval.h"
+#include "kernel/split/kernel_holdout_emission_blurring_pathtermination_ao.h"
+#include "kernel/split/kernel_subsurface_scatter.h"
+#include "kernel/split/kernel_direct_lighting.h"
+#include "kernel/split/kernel_shadow_blocked_ao.h"
+#include "kernel/split/kernel_shadow_blocked_dl.h"
+#include "kernel/split/kernel_next_iteration_setup.h"
+#include "kernel/split/kernel_indirect_subsurface.h"
+#include "kernel/split/kernel_buffer_update.h"

-#include "../../kernel_film.h"
+#include "kernel/kernel_film.h"

 /* kernels */
 extern "C" __global__ void
--- a/intern/cycles/kernel/kernels/opencl/kernel.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel.cl
@@ -16,34 +16,34 @@

 /* OpenCL kernel entry points - unfinished */

-#include "../../kernel_compat_opencl.h"
-#include "../../kernel_math.h"
-#include "../../kernel_types.h"
-#include "../../kernel_globals.h"
-#include "../../kernel_image_opencl.h"
+#include "kernel/kernel_compat_opencl.h"
+#include "kernel/kernel_math.h"
+#include "kernel/kernel_types.h"
+#include "kernel/kernel_globals.h"
+#include "kernel/kernel_image_opencl.h"

-#include "../../kernel_film.h"
+#include "kernel/kernel_film.h"

 #if defined(__COMPILE_ONLY_MEGAKERNEL__) || !defined(__NO_BAKING__)
-#  include "../../kernel_path.h"
-#  include "../../kernel_path_branched.h"
+#  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_random.h"
-#  include "../../kernel_differential.h"
-#  include "../../kernel_montecarlo.h"
-#  include "../../kernel_projection.h"
-#  include "../../geom/geom.h"
-#  include "../../bvh/bvh.h"
+#  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_accumulate.h"
-#  include "../../kernel_camera.h"
-#  include "../../kernel_shader.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_bake.h"
+#include "kernel/kernel_bake.h"

 #ifdef __COMPILE_ONLY_MEGAKERNEL__

@@ -54,7 +54,7 @@ __kernel void kernel_ocl_path_trace(

 #define KERNEL_TEX(type, ttype, name) \
 	ccl_global type *name,
-#include "../../kernel_textures.h"
+#include "kernel/kernel_textures.h"

 	int sample,
 	int sx, int sy, int sw, int sh, int offset, int stride)
@@ -65,7 +65,7 @@ __kernel void kernel_ocl_path_trace(

 #define KERNEL_TEX(type, ttype, name) \
 	kg->name = name;
-#include "../../kernel_textures.h"
+#include "kernel/kernel_textures.h"

 	int x = sx + ccl_global_id(0);
 	int y = sy + ccl_global_id(1);
@@ -84,7 +84,7 @@ __kernel void kernel_ocl_shader(

 #define KERNEL_TEX(type, ttype, name) \
 	ccl_global type *name,
-#include "../../kernel_textures.h"
+#include "kernel/kernel_textures.h"

 	int type, int sx, int sw, int offset, int sample)
 {
@@ -94,7 +94,7 @@ __kernel void kernel_ocl_shader(

 #define KERNEL_TEX(type, ttype, name) \
 	kg->name = name;
-#include "../../kernel_textures.h"
+#include "kernel/kernel_textures.h"

 	int x = sx + ccl_global_id(0);

@@ -116,7 +116,7 @@ __kernel void kernel_ocl_bake(

 #define KERNEL_TEX(type, ttype, name) \
 	ccl_global type *name,
-#include "../../kernel_textures.h"
+#include "kernel/kernel_textures.h"

 	int type, int filter, int sx, int sw, int offset, int sample)
 {
@@ -126,7 +126,7 @@ __kernel void kernel_ocl_bake(

 #define KERNEL_TEX(type, ttype, name) \
 	kg->name = name;
-#include "../../kernel_textures.h"
+#include "kernel/kernel_textures.h"

 	int x = sx + ccl_global_id(0);

@@ -146,7 +146,7 @@ __kernel void kernel_ocl_convert_to_byte(

 #define KERNEL_TEX(type, ttype, name) \
 	ccl_global type *name,
-#include "../../kernel_textures.h"
+#include "kernel/kernel_textures.h"

 	float sample_scale,
 	int sx, int sy, int sw, int sh, int offset, int stride)
@@ -157,7 +157,7 @@ __kernel void kernel_ocl_convert_to_byte(

 #define KERNEL_TEX(type, ttype, name) \
 	kg->name = name;
-#include "../../kernel_textures.h"
+#include "kernel/kernel_textures.h"

 	int x = sx + ccl_global_id(0);
 	int y = sy + ccl_global_id(1);
@@ -173,7 +173,7 @@ __kernel void kernel_ocl_convert_to_half_float(

 #define KERNEL_TEX(type, ttype, name) \
 	ccl_global type *name,
-#include "../../kernel_textures.h"
+#include "kernel/kernel_textures.h"

 	float sample_scale,
 	int sx, int sy, int sw, int sh, int offset, int stride)
@@ -184,7 +184,7 @@ __kernel void kernel_ocl_convert_to_half_float(

 #define KERNEL_TEX(type, ttype, name) \
 	kg->name = name;
-#include "../../kernel_textures.h"
+#include "kernel/kernel_textures.h"

 	int x = sx + ccl_global_id(0);
 	int y = sy + ccl_global_id(1);
--- a/intern/cycles/kernel/kernels/opencl/kernel_buffer_update.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_buffer_update.cl
@@ -14,9 +14,9 @@
 * limitations under the License.
 */

-#include "kernel_compat_opencl.h"
-#include "split/kernel_split_common.h"
-#include "split/kernel_buffer_update.h"
+#include "kernel/kernel_compat_opencl.h"
+#include "kernel/split/kernel_split_common.h"
+#include "kernel/split/kernel_buffer_update.h"

 __kernel void kernel_ocl_path_trace_buffer_update(
        ccl_global char *kg,
--- a/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl
@@ -14,9 +14,9 @@
 * limitations under the License.
 */

-#include "kernel_compat_opencl.h"
-#include "split/kernel_split_common.h"
-#include "split/kernel_data_init.h"
+#include "kernel/kernel_compat_opencl.h"
+#include "kernel/split/kernel_split_common.h"
+#include "kernel/split/kernel_data_init.h"

 __kernel void kernel_ocl_path_trace_data_init(
        ccl_global char *kg,
@@ -28,7 +28,7 @@ __kernel void kernel_ocl_path_trace_data_init(

 #define KERNEL_TEX(type, ttype, name)                                   \
        ccl_global type *name,
-#include "../../kernel_textures.h"
+#include "kernel/kernel_textures.h"

        int start_sample,
        int end_sample,
@@ -48,7 +48,7 @@ __kernel void kernel_ocl_path_trace_data_init(
 	                 rng_state,

 #define KERNEL_TEX(type, ttype, name) name,
-#include "../../kernel_textures.h"
+#include "kernel/kernel_textures.h"

 	                 start_sample,
 	                 end_sample,
--- a/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl
@@ -14,9 +14,9 @@
 * limitations under the License.
 */

-#include "kernel_compat_opencl.h"
-#include "split/kernel_split_common.h"
-#include "split/kernel_direct_lighting.h"
+#include "kernel/kernel_compat_opencl.h"
+#include "kernel/split/kernel_split_common.h"
+#include "kernel/split/kernel_direct_lighting.h"

 __kernel void kernel_ocl_path_trace_direct_lighting(
        ccl_global char *kg,
--- a/intern/cycles/kernel/kernels/opencl/kernel_do_volume.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_do_volume.cl
@@ -14,9 +14,9 @@
 * limitations under the License.
 */

-#include "kernel_compat_opencl.h"
-#include "split/kernel_split_common.h"
-#include "split/kernel_do_volume.h"
+#include "kernel/kernel_compat_opencl.h"
+#include "kernel/split/kernel_split_common.h"
+#include "kernel/split/kernel_do_volume.h"

 __kernel void kernel_ocl_path_trace_do_volume(
        ccl_global char *kg,
--- a/intern/cycles/kernel/kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl
@@ -14,9 +14,9 @@
 * limitations under the License.
 */

-#include "kernel_compat_opencl.h"
-#include "split/kernel_split_common.h"
-#include "split/kernel_holdout_emission_blurring_pathtermination_ao.h"
+#include "kernel/kernel_compat_opencl.h"
+#include "kernel/split/kernel_split_common.h"
+#include "kernel/split/kernel_holdout_emission_blurring_pathtermination_ao.h"

 __kernel void kernel_ocl_path_trace_holdout_emission_blurring_pathtermination_ao(
        ccl_global char *kg,
--- a/intern/cycles/kernel/kernels/opencl/kernel_indirect_background.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_indirect_background.cl
@@ -14,9 +14,9 @@
 * limitations under the License.
 */

-#include "kernel_compat_opencl.h"
-#include "split/kernel_split_common.h"
-#include "split/kernel_indirect_background.h"
+#include "kernel/kernel_compat_opencl.h"
+#include "kernel/split/kernel_split_common.h"
+#include "kernel/split/kernel_indirect_background.h"

 __kernel void kernel_ocl_path_trace_indirect_background(
        ccl_global char *kg,
--- a/intern/cycles/kernel/kernels/opencl/kernel_indirect_subsurface.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_indirect_subsurface.cl
@@ -14,9 +14,9 @@
 * limitations under the License.
 */

-#include "kernel_compat_opencl.h"
-#include "split/kernel_split_common.h"
-#include "split/kernel_indirect_subsurface.h"
+#include "kernel/kernel_compat_opencl.h"
+#include "kernel/split/kernel_split_common.h"
+#include "kernel/split/kernel_indirect_subsurface.h"

 __kernel void kernel_ocl_path_trace_indirect_subsurface(
        ccl_global char *kg,
--- a/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl
@@ -14,9 +14,9 @@
 * limitations under the License.
 */

-#include "kernel_compat_opencl.h"
-#include "split/kernel_split_common.h"
-#include "split/kernel_lamp_emission.h"
+#include "kernel/kernel_compat_opencl.h"
+#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,
--- a/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl
@@ -14,9 +14,9 @@
 * limitations under the License.
 */

-#include "kernel_compat_opencl.h"
-#include "split/kernel_split_common.h"
-#include "split/kernel_next_iteration_setup.h"
+#include "kernel/kernel_compat_opencl.h"
+#include "kernel/split/kernel_split_common.h"
+#include "kernel/split/kernel_next_iteration_setup.h"

 __kernel void kernel_ocl_path_trace_next_iteration_setup(
        ccl_global char *kg,
--- a/intern/cycles/kernel/kernels/opencl/kernel_path_init.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_path_init.cl
@@ -14,9 +14,9 @@
 * limitations under the License.
 */

-#include "kernel_compat_opencl.h"
-#include "split/kernel_split_common.h"
-#include "split/kernel_path_init.h"
+#include "kernel/kernel_compat_opencl.h"
+#include "kernel/split/kernel_split_common.h"
+#include "kernel/split/kernel_path_init.h"

 __kernel void kernel_ocl_path_trace_path_init(
        ccl_global char *kg,
--- a/intern/cycles/kernel/kernels/opencl/kernel_queue_enqueue.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_queue_enqueue.cl
@@ -14,9 +14,9 @@
 * limitations under the License.
 */

-#include "kernel_compat_opencl.h"
-#include "split/kernel_split_common.h"
-#include "split/kernel_queue_enqueue.h"
+#include "kernel/kernel_compat_opencl.h"
+#include "kernel/split/kernel_split_common.h"
+#include "kernel/split/kernel_queue_enqueue.h"

 __kernel void kernel_ocl_path_trace_queue_enqueue(
        ccl_global char *kg,
--- a/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl
@@ -14,9 +14,9 @@
 * limitations under the License.
 */

-#include "kernel_compat_opencl.h"
-#include "split/kernel_split_common.h"
-#include "split/kernel_scene_intersect.h"
+#include "kernel/kernel_compat_opencl.h"
+#include "kernel/split/kernel_split_common.h"
+#include "kernel/split/kernel_scene_intersect.h"

 __kernel void kernel_ocl_path_trace_scene_intersect(
        ccl_global char *kg,
--- a/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl
@@ -14,9 +14,9 @@
 * limitations under the License.
 */

-#include "kernel_compat_opencl.h"
-#include "split/kernel_split_common.h"
-#include "split/kernel_shader_eval.h"
+#include "kernel/kernel_compat_opencl.h"
+#include "kernel/split/kernel_split_common.h"
+#include "kernel/split/kernel_shader_eval.h"

 __kernel void kernel_ocl_path_trace_shader_eval(
        ccl_global char *kg,
--- a/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_ao.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_ao.cl
@@ -14,9 +14,9 @@
 * limitations under the License.
 */

-#include "kernel_compat_opencl.h"
-#include "split/kernel_split_common.h"
-#include "split/kernel_shadow_blocked_ao.h"
+#include "kernel/kernel_compat_opencl.h"
+#include "kernel/split/kernel_split_common.h"
+#include "kernel/split/kernel_shadow_blocked_ao.h"

 __kernel void kernel_ocl_path_trace_shadow_blocked_ao(
        ccl_global char *kg,
--- a/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_dl.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_dl.cl
@@ -14,9 +14,9 @@
 * limitations under the License.
 */

-#include "kernel_compat_opencl.h"
-#include "split/kernel_split_common.h"
-#include "split/kernel_shadow_blocked_dl.h"
+#include "kernel/kernel_compat_opencl.h"
+#include "kernel/split/kernel_split_common.h"
+#include "kernel/split/kernel_shadow_blocked_dl.h"

 __kernel void kernel_ocl_path_trace_shadow_blocked_dl(
        ccl_global char *kg,
--- a/intern/cycles/kernel/kernels/opencl/kernel_split.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_split.cl
@@ -14,22 +14,22 @@
 * limitations under the License.
 */

-#include "kernel_state_buffer_size.cl"
-#include "kernel_data_init.cl"
-#include "kernel_path_init.cl"
+#include "kernel/kernels/opencl/kernel_state_buffer_size.cl"
+#include "kernel/kernels/opencl/kernel_data_init.cl"
+#include "kernel/kernels/opencl/kernel_path_init.cl"

-#include "kernel_scene_intersect.cl"
-#include "kernel_lamp_emission.cl"
-#include "kernel_do_volume.cl"
-#include "kernel_indirect_background.cl"
-#include "kernel_queue_enqueue.cl"
-#include "kernel_shader_eval.cl"
-#include "kernel_holdout_emission_blurring_pathtermination_ao.cl"
-#include "kernel_subsurface_scatter.cl"
-#include "kernel_direct_lighting.cl"
-#include "kernel_shadow_blocked_ao.cl"
-#include "kernel_shadow_blocked_dl.cl"
-#include "kernel_next_iteration_setup.cl"
-#include "kernel_indirect_subsurface.cl"
-#include "kernel_buffer_update.cl"
+#include "kernel/kernels/opencl/kernel_scene_intersect.cl"
+#include "kernel/kernels/opencl/kernel_lamp_emission.cl"
+#include "kernel/kernels/opencl/kernel_do_volume.cl"
+#include "kernel/kernels/opencl/kernel_indirect_background.cl"
+#include "kernel/kernels/opencl/kernel_queue_enqueue.cl"
+#include "kernel/kernels/opencl/kernel_shader_eval.cl"
+#include "kernel/kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl"
+#include "kernel/kernels/opencl/kernel_subsurface_scatter.cl"
+#include "kernel/kernels/opencl/kernel_direct_lighting.cl"
+#include "kernel/kernels/opencl/kernel_shadow_blocked_ao.cl"
+#include "kernel/kernels/opencl/kernel_shadow_blocked_dl.cl"
+#include "kernel/kernels/opencl/kernel_next_iteration_setup.cl"
+#include "kernel/kernels/opencl/kernel_indirect_subsurface.cl"
+#include "kernel/kernels/opencl/kernel_buffer_update.cl"

--- a/intern/cycles/kernel/kernels/opencl/kernel_state_buffer_size.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_state_buffer_size.cl
@@ -14,8 +14,8 @@
 * limitations under the License.
 */

-#include "kernel_compat_opencl.h"
-#include "split/kernel_split_common.h"
+#include "kernel/kernel_compat_opencl.h"
+#include "kernel/split/kernel_split_common.h"

 __kernel void kernel_ocl_path_trace_state_buffer_size(
        ccl_global char *kg,
--- a/intern/cycles/kernel/kernels/opencl/kernel_subsurface_scatter.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_subsurface_scatter.cl
@@ -14,9 +14,9 @@
 * limitations under the License.
 */

-#include "kernel_compat_opencl.h"
-#include "split/kernel_split_common.h"
-#include "split/kernel_subsurface_scatter.h"
+#include "kernel/kernel_compat_opencl.h"
+#include "kernel/split/kernel_split_common.h"
+#include "kernel/split/kernel_subsurface_scatter.h"

 __kernel void kernel_ocl_path_trace_subsurface_scatter(
        ccl_global char *kg,