[Bf-blender-cvs] [3c10ec9] master: Cycles: Enable object motion blur on Intel OpenCL platform
Sergey Sharybin
noreply at git.blender.org
Thu May 14 22:03:13 CEST 2015
Commit: 3c10ec96b5ec1c705362ccdaeaac8007050d081f
Author: Sergey Sharybin
Date: Fri May 15 00:48:12 2015 +0500
Branches: master
https://developer.blender.org/rB3c10ec96b5ec1c705362ccdaeaac8007050d081f
Cycles: Enable object motion blur on Intel OpenCL platform
This required allocating some memory related on object transform needed
by ShaderData and currently it is done for all the platforms. Since we're
targeting full feature-complete platforms this is rather acceptable at
this point and in the future we'll do selective NO_HAIR/NO_SSS/NO_BLUR
kernels.
This is experimental still and in fact there're some major issues on
NVidia platform and it's not really clear if it's a bug in compiler,
some uninitizlied variable or other kind of issue.
===================================================================
M intern/cycles/device/device_opencl.cpp
M intern/cycles/kernel/kernel_data_init.cl
M intern/cycles/kernel/kernel_types.h
===================================================================
diff --git a/intern/cycles/device/device_opencl.cpp b/intern/cycles/device/device_opencl.cpp
index 81edd25..d27cf8d 100644
--- a/intern/cycles/device/device_opencl.cpp
+++ b/intern/cycles/device/device_opencl.cpp
@@ -1580,6 +1580,10 @@ public:
cl_mem dPdu_sd, dPdv_sd;
cl_mem dPdu_sd_DL_shadow, dPdv_sd_DL_shadow;
+ /* Object motion. */
+ cl_mem ob_tfm_sd, ob_itfm_sd;
+ cl_mem ob_tfm_sd_DL_shadow, ob_itfm_sd_DL_shadow;
+
cl_mem closure_sd;
cl_mem closure_sd_DL_shadow;
cl_mem num_closure_sd;
@@ -1737,6 +1741,12 @@ public:
dPdu_sd_DL_shadow = NULL;
dPdv_sd_DL_shadow = NULL;
+ /* Object motion. */
+ ob_tfm_sd = NULL;
+ ob_itfm_sd = NULL;
+ ob_tfm_sd_DL_shadow = NULL;
+ ob_itfm_sd_DL_shadow = NULL;
+
closure_sd = NULL;
closure_sd_DL_shadow = NULL;
num_closure_sd = NULL;
@@ -2105,6 +2115,13 @@ public:
release_mem_object_safe(dPdv_sd);
release_mem_object_safe(dPdv_sd_DL_shadow);
+ /* Object motion. */
+ release_mem_object_safe(ob_tfm_sd);
+ release_mem_object_safe(ob_itfm_sd);
+
+ release_mem_object_safe(ob_tfm_sd_DL_shadow);
+ release_mem_object_safe(ob_itfm_sd_DL_shadow);
+
release_mem_object_safe(closure_sd);
release_mem_object_safe(closure_sd_DL_shadow);
release_mem_object_safe(num_closure_sd);
@@ -2300,6 +2317,12 @@ public:
dPdv_sd = mem_alloc(num_global_elements * sizeof(float3));
dPdv_sd_DL_shadow = mem_alloc(num_global_elements * 2 * sizeof(float3));
+ /* Object motion. */
+ ob_tfm_sd = mem_alloc(num_global_elements * sizeof(Transform));
+ ob_tfm_sd_DL_shadow = mem_alloc(num_global_elements * 2 * sizeof(Transform));
+ ob_itfm_sd = mem_alloc(num_global_elements * sizeof(float3));
+ ob_itfm_sd_DL_shadow = mem_alloc(num_global_elements * 2 * sizeof(Transform));
+
closure_sd = mem_alloc(num_global_elements * ShaderClosure_size);
closure_sd_DL_shadow = mem_alloc(num_global_elements * 2 * ShaderClosure_size);
num_closure_sd = mem_alloc(num_global_elements * sizeof(int));
@@ -2385,10 +2408,10 @@ public:
transparent_depth_sd,
transparent_depth_sd_DL_shadow);
+ /* Ray differentials. */
start_arg_index +=
kernel_set_args(ckPathTraceKernel_data_init,
start_arg_index,
- /* Ray differentials. */
dP_sd,
dP_sd_DL_shadow,
dI_sd,
@@ -2396,14 +2419,29 @@ public:
du_sd,
du_sd_DL_shadow,
dv_sd,
- dv_sd_DL_shadow,
+ dv_sd_DL_shadow);
- /* Dp/Du */
+ /* Dp/Du */
+ start_arg_index +=
+ kernel_set_args(ckPathTraceKernel_data_init,
+ start_arg_index,
dPdu_sd,
dPdu_sd_DL_shadow,
dPdv_sd,
- dPdv_sd_DL_shadow,
+ dPdv_sd_DL_shadow);
+
+ /* Object motion. */
+ start_arg_index +=
+ kernel_set_args(ckPathTraceKernel_data_init,
+ start_arg_index,
+ ob_tfm_sd,
+ ob_tfm_sd_DL_shadow,
+ ob_itfm_sd,
+ ob_itfm_sd_DL_shadow);
+ start_arg_index +=
+ kernel_set_args(ckPathTraceKernel_data_init,
+ start_arg_index,
closure_sd,
closure_sd_DL_shadow,
num_closure_sd,
diff --git a/intern/cycles/kernel/kernel_data_init.cl b/intern/cycles/kernel/kernel_data_init.cl
index 7982573..4c5f086 100644
--- a/intern/cycles/kernel/kernel_data_init.cl
+++ b/intern/cycles/kernel/kernel_data_init.cl
@@ -121,6 +121,13 @@ __kernel void kernel_ocl_path_trace_data_initialization(
ccl_global float3 *dPdv_sd,
ccl_global float3 *dPdv_sd_DL_shadow,
+ /* Object motion. */
+ ccl_global Transform *ob_tfm_sd,
+ ccl_global Transform *ob_tfm_sd_DL_shadow,
+
+ ccl_global Transform *ob_itfm_sd,
+ ccl_global Transform *ob_itfm_sd_DL_shadow,
+
ShaderClosure *closure_sd,
ShaderClosure *closure_sd_DL_shadow,
@@ -249,6 +256,14 @@ __kernel void kernel_ocl_path_trace_data_initialization(
#endif
#endif
+#ifdef __OBJECT_MOTION__
+ sd->ob_tfm = ob_tfm_sd;
+ sd_DL_shadow->ob_tfm = ob_tfm_sd_DL_shadow;
+
+ sd->ob_itfm = ob_itfm_sd;
+ sd_DL_shadow->ob_itfm = ob_itfm_sd_DL_shadow;
+#endif
+
sd->closure = closure_sd;
sd_DL_shadow->closure = closure_sd_DL_shadow;
diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h
index 7108d7b..8f91937 100644
--- a/intern/cycles/kernel/kernel_types.h
+++ b/intern/cycles/kernel/kernel_types.h
@@ -138,21 +138,7 @@ CCL_NAMESPACE_BEGIN
#ifdef __KERNEL_OPENCL_INTEL_CPU__
#define __CL_USE_NATIVE__
#define __KERNEL_SHADING__
-/* TODO(sergey): Advanced shading code still requires work
- * for split kernel.
- */
-# ifndef __SPLIT_KERNEL__
-# define __KERNEL_ADV_SHADING__
-# else
-# define __MULTI_CLOSURE__
-# define __TRANSPARENT_SHADOWS__
-# define __PASSES__
-# define __BACKGROUND_MIS__
-# define __LAMP_MIS__
-# define __AO__
-# define __HAIR__
-# define __CAMERA_MOTION__
-# endif
+#define __KERNEL_ADV_SHADING__
#endif
#endif // __KERNEL_OPENCL__
More information about the Bf-blender-cvs
mailing list