[Bf-blender-cvs] [a3ef51b] master: Fix T44833, OpenCL compile error on AMD.

Thomas Dinges noreply at git.blender.org
Mon May 25 01:05:42 CEST 2015 

Commit: a3ef51bba5f0494f5ff87acf52222c3f48eb3684
Author: Thomas Dinges
Date:   Mon May 25 01:02:06 2015 +0200
Branches: master
https://developer.blender.org/rBa3ef51bba5f0494f5ff87acf52222c3f48eb3684

Fix T44833, OpenCL compile error on AMD.

This was broken after the kernel file restructure.
Variables allocated in the __local address space can only be defined
inside a __kernel function.

We probably need to solve this a bit differently once we do the CUDA
kernel split, but this fix shoud be good enough until then.

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

M	intern/cycles/kernel/split/kernel_background_buffer_update.h
M	intern/cycles/kernel/split/kernel_data_init.h
M	intern/cycles/kernel/split/kernel_direct_lighting.h
M	intern/cycles/kernel/split/kernel_holdout_emission_blurring_pathtermination_ao.h
M	intern/cycles/kernel/split/kernel_lamp_emission.h
M	intern/cycles/kernel/split/kernel_next_iteration_setup.h
M	intern/cycles/kernel/split/kernel_queue_enqueue.h
M	intern/cycles/kernel/split/kernel_scene_intersect.h
M	intern/cycles/kernel/split/kernel_shader_eval.h
M	intern/cycles/kernel/split/kernel_shadow_blocked.h
M	intern/cycles/kernel/split/kernel_sum_all_radiance.h

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

diff --git a/intern/cycles/kernel/split/kernel_background_buffer_update.h b/intern/cycles/kernel/split/kernel_background_buffer_update.h
index 95de1a4..32c7e6f 100644
--- a/intern/cycles/kernel/split/kernel_background_buffer_update.h
+++ b/intern/cycles/kernel/split/kernel_background_buffer_update.h
@@ -70,7 +70,7 @@
  * QUEUE_ACTIVE_AND_REGENERATED_RAYS will be filled with RAY_ACTIVE and RAY_REGENERATED rays
  * QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS will be empty
  */
-ccl_device void kernel_background_buffer_update(
+__kernel void kernel_background_buffer_update(
 	ccl_global char *globals,
 	ccl_constant KernelData *data,
 	ccl_global char *shader_data,
diff --git a/intern/cycles/kernel/split/kernel_data_init.h b/intern/cycles/kernel/split/kernel_data_init.h
index b7a4d84..006f2c8 100644
--- a/intern/cycles/kernel/split/kernel_data_init.h
+++ b/intern/cycles/kernel/split/kernel_data_init.h
@@ -51,7 +51,7 @@
  * All slots in queues are initialized to queue empty slot;
  * The number of elements in the queues is initialized to 0;
  */
-ccl_device void kernel_data_init(
+__kernel void kernel_data_init(
 	ccl_global char *globals,
 	ccl_global char *shader_data_sd,                  /* Arguments related to ShaderData */
 	ccl_global char *shader_data_sd_DL_shadow,     /* Arguments related to ShaderData */
diff --git a/intern/cycles/kernel/split/kernel_direct_lighting.h b/intern/cycles/kernel/split/kernel_direct_lighting.h
index 6b83d89..91c3ef1 100644
--- a/intern/cycles/kernel/split/kernel_direct_lighting.h
+++ b/intern/cycles/kernel/split/kernel_direct_lighting.h
@@ -49,7 +49,7 @@
  * QUEUE_SHADOW_RAY_CAST_DL_RAYS queue will be filled with rays for which a shadow_blocked function must be executed, after this
  * kernel call. Before this kernel call the QUEUE_SHADOW_RAY_CAST_DL_RAYS will be empty.
  */
-ccl_device void kernel_direct_lighting(
+__kernel void kernel_direct_lighting(
 	ccl_global char *globals,
 	ccl_constant KernelData *data,
 	ccl_global char *shader_data,           /* Required for direct lighting */
diff --git a/intern/cycles/kernel/split/kernel_holdout_emission_blurring_pathtermination_ao.h b/intern/cycles/kernel/split/kernel_holdout_emission_blurring_pathtermination_ao.h
index 393ea4b..174070a 100644
--- a/intern/cycles/kernel/split/kernel_holdout_emission_blurring_pathtermination_ao.h
+++ b/intern/cycles/kernel/split/kernel_holdout_emission_blurring_pathtermination_ao.h
@@ -72,7 +72,7 @@
  * QUEUE_SHADOW_RAY_CAST_AO_RAYS will be filled with rays marked with flag RAY_SHADOW_RAY_CAST_AO
  */
 
-ccl_device void kernel_holdout_emission_blurring_pathtermination_ao(
+__kernel void kernel_holdout_emission_blurring_pathtermination_ao(
 	ccl_global char *globals,
 	ccl_constant KernelData *data,
 	ccl_global char *shader_data,               /* Required throughout the kernel except probabilistic path termination and AO */
diff --git a/intern/cycles/kernel/split/kernel_lamp_emission.h b/intern/cycles/kernel/split/kernel_lamp_emission.h
index f400a99..b804bfc 100644
--- a/intern/cycles/kernel/split/kernel_lamp_emission.h
+++ b/intern/cycles/kernel/split/kernel_lamp_emission.h
@@ -40,7 +40,7 @@
  *
  * note : shader_data is neither input nor output. Its just filled and consumed in the same, kernel_lamp_emission, kernel.
  */
-ccl_device void kernel_lamp_emission(
+__kernel void kernel_lamp_emission(
 	ccl_global char *globals,
 	ccl_constant KernelData *data,
 	ccl_global char *shader_data,               /* Required for lamp emission */
diff --git a/intern/cycles/kernel/split/kernel_next_iteration_setup.h b/intern/cycles/kernel/split/kernel_next_iteration_setup.h
index 343dbb0..6ce56e4 100644
--- a/intern/cycles/kernel/split/kernel_next_iteration_setup.h
+++ b/intern/cycles/kernel/split/kernel_next_iteration_setup.h
@@ -61,7 +61,7 @@
  * QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS will be filled with RAY_TO_REGENERATE and more RAY_UPDATE_BUFFER rays
  */
 
-ccl_device void kernel_next_iteration_setup(
+__kernel void kernel_next_iteration_setup(
 	ccl_global char *globals,
 	ccl_constant KernelData *data,
 	ccl_global char *shader_data,               /* Required for setting up ray for next iteration */
diff --git a/intern/cycles/kernel/split/kernel_queue_enqueue.h b/intern/cycles/kernel/split/kernel_queue_enqueue.h
index 9bcf8f5..5a9838f 100644
--- a/intern/cycles/kernel/split/kernel_queue_enqueue.h
+++ b/intern/cycles/kernel/split/kernel_queue_enqueue.h
@@ -52,7 +52,7 @@
  * QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS will be filled with RAY_TO_REGENERATE, RAY_UPDATE_BUFFER, RAY_HIT_BACKGROUND rays.
  */
 
-ccl_device void kernel_queue_enqueue(
+__kernel void kernel_queue_enqueue(
 	ccl_global int *Queue_data,   /* Queue memory */
 	ccl_global int *Queue_index,  /* Tracks the number of elements in each queue */
 	ccl_global char *ray_state,   /* Denotes the state of each ray */
diff --git a/intern/cycles/kernel/split/kernel_scene_intersect.h b/intern/cycles/kernel/split/kernel_scene_intersect.h
index 01e0b1f..6cb3cdc 100644
--- a/intern/cycles/kernel/split/kernel_scene_intersect.h
+++ b/intern/cycles/kernel/split/kernel_scene_intersect.h
@@ -63,7 +63,7 @@
  * QUEUE_HITBF_BUFF_UPDATE_TOREGEN_RAYS - no change
  */
 
-ccl_device void kernel_scene_intersect(
+__kernel void kernel_scene_intersect(
 	ccl_global char *globals,
 	ccl_constant KernelData *data,
 	ccl_global uint *rng_coop,
diff --git a/intern/cycles/kernel/split/kernel_shader_eval.h b/intern/cycles/kernel/split/kernel_shader_eval.h
index 0a8d77f..924c7fa 100644
--- a/intern/cycles/kernel/split/kernel_shader_eval.h
+++ b/intern/cycles/kernel/split/kernel_shader_eval.h
@@ -46,7 +46,7 @@
  * QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS will be filled with RAY_TO_REGENERATE rays
  */
 
-ccl_device void kernel_shader_eval(
+__kernel void kernel_shader_eval(
 	ccl_global char *globals,
 	ccl_constant KernelData *data,
 	ccl_global char *shader_data,               /* Output ShaderData structure to be filled */
diff --git a/intern/cycles/kernel/split/kernel_shadow_blocked.h b/intern/cycles/kernel/split/kernel_shadow_blocked.h
index 71fab19..52bd9eb 100644
--- a/intern/cycles/kernel/split/kernel_shadow_blocked.h
+++ b/intern/cycles/kernel/split/kernel_shadow_blocked.h
@@ -47,7 +47,7 @@
  * QUEUE_SHADOW_RAY_CAST_AO_RAYS and QUEUE_SHADOW_RAY_CAST_DL_RAYS will be empty at kernel exit.
  */
 
-ccl_device void kernel_shadow_blocked(
+__kernel void kernel_shadow_blocked(
 	ccl_global char *globals,
 	ccl_constant KernelData *data,
 	ccl_global char *shader_shadow,             /* Required for shadow blocked */
diff --git a/intern/cycles/kernel/split/kernel_sum_all_radiance.h b/intern/cycles/kernel/split/kernel_sum_all_radiance.h
index eeb7da7..faa4162 100644
--- a/intern/cycles/kernel/split/kernel_sum_all_radiance.h
+++ b/intern/cycles/kernel/split/kernel_sum_all_radiance.h
@@ -25,7 +25,7 @@
 * by all different samples and stores them in the RenderTile's output buffer.
 */
 
-ccl_device void kernel_sum_all_radiance(
+__kernel void kernel_sum_all_radiance(
 	ccl_constant KernelData *data,               /* To get pass_stride to offet into buffer */
 	ccl_global float *buffer,                    /* Output buffer of RenderTile */
 	ccl_global float *per_sample_output_buffer,  /* Radiance contributed by all samples */

More information about the Bf-blender-cvs mailing list