sandey/blender
1
0
Fork 0
forked from bartvdbraak/blender
Code Pull Requests Activity

Fix T44833, OpenCL compile error on AMD.

Browse Source 
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.
This commit is contained in:
Thomas Dinges 2015-05-25 01:02:06 +02:00
parent c3ab5b3089
commit a3ef51bba5
11 changed files with 11 additions and 11 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
intern/cycles/kernel/split/kernel_background_buffer_update.h
View File

@ -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,

2
intern/cycles/kernel/split/kernel_data_init.h
View File

@ -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 */

2
intern/cycles/kernel/split/kernel_direct_lighting.h
View File

@ -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 */

2
intern/cycles/kernel/split/kernel_holdout_emission_blurring_pathtermination_ao.h
View File

@ -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 */

2
intern/cycles/kernel/split/kernel_lamp_emission.h
View File

@ -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 */

2
intern/cycles/kernel/split/kernel_next_iteration_setup.h
View File

@ -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 */

2
intern/cycles/kernel/split/kernel_queue_enqueue.h
View File

@ -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 */

2
intern/cycles/kernel/split/kernel_scene_intersect.h
View File

@ -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,

2
intern/cycles/kernel/split/kernel_shader_eval.h
View File

@ -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 */

2
intern/cycles/kernel/split/kernel_shadow_blocked.h
View File

@ -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 */

2
intern/cycles/kernel/split/kernel_sum_all_radiance.h
View File

@ -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 */