[Bf-blender-cvs] [712f7c3640] master: Cycles: Make it possible to access KernelGlobals from split data initialization function
Sergey Sharybin
noreply at git.blender.org
Wed Mar 8 11:14:14 CET 2017
Commit: 712f7c36402bd9b195d9e32dd72552e3dfc9ca24
Author: Sergey Sharybin
Date: Wed Mar 8 11:02:54 2017 +0100
Branches: master
https://developer.blender.org/rB712f7c36402bd9b195d9e32dd72552e3dfc9ca24
Cycles: Make it possible to access KernelGlobals from split data initialization function
===================================================================
M intern/cycles/device/device_cuda.cpp
M intern/cycles/device/device_split_kernel.cpp
M intern/cycles/device/opencl/opencl_split.cpp
M intern/cycles/kernel/CMakeLists.txt
M intern/cycles/kernel/kernel_globals.h
M intern/cycles/kernel/osl/osl_closures.cpp
M intern/cycles/kernel/osl/osl_services.cpp
M intern/cycles/kernel/osl/osl_shader.cpp
M intern/cycles/kernel/split/kernel_split_data.h
A intern/cycles/kernel/split/kernel_split_data_types.h
===================================================================
diff --git a/intern/cycles/device/device_cuda.cpp b/intern/cycles/device/device_cuda.cpp
index 0fb369ba50..a630a3d118 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 1d58bfda11..b9705077fb 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 aadd9f778b..b651b4a848 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 6867ab0231..1c740b5c6e 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 12d3578746..1c3884890b 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 4cb46254bc..fe61587d17 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 2d5b9c1c1c..b08353e82d 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 78848d7dfc..c7e9f57b18 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 9c2793f694..5380c0c5de 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__
-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.
- */
+#include "kernel_split_data_types.h"
+#include "kernel_globals.h"
-/* 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;
+CCL_NAMESPACE_BEGIN
-/* 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.h b/intern/cycles/kernel/split/kernel_split_data_types.h
similarity index 75%
copy from intern/cycles/kernel/split/kernel_split_data.h
copy to intern/cycles/kernel/split/kernel_split_data_types.h
index 9c2793f694..62e3ea45ae 100644
--- a/intern/cycles/kernel/split/kernel_split_data.h
+++ b/intern/cycles/kernel/split/kernel_split_data_types.h
@@ -14,8 +14,8 @@
* limitations under the License.
*/
-#ifndef __KERNEL_SPLIT_DATA_H__
-#define __KERNEL_SPLIT_DATA_H__
+#ifndef __KERNEL_SPLIT_DATA_TYPES_H__
+#define __KERNEL_SPLIT_DATA_TYPES_H__
CCL_NAMESPACE_BEGIN
@@ -94,37 +94,6 @@ typedef struct SplitData {
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)
-{
- (void)kg; /* Unused on CPU. */
-
- size_t size = 0;
-#define SPLIT_DATA_ENTRY(type, name, num) + align_up(num_elements * num * sizeof(type), 16)
- size = size SPLIT_DATA_ENTRIES;
-#undef SPLIT_DATA_ENTRY
-
- return size;
-}
-
-ccl_device_inline void split_data_init(ccl_global void *kg,
- ccl_global SplitData *split_data,
- size_t num_elements,
- ccl_global void *data,
- ccl_global char *ray_state)
-{
- (void)kg; /* Unused on CPU. */
-
- ccl_global char *p = (ccl_global char*)data;
-
-#define SPLIT_DATA_ENTRY(type, name, num) \
- split_data->name = (type*)p; p += align_up(num_elements * num * sizeof(type), 16);
- SPLIT_DATA_ENTRIES
-#undef SPLIT_DATA_ENTRY
-
- split_data->ray_state = ray_state;
-}
-
#ifndef __KERNEL_CUDA__
# define kernel_split_state (kg->split_data)
# define kernel_split_params (kg->split_param_data)
@@ -137,7 +106,4 @@ __device__ SplitParams __split_param_data;
CCL_NAMESPACE_END
-#endif /* __KERNEL_SPLIT_DATA_H__ */
-
-
-
+#endif /* __KERNEL_SPLIT_DATA_TYPES_H__ */
More information about the Bf-blender-cvs
mailing list