Cycles: Remove few function arguments needed only for the split kernel
[blender.git] / intern / cycles / kernel / kernels / opencl / kernel_shadow_blocked.cl
1 /*
2  * Copyright 2011-2015 Blender Foundation
3  *
4  * Licensed under the Apache License, Version 2.0 (the "License");
5  * you may not use this file except in compliance with the License.
6  * You may obtain a copy of the License at
7  *
8  * http://www.apache.org/licenses/LICENSE-2.0
9  *
10  * Unless required by applicable law or agreed to in writing, software
11  * distributed under the License is distributed on an "AS IS" BASIS,
12  * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13  * See the License for the specific language governing permissions and
14  * limitations under the License.
15  */
16
17 #include "split/kernel_shadow_blocked.h"
18
19 __kernel void kernel_ocl_path_trace_shadow_blocked(
20         ccl_global char *kg,
21         ccl_constant KernelData *data,
22         ccl_global char *sd_shadow,            /* Required for shadow blocked */
23         ccl_global PathState *PathState_coop,  /* Required for shadow blocked */
24         ccl_global Ray *LightRay_dl_coop,      /* Required for direct lighting's shadow blocked */
25         ccl_global Ray *LightRay_ao_coop,      /* Required for AO's shadow blocked */
26         ccl_global char *ray_state,
27         ccl_global int *Queue_data,            /* Queue memory */
28         ccl_global int *Queue_index,           /* Tracks the number of elements in each queue */
29         int queuesize,                         /* Size (capacity) of each queue */
30         int total_num_rays)
31 {
32 #if 0
33         /* We will make the Queue_index entries '0' in the next kernel. */
34         if(get_global_id(0) == 0 && get_global_id(1) == 0) {
35                 /* We empty this queue here */
36                 Queue_index[QUEUE_SHADOW_RAY_CAST_AO_RAYS] = 0;
37                 Queue_index[QUEUE_SHADOW_RAY_CAST_DL_RAYS] = 0;
38         }
39 #endif
40
41         int lidx = get_local_id(1) * get_local_id(0) + get_local_id(0);
42
43         ccl_local unsigned int ao_queue_length;
44         ccl_local unsigned int dl_queue_length;
45         if(lidx == 0) {
46                 ao_queue_length = Queue_index[QUEUE_SHADOW_RAY_CAST_AO_RAYS];
47                 dl_queue_length = Queue_index[QUEUE_SHADOW_RAY_CAST_DL_RAYS];
48         }
49         barrier(CLK_LOCAL_MEM_FENCE);
50
51         /* flag determining if the current ray is to process shadow ray for AO or DL */
52         char shadow_blocked_type = -1;
53
54         int ray_index = QUEUE_EMPTY_SLOT;
55         int thread_index = get_global_id(1) * get_global_size(0) + get_global_id(0);
56         if(thread_index < ao_queue_length + dl_queue_length) {
57                 if(thread_index < ao_queue_length) {
58                         ray_index = get_ray_index(thread_index, QUEUE_SHADOW_RAY_CAST_AO_RAYS, Queue_data, queuesize, 1);
59                         shadow_blocked_type = RAY_SHADOW_RAY_CAST_AO;
60                 } else {
61                         ray_index = get_ray_index(thread_index - ao_queue_length, QUEUE_SHADOW_RAY_CAST_DL_RAYS, Queue_data, queuesize, 1);
62                         shadow_blocked_type = RAY_SHADOW_RAY_CAST_DL;
63                 }
64         }
65
66         if(ray_index == QUEUE_EMPTY_SLOT)
67                 return;
68
69         kernel_shadow_blocked((KernelGlobals *)kg,
70                               (ShaderData *)sd_shadow,
71                               PathState_coop,
72                               LightRay_dl_coop,
73                               LightRay_ao_coop,
74                               ray_state,
75                               total_num_rays,
76                               shadow_blocked_type,
77                               ray_index);
78 }