diff --git a/intern/cycles/device/device_cuda.cpp b/intern/cycles/device/device_cuda.cpp index 0fb369ba50d..a630a3d1183 100644 --- a/intern/cycles/device/device_cuda.cpp +++ b/intern/cycles/device/device_cuda.cpp @@ -44,7 +44,7 @@ #include "util_types.h" #include "util_time.h" -#include "split/kernel_split_data.h" +#include "split/kernel_split_data_types.h" CCL_NAMESPACE_BEGIN diff --git a/intern/cycles/device/device_split_kernel.cpp b/intern/cycles/device/device_split_kernel.cpp index 1d58bfda117..b9705077fbf 100644 --- a/intern/cycles/device/device_split_kernel.cpp +++ b/intern/cycles/device/device_split_kernel.cpp @@ -17,7 +17,7 @@ #include "device_split_kernel.h" #include "kernel_types.h" -#include "kernel_split_data.h" +#include "kernel_split_data_types.h" #include "util_time.h" diff --git a/intern/cycles/device/opencl/opencl_split.cpp b/intern/cycles/device/opencl/opencl_split.cpp index aadd9f778b5..b651b4a848e 100644 --- a/intern/cycles/device/opencl/opencl_split.cpp +++ b/intern/cycles/device/opencl/opencl_split.cpp @@ -21,7 +21,7 @@ #include "buffers.h" #include "kernel_types.h" -#include "kernel_split_data.h" +#include "kernel_split_data_types.h" #include "device_split_kernel.h" diff --git a/intern/cycles/kernel/CMakeLists.txt b/intern/cycles/kernel/CMakeLists.txt index 6867ab02318..1c740b5c6eb 100644 --- a/intern/cycles/kernel/CMakeLists.txt +++ b/intern/cycles/kernel/CMakeLists.txt @@ -209,6 +209,7 @@ set(SRC_SPLIT_HEADERS split/kernel_shadow_blocked.h split/kernel_split_common.h split/kernel_split_data.h + split/kernel_split_data_types.h ) # CUDA module diff --git a/intern/cycles/kernel/kernel_globals.h b/intern/cycles/kernel/kernel_globals.h index 12d35787462..1c3884890bf 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 @@ -153,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__ */ diff --git a/intern/cycles/kernel/osl/osl_closures.cpp b/intern/cycles/kernel/osl/osl_closures.cpp index 4cb46254bc7..fe61587d179 100644 --- a/intern/cycles/kernel/osl/osl_closures.cpp +++ b/intern/cycles/kernel/osl/osl_closures.cpp @@ -42,7 +42,7 @@ #include "kernel_types.h" #include "kernel_compat_cpu.h" -#include "split/kernel_split_data.h" +#include "split/kernel_split_data_types.h" #include "kernel_globals.h" #include "kernel_montecarlo.h" #include "kernel_random.h" diff --git a/intern/cycles/kernel/osl/osl_services.cpp b/intern/cycles/kernel/osl/osl_services.cpp index 2d5b9c1c1cc..b08353e82d1 100644 --- a/intern/cycles/kernel/osl/osl_services.cpp +++ b/intern/cycles/kernel/osl/osl_services.cpp @@ -39,7 +39,7 @@ #include "util_string.h" #include "kernel_compat_cpu.h" -#include "split/kernel_split_data.h" +#include "split/kernel_split_data_types.h" #include "kernel_globals.h" #include "kernel_random.h" #include "kernel_projection.h" diff --git a/intern/cycles/kernel/osl/osl_shader.cpp b/intern/cycles/kernel/osl/osl_shader.cpp index 78848d7dfc9..c7e9f57b18a 100644 --- a/intern/cycles/kernel/osl/osl_shader.cpp +++ b/intern/cycles/kernel/osl/osl_shader.cpp @@ -19,7 +19,7 @@ #include "kernel_compat_cpu.h" #include "kernel_montecarlo.h" #include "kernel_types.h" -#include "split/kernel_split_data.h" +#include "split/kernel_split_data_types.h" #include "kernel_globals.h" #include "geom/geom_object.h" diff --git a/intern/cycles/kernel/split/kernel_split_data.h b/intern/cycles/kernel/split/kernel_split_data.h index 9c2793f6941..5380c0c5de6 100644 --- a/intern/cycles/kernel/split/kernel_split_data.h +++ b/intern/cycles/kernel/split/kernel_split_data.h @@ -17,85 +17,12 @@ #ifndef __KERNEL_SPLIT_DATA_H__ #define __KERNEL_SPLIT_DATA_H__ +#include "kernel_split_data_types.h" +#include "kernel_globals.h" + CCL_NAMESPACE_BEGIN -/* parameters used by the split kernels, we use a single struct to avoid passing these to each kernel */ - -typedef struct SplitParams { - int x; - int y; - int w; - int h; - - int offset; - int stride; - - ccl_global uint *rng_state; - - int start_sample; - int end_sample; - - ccl_global unsigned int *work_pools; - unsigned int num_samples; - - ccl_global int *queue_index; - int queue_size; - ccl_global char *use_queues_flag; - - ccl_global float *buffer; -} SplitParams; - -/* Global memory variables [porting]; These memory is used for - * co-operation between different kernels; Data written by one - * kernel will be available to another kernel via this global - * memory. - */ - -/* SPLIT_DATA_ENTRY(type, name, num) */ - -#if defined(WITH_CYCLES_DEBUG) || defined(__KERNEL_DEBUG__) -/* DebugData memory */ -# define SPLIT_DATA_DEBUG_ENTRIES \ - SPLIT_DATA_ENTRY(DebugData, debug_data, 1) -#else -# define SPLIT_DATA_DEBUG_ENTRIES -#endif - -#define SPLIT_DATA_ENTRIES \ - SPLIT_DATA_ENTRY(ccl_global RNG, rng, 1) \ - SPLIT_DATA_ENTRY(ccl_global float3, throughput, 1) \ - SPLIT_DATA_ENTRY(ccl_global float, L_transparent, 1) \ - SPLIT_DATA_ENTRY(PathRadiance, path_radiance, 1) \ - SPLIT_DATA_ENTRY(ccl_global Ray, ray, 1) \ - SPLIT_DATA_ENTRY(ccl_global PathState, path_state, 1) \ - SPLIT_DATA_ENTRY(Intersection, isect, 1) \ - SPLIT_DATA_ENTRY(ccl_global float3, ao_alpha, 1) \ - SPLIT_DATA_ENTRY(ccl_global float3, ao_bsdf, 1) \ - SPLIT_DATA_ENTRY(ccl_global Ray, ao_light_ray, 1) \ - SPLIT_DATA_ENTRY(ccl_global BsdfEval, bsdf_eval, 1) \ - SPLIT_DATA_ENTRY(ccl_global int, is_lamp, 1) \ - SPLIT_DATA_ENTRY(ccl_global Ray, light_ray, 1) \ - SPLIT_DATA_ENTRY(Intersection, isect_shadow, 2) \ - SPLIT_DATA_ENTRY(ccl_global int, queue_data, (NUM_QUEUES*2)) /* TODO(mai): this is too large? */ \ - SPLIT_DATA_ENTRY(ccl_global uint, work_array, 1) \ - SPLIT_DATA_ENTRY(ShaderData, sd, 1) \ - SPLIT_DATA_ENTRY(ShaderData, sd_DL_shadow, 2) \ - SPLIT_DATA_DEBUG_ENTRIES \ - -/* struct that holds pointers to data in the shared state buffer */ -typedef struct SplitData { -#define SPLIT_DATA_ENTRY(type, name, num) type *name; - SPLIT_DATA_ENTRIES -#undef SPLIT_DATA_ENTRY - - /* this is actually in a separate buffer from the rest of the split state data (so it can be read back from - * the host easily) but is still used the same as the other data so we have it here in this struct as well - */ - ccl_global char *ray_state; -} SplitData; - -/* TODO: find a way to get access to kg here */ -ccl_device_inline size_t split_data_buffer_size(ccl_global void *kg, size_t num_elements) +ccl_device_inline size_t split_data_buffer_size(KernelGlobals *kg, size_t num_elements) { (void)kg; /* Unused on CPU. */ @@ -107,7 +34,7 @@ ccl_device_inline size_t split_data_buffer_size(ccl_global void *kg, size_t num_ return size; } -ccl_device_inline void split_data_init(ccl_global void *kg, +ccl_device_inline void split_data_init(KernelGlobals *kg, ccl_global SplitData *split_data, size_t num_elements, ccl_global void *data, @@ -125,19 +52,6 @@ ccl_device_inline void split_data_init(ccl_global void *kg, split_data->ray_state = ray_state; } -#ifndef __KERNEL_CUDA__ -# define kernel_split_state (kg->split_data) -# define kernel_split_params (kg->split_param_data) -#else -__device__ SplitData __split_data; -# define kernel_split_state (__split_data) -__device__ SplitParams __split_param_data; -# define kernel_split_params (__split_param_data) -#endif /* __KERNEL_CUDA__ */ - CCL_NAMESPACE_END #endif /* __KERNEL_SPLIT_DATA_H__ */ - - - diff --git a/intern/cycles/kernel/split/kernel_split_data_types.h b/intern/cycles/kernel/split/kernel_split_data_types.h new file mode 100644 index 00000000000..62e3ea45ae2 --- /dev/null +++ b/intern/cycles/kernel/split/kernel_split_data_types.h @@ -0,0 +1,109 @@ +/* + * Copyright 2011-2016 Blender Foundation + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#ifndef __KERNEL_SPLIT_DATA_TYPES_H__ +#define __KERNEL_SPLIT_DATA_TYPES_H__ + +CCL_NAMESPACE_BEGIN + +/* parameters used by the split kernels, we use a single struct to avoid passing these to each kernel */ + +typedef struct SplitParams { + int x; + int y; + int w; + int h; + + int offset; + int stride; + + ccl_global uint *rng_state; + + int start_sample; + int end_sample; + + ccl_global unsigned int *work_pools; + unsigned int num_samples; + + ccl_global int *queue_index; + int queue_size; + ccl_global char *use_queues_flag; + + ccl_global float *buffer; +} SplitParams; + +/* Global memory variables [porting]; These memory is used for + * co-operation between different kernels; Data written by one + * kernel will be available to another kernel via this global + * memory. + */ + +/* SPLIT_DATA_ENTRY(type, name, num) */ + +#if defined(WITH_CYCLES_DEBUG) || defined(__KERNEL_DEBUG__) +/* DebugData memory */ +# define SPLIT_DATA_DEBUG_ENTRIES \ + SPLIT_DATA_ENTRY(DebugData, debug_data, 1) +#else +# define SPLIT_DATA_DEBUG_ENTRIES +#endif + +#define SPLIT_DATA_ENTRIES \ + SPLIT_DATA_ENTRY(ccl_global RNG, rng, 1) \ + SPLIT_DATA_ENTRY(ccl_global float3, throughput, 1) \ + SPLIT_DATA_ENTRY(ccl_global float, L_transparent, 1) \ + SPLIT_DATA_ENTRY(PathRadiance, path_radiance, 1) \ + SPLIT_DATA_ENTRY(ccl_global Ray, ray, 1) \ + SPLIT_DATA_ENTRY(ccl_global PathState, path_state, 1) \ + SPLIT_DATA_ENTRY(Intersection, isect, 1) \ + SPLIT_DATA_ENTRY(ccl_global float3, ao_alpha, 1) \ + SPLIT_DATA_ENTRY(ccl_global float3, ao_bsdf, 1) \ + SPLIT_DATA_ENTRY(ccl_global Ray, ao_light_ray, 1) \ + SPLIT_DATA_ENTRY(ccl_global BsdfEval, bsdf_eval, 1) \ + SPLIT_DATA_ENTRY(ccl_global int, is_lamp, 1) \ + SPLIT_DATA_ENTRY(ccl_global Ray, light_ray, 1) \ + SPLIT_DATA_ENTRY(Intersection, isect_shadow, 2) \ + SPLIT_DATA_ENTRY(ccl_global int, queue_data, (NUM_QUEUES*2)) /* TODO(mai): this is too large? */ \ + SPLIT_DATA_ENTRY(ccl_global uint, work_array, 1) \ + SPLIT_DATA_ENTRY(ShaderData, sd, 1) \ + SPLIT_DATA_ENTRY(ShaderData, sd_DL_shadow, 2) \ + SPLIT_DATA_DEBUG_ENTRIES \ + +/* struct that holds pointers to data in the shared state buffer */ +typedef struct SplitData { +#define SPLIT_DATA_ENTRY(type, name, num) type *name; + SPLIT_DATA_ENTRIES +#undef SPLIT_DATA_ENTRY + + /* this is actually in a separate buffer from the rest of the split state data (so it can be read back from + * the host easily) but is still used the same as the other data so we have it here in this struct as well + */ + ccl_global char *ray_state; +} SplitData; + +#ifndef __KERNEL_CUDA__ +# define kernel_split_state (kg->split_data) +# define kernel_split_params (kg->split_param_data) +#else +__device__ SplitData __split_data; +# define kernel_split_state (__split_data) +__device__ SplitParams __split_param_data; +# define kernel_split_params (__split_param_data) +#endif /* __KERNEL_CUDA__ */ + +CCL_NAMESPACE_END + +#endif /* __KERNEL_SPLIT_DATA_TYPES_H__ */