[Bf-blender-cvs] [1ff753baa4] master: Cycles: Workaround for compilation error caused by passing KernelGlobals

Sergey Sharybin noreply at git.blender.org
Thu Mar 16 11:29:51 CET 2017 

Commit: 1ff753baa4bbf9aeb2c65e0d697840545bfbea24
Author: Sergey Sharybin
Date:   Wed Mar 8 12:49:04 2017 +0100
Branches: master
https://developer.blender.org/rB1ff753baa4bbf9aeb2c65e0d697840545bfbea24

Cycles: Workaround for compilation error caused by passing KernelGlobals

Pass globals as a bare pointer, same as it sued to be prior to split kernel rework.

AMD CPU platform and Intel OpenCL were complaining about this.

Perhaps we shouldn't pass globals as pointer at all, this isn't something what is
really portable and can cause issues on 32 bit perhaps.

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

M	intern/cycles/kernel/kernels/opencl/kernel_buffer_update.cl
M	intern/cycles/kernel/kernels/opencl/kernel_data_init.cl
M	intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl
M	intern/cycles/kernel/kernels/opencl/kernel_do_volume.cl
M	intern/cycles/kernel/kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl
M	intern/cycles/kernel/kernels/opencl/kernel_indirect_background.cl
M	intern/cycles/kernel/kernels/opencl/kernel_indirect_subsurface.cl
M	intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl
M	intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl
M	intern/cycles/kernel/kernels/opencl/kernel_path_init.cl
M	intern/cycles/kernel/kernels/opencl/kernel_queue_enqueue.cl
M	intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl
M	intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl
M	intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_ao.cl
M	intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_dl.cl
M	intern/cycles/kernel/kernels/opencl/kernel_state_buffer_size.cl
M	intern/cycles/kernel/kernels/opencl/kernel_subsurface_scatter.cl

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

diff --git a/intern/cycles/kernel/kernels/opencl/kernel_buffer_update.cl b/intern/cycles/kernel/kernels/opencl/kernel_buffer_update.cl
index d3058501f2..3c25d1d85a 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_buffer_update.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_buffer_update.cl
@@ -19,8 +19,8 @@
 #include "split/kernel_buffer_update.h"
 
 __kernel void kernel_ocl_path_trace_buffer_update(
-        KernelGlobals *kg,
+        ccl_global char *kg,
         ccl_constant KernelData *data)
 {
-	kernel_buffer_update(kg);
+	kernel_buffer_update((KernelGlobals*)kg);
 }
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl b/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl
index 1e3c4fa28c..54d4a577e1 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl
@@ -19,7 +19,7 @@
 #include "split/kernel_data_init.h"
 
 __kernel void kernel_ocl_path_trace_data_init(
-        KernelGlobals *kg,
+        ccl_global char *kg,
         ccl_constant KernelData *data,
         ccl_global void *split_data_buffer,
         int num_elements,
@@ -40,7 +40,7 @@ __kernel void kernel_ocl_path_trace_data_init(
         unsigned int num_samples,                    /* Total number of samples per pixel */
         ccl_global float *buffer)
 {
-	kernel_data_init(kg,
+	kernel_data_init((KernelGlobals*)kg,
 	                 data,
 	                 split_data_buffer,
 	                 num_elements,
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl b/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl
index 5d2f46b319..942a80f94f 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl
@@ -19,8 +19,8 @@
 #include "split/kernel_direct_lighting.h"
 
 __kernel void kernel_ocl_path_trace_direct_lighting(
-        KernelGlobals *kg,
+        ccl_global char *kg,
         ccl_constant KernelData *data)
 {
-	kernel_direct_lighting(kg);
+	kernel_direct_lighting((KernelGlobals*)kg);
 }
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_do_volume.cl b/intern/cycles/kernel/kernels/opencl/kernel_do_volume.cl
index 6380e9cb74..08187b0e03 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_do_volume.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_do_volume.cl
@@ -19,8 +19,8 @@
 #include "split/kernel_do_volume.h"
 
 __kernel void kernel_ocl_path_trace_do_volume(
-        KernelGlobals *kg,
+        ccl_global char *kg,
         ccl_constant KernelData *data)
 {
-	kernel_do_volume(kg);
+	kernel_do_volume((KernelGlobals*)kg);
 }
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl b/intern/cycles/kernel/kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl
index 7724b8a0bd..209080fecd 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl
@@ -19,8 +19,8 @@
 #include "split/kernel_holdout_emission_blurring_pathtermination_ao.h"
 
 __kernel void kernel_ocl_path_trace_holdout_emission_blurring_pathtermination_ao(
-        KernelGlobals *kg,
+        ccl_global char *kg,
         ccl_constant KernelData *data)
 {
-	kernel_holdout_emission_blurring_pathtermination_ao(kg);
+	kernel_holdout_emission_blurring_pathtermination_ao((KernelGlobals*)kg);
 }
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_indirect_background.cl b/intern/cycles/kernel/kernels/opencl/kernel_indirect_background.cl
index 671501bf23..b18fba4c01 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_indirect_background.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_indirect_background.cl
@@ -19,8 +19,8 @@
 #include "split/kernel_indirect_background.h"
 
 __kernel void kernel_ocl_path_trace_indirect_background(
-        KernelGlobals *kg,
+        ccl_global char *kg,
         ccl_constant KernelData *data)
 {
-	kernel_indirect_background(kg);
+	kernel_indirect_background((KernelGlobals*)kg);
 }
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_indirect_subsurface.cl b/intern/cycles/kernel/kernels/opencl/kernel_indirect_subsurface.cl
index b5e52e81eb..ce2e96ad78 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_indirect_subsurface.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_indirect_subsurface.cl
@@ -19,8 +19,8 @@
 #include "split/kernel_indirect_subsurface.h"
 
 __kernel void kernel_ocl_path_trace_indirect_subsurface(
-        KernelGlobals *kg,
+        ccl_global char *kg,
         ccl_constant KernelData *data)
 {
-	kernel_indirect_subsurface(kg);
+	kernel_indirect_subsurface((KernelGlobals*)kg);
 }
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl b/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl
index 2b84d0ea43..830e4e373a 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl
@@ -19,8 +19,8 @@
 #include "split/kernel_lamp_emission.h"
 
 __kernel void kernel_ocl_path_trace_lamp_emission(
-        KernelGlobals *kg,
+        ccl_global char *kg,
         ccl_constant KernelData *data)
 {
-	kernel_lamp_emission(kg);
+	kernel_lamp_emission((KernelGlobals*)kg);
 }
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl b/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl
index e87e367fb9..2a007e39c3 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl
@@ -19,8 +19,8 @@
 #include "split/kernel_next_iteration_setup.h"
 
 __kernel void kernel_ocl_path_trace_next_iteration_setup(
-        KernelGlobals *kg,
+        ccl_global char *kg,
         ccl_constant KernelData *data)
 {
-	kernel_next_iteration_setup(kg);
+	kernel_next_iteration_setup((KernelGlobals*)kg);
 }
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_path_init.cl b/intern/cycles/kernel/kernels/opencl/kernel_path_init.cl
index 7e9e4a0252..8194f5d22c 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_path_init.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_path_init.cl
@@ -19,8 +19,8 @@
 #include "split/kernel_path_init.h"
 
 __kernel void kernel_ocl_path_trace_path_init(
-        KernelGlobals *kg,
+        ccl_global char *kg,
         ccl_constant KernelData *data)
 {
-	kernel_path_init(kg);
+	kernel_path_init((KernelGlobals*)kg);
 }
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_queue_enqueue.cl b/intern/cycles/kernel/kernels/opencl/kernel_queue_enqueue.cl
index 9ceb6a5c3d..19074db1b8 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_queue_enqueue.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_queue_enqueue.cl
@@ -19,8 +19,8 @@
 #include "split/kernel_queue_enqueue.h"
 
 __kernel void kernel_ocl_path_trace_queue_enqueue(
-        KernelGlobals *kg,
+        ccl_global char *kg,
         ccl_constant KernelData *data)
 {
-	kernel_queue_enqueue(kg);
+	kernel_queue_enqueue((KernelGlobals*)kg);
 }
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl b/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl
index 4e083e87d1..c675640c59 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl
@@ -19,8 +19,8 @@
 #include "split/kernel_scene_intersect.h"
 
 __kernel void kernel_ocl_path_trace_scene_intersect(
-        KernelGlobals *kg,
+        ccl_global char *kg,
         ccl_constant KernelData *data)
 {
-	kernel_scene_intersect(kg);
+	kernel_scene_intersect((KernelGlobals*)kg);
 }
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl b/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl
index a2b48b1592..534d37f695 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl
@@ -19,8 +19,8 @@
 #include "split/kernel_shader_eval.h"
 
 __kernel void kernel_ocl_path_trace_shader_eval(
-        KernelGlobals *kg,
+        ccl_global char *kg,
         ccl_constant KernelData *data)
 {
-	kernel_shader_eval(kg);
+	kernel_shader_eval((KernelGlobals*)kg);
 }
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_ao.cl b/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_ao.cl
index 1c96d67fec..3782409703 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_ao.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_ao.cl
@@ -19,8 +19,8 @@
 #include "split/kernel_shadow_blocked_ao.h"
 
 __kernel void kernel_ocl_path_trace_shadow_blocked_ao(
-        KernelGlobals *kg,
+        ccl_global char *kg,
         ccl_constant KernelData *data)
 {
-	kernel_shadow_blocked_ao(kg);
+	kernel_shadow_blocked_ao((KernelGlobals*)kg);
 }
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_dl.cl b/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_dl.cl
index 2231f767c0..4889f49d8d 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_dl.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked_dl.cl
@@ -19,8 +19,8 @@
 #include "split/kernel_shadow_blocked_dl.h"
 
 __kernel void kernel_ocl_path_trace_shadow_blocked_dl(
-        KernelGlobals *kg,
+        ccl_global char *kg,
         ccl_constant KernelData *data)
 {
-	kernel_shadow_blocked_dl(kg);
+	kernel_shadow_blocked_dl((KernelGlobals*)kg);
 }
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_state_buffer_size.cl b/intern/cycles/kernel/kernels/opencl/kernel_state_buffer_size.cl
index 4c9bf63ef5..b23ff33786 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_state_buffer_size.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_state_buffer_size.cl
@@ -18,12 +18,12 @@
 #include "split/kernel_split_common.h"
 
 __kernel void kernel_ocl_path_trace_state_buffer_size(
-        KernelGlobals *kg,
+        ccl_global char *kg,
         ccl_constant KernelData *data,
         uint num_threads,
         ccl_global uint64_t *size)
 {
-	kg->data = data;

@@ Diff output truncated at 10240 characters. @@

More information about the Bf-blender-cvs mailing list