diff options
Diffstat (limited to 'intern/cycles/kernel/kernel_globals.h')
| -rw-r--r-- | intern/cycles/kernel/kernel_globals.h | 26 |
1 files changed, 20 insertions, 6 deletions
diff --git a/intern/cycles/kernel/kernel_globals.h b/intern/cycles/kernel/kernel_globals.h index 8e66a3a0340..c9c97ea977e 100644 --- a/intern/cycles/kernel/kernel_globals.h +++ b/intern/cycles/kernel/kernel_globals.h @@ -16,6 +16,9 @@ /* Constant Globals */ +#ifndef __KERNEL_GLOBALS_H__ +#define __KERNEL_GLOBALS_H__ + CCL_NAMESPACE_BEGIN /* On the CPU, we pass along the struct KernelGlobals to nearly everywhere in @@ -44,7 +47,7 @@ typedef struct KernelGlobals { # define KERNEL_TEX(type, ttype, name) ttype name; # define KERNEL_IMAGE_TEX(type, ttype, name) -# include "kernel_textures.h" +# include "kernel/kernel_textures.h" KernelData __data; @@ -64,6 +67,13 @@ typedef struct KernelGlobals { /* Storage for decoupled volume steps. */ VolumeStep *decoupled_volume_steps[2]; int decoupled_volume_steps_index; + + /* split kernel */ + SplitData split_data; + SplitParams split_param_data; + + int2 global_size; + int2 global_id; } KernelGlobals; #endif /* __KERNEL_CPU__ */ @@ -76,7 +86,10 @@ typedef struct KernelGlobals { #ifdef __KERNEL_CUDA__ __constant__ KernelData __data; -typedef struct KernelGlobals {} KernelGlobals; +typedef struct KernelGlobals { + /* NOTE: Keep the size in sync with SHADOW_STACK_MAX_HITS. */ + Intersection hits_stack[64]; +} KernelGlobals; # ifdef __KERNEL_CUDA_TEX_STORAGE__ # define KERNEL_TEX(type, ttype, name) ttype name; @@ -84,7 +97,7 @@ typedef struct KernelGlobals {} KernelGlobals; # define KERNEL_TEX(type, ttype, name) const __constant__ __device__ type *name; # endif # define KERNEL_IMAGE_TEX(type, ttype, name) ttype name; -# include "kernel_textures.h" +# include "kernel/kernel_textures.h" #endif /* __KERNEL_CUDA__ */ @@ -97,11 +110,11 @@ typedef ccl_addr_space struct KernelGlobals { # define KERNEL_TEX(type, ttype, name) \ ccl_global type *name; -# include "kernel_textures.h" +# include "kernel/kernel_textures.h" # ifdef __SPLIT_KERNEL__ - ShaderData *sd_input; - Intersection *isect_shadow; + SplitData split_data; + SplitParams split_param_data; # endif } KernelGlobals; @@ -143,3 +156,4 @@ ccl_device float lookup_table_read_2D(KernelGlobals *kg, float x, float y, int o CCL_NAMESPACE_END +#endif /* __KERNEL_GLOBALS_H__ */ |