[Bf-blender-cvs] [cb5b0e0] cycles_kernel_split: Cycles kernel split: Move address space specification to DebugData typedef

Sergey Sharybin noreply at git.blender.org
Thu May 7 11:57:49 CEST 2015


Commit: cb5b0e037b4fd872778720fff9b6712b65ac9c34
Author: Sergey Sharybin
Date:   Thu May 7 14:56:13 2015 +0500
Branches: cycles_kernel_split
https://developer.blender.org/rBcb5b0e037b4fd872778720fff9b6712b65ac9c34

Cycles kernel split: Move address space specification to DebugData typedef

Debug pass does not really work, but it isn't caused by this change and would
need to have separate look.

===================================================================

M	intern/cycles/kernel/kernel_Background_BufferUpdate.cl
M	intern/cycles/kernel/kernel_DataInit.cl
M	intern/cycles/kernel/kernel_SceneIntersect.cl
M	intern/cycles/kernel/kernel_debug.h
M	intern/cycles/kernel/kernel_types.h

===================================================================

diff --git a/intern/cycles/kernel/kernel_Background_BufferUpdate.cl b/intern/cycles/kernel/kernel_Background_BufferUpdate.cl
index 9634e60..f5621f2 100644
--- a/intern/cycles/kernel/kernel_Background_BufferUpdate.cl
+++ b/intern/cycles/kernel/kernel_Background_BufferUpdate.cl
@@ -123,7 +123,7 @@ __kernel void kernel_ocl_path_trace_Background_BufferUpdate_SPLIT_KERNEL(
 	unsigned int num_samples,
 #endif
 #ifdef __KERNEL_DEBUG__
-	ccl_global DebugData *debugdata_coop,
+	DebugData *debugdata_coop,
 #endif
 	int parallel_samples                         /* Number of samples to be processed in parallel */
 	)
@@ -160,7 +160,7 @@ __kernel void kernel_ocl_path_trace_Background_BufferUpdate_SPLIT_KERNEL(
 		ShaderData *sd = (ShaderData *)shader_data;
 
 #ifdef __KERNEL_DEBUG__
-		ccl_global DebugData *debug_data = &debugdata_coop[ray_index];
+		DebugData *debug_data = &debugdata_coop[ray_index];
 #endif
 		ccl_global PathState *state = &PathState_coop[ray_index];
 		PathRadiance *L = L = &PathRadiance_coop[ray_index];
diff --git a/intern/cycles/kernel/kernel_DataInit.cl b/intern/cycles/kernel/kernel_DataInit.cl
index 8019f3d..15897dd 100644
--- a/intern/cycles/kernel/kernel_DataInit.cl
+++ b/intern/cycles/kernel/kernel_DataInit.cl
@@ -190,7 +190,7 @@ __kernel void kernel_ocl_path_trace_data_initialization_SPLIT_KERNEL(
 	unsigned int num_samples,                    /* Total number of samples per pixel */
 #endif
 #ifdef __KERNEL_DEBUG__
-	ccl_global DebugData *debugdata_coop,
+	DebugData *debugdata_coop,
 #endif
 	int parallel_samples                         /* Number of samples to be processed in parallel */
 	)
diff --git a/intern/cycles/kernel/kernel_SceneIntersect.cl b/intern/cycles/kernel/kernel_SceneIntersect.cl
index 34bb359..111fa25 100644
--- a/intern/cycles/kernel/kernel_SceneIntersect.cl
+++ b/intern/cycles/kernel/kernel_SceneIntersect.cl
@@ -77,7 +77,7 @@ __kernel void kernel_ocl_path_trace_SceneIntersect_SPLIT_KERNEL(
 	int queuesize,                              /* Size (capacity) of queues */
 	ccl_global char *use_queues_flag,           /* used to decide if this kernel should use queues to fetch ray index */
 #ifdef __KERNEL_DEBUG__
-	ccl_global DebugData *debugdata_coop,
+	DebugData *debugdata_coop,
 #endif
 	int parallel_samples                        /* Number of samples to be processed in parallel */
 	)
@@ -119,7 +119,7 @@ __kernel void kernel_ocl_path_trace_SceneIntersect_SPLIT_KERNEL(
 	KernelGlobals *kg = (KernelGlobals *)globals;
 
 #ifdef __KERNEL_DEBUG__
-	ccl_global DebugData *debug_data = &debugdata_coop[ray_index];
+	DebugData *debug_data = &debugdata_coop[ray_index];
 #endif
 	Intersection *isect = &Intersection_coop[ray_index];
 	PathState state = PathState_coop[ray_index];
diff --git a/intern/cycles/kernel/kernel_debug.h b/intern/cycles/kernel/kernel_debug.h
index 9c1de97..94ede39 100644
--- a/intern/cycles/kernel/kernel_debug.h
+++ b/intern/cycles/kernel/kernel_debug.h
@@ -16,7 +16,7 @@
 
 CCL_NAMESPACE_BEGIN
 
-ccl_device_inline void debug_data_init(ccl_addr_space DebugData *debug_data)
+ccl_device_inline void debug_data_init(DebugData *debug_data)
 {
 	debug_data->num_bvh_traversal_steps = 0;
 }
@@ -24,7 +24,7 @@ ccl_device_inline void debug_data_init(ccl_addr_space DebugData *debug_data)
 ccl_device_inline void kernel_write_debug_passes(KernelGlobals *kg,
                                                  ccl_global float *buffer,
                                                  ccl_addr_space PathState *state,
-                                                 ccl_addr_space DebugData *debug_data,
+                                                 DebugData *debug_data,
                                                  int sample)
 {
 	int flag = kernel_data.film.pass_flag;
diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h
index 434f792..ab20127 100644
--- a/intern/cycles/kernel/kernel_types.h
+++ b/intern/cycles/kernel/kernel_types.h
@@ -962,7 +962,7 @@ typedef struct KernelData {
 } KernelData;
 
 #ifdef __KERNEL_DEBUG__
-typedef struct DebugData {
+typedef ccl_addr_space struct DebugData {
 	// Total number of BVH node traversal steps and primitives intersections
 	// for the camera rays.
 	int num_bvh_traversal_steps;




More information about the Bf-blender-cvs mailing list