03886c0a030756417dd5bd771cdac13b27cd1d27
[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 *globals,
21         ccl_constant KernelData *data,
22         ccl_global char *shader_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         Intersection *Intersection_coop_AO,
27         Intersection *Intersection_coop_DL,
28         ccl_global char *ray_state,
29         ccl_global int *Queue_data,            /* Queue memory */
30         ccl_global int *Queue_index,           /* Tracks the number of elements in each queue */
31         int queuesize,                         /* Size (capacity) of each queue */
32         int total_num_rays)
33 {
34 #if 0
35         /* We will make the Queue_index entries '0' in the next kernel. */
36         if(get_global_id(0) == 0 && get_global_id(1) == 0) {
37                 /* We empty this queue here */
38                 Queue_index[QUEUE_SHADOW_RAY_CAST_AO_RAYS] = 0;
39                 Queue_index[QUEUE_SHADOW_RAY_CAST_DL_RAYS] = 0;
40         }
41 #endif
42
43         int lidx = get_local_id(1) * get_local_id(0) + get_local_id(0);
44
45         ccl_local unsigned int ao_queue_length;
46         ccl_local unsigned int dl_queue_length;
47         if(lidx == 0) {
48                 ao_queue_length = Queue_index[QUEUE_SHADOW_RAY_CAST_AO_RAYS];
49                 dl_queue_length = Queue_index[QUEUE_SHADOW_RAY_CAST_DL_RAYS];
50         }
51         barrier(CLK_LOCAL_MEM_FENCE);
52
53         /* flag determining if the current ray is to process shadow ray for AO or DL */
54         char shadow_blocked_type = -1;
55
56         int ray_index = QUEUE_EMPTY_SLOT;
57         int thread_index = get_global_id(1) * get_global_size(0) + get_global_id(0);
58         if(thread_index < ao_queue_length + dl_queue_length) {
59                 if(thread_index < ao_queue_length) {
60                         ray_index = get_ray_index(thread_index, QUEUE_SHADOW_RAY_CAST_AO_RAYS, Queue_data, queuesize, 1);
61                         shadow_blocked_type = RAY_SHADOW_RAY_CAST_AO;
62                 } else {
63                         ray_index = get_ray_index(thread_index - ao_queue_length, QUEUE_SHADOW_RAY_CAST_DL_RAYS, Queue_data, queuesize, 1);
64                         shadow_blocked_type = RAY_SHADOW_RAY_CAST_DL;
65                 }
66         }
67
68         if(ray_index == QUEUE_EMPTY_SLOT)
69                 return;
70
71         kernel_shadow_blocked(globals,
72                               data,
73                               shader_shadow,
74                               PathState_coop,
75                               LightRay_dl_coop,
76                               LightRay_ao_coop,
77                               Intersection_coop_AO,
78                               Intersection_coop_DL,
79                               ray_state,
80                               total_num_rays,
81                               shadow_blocked_type,
82                               ray_index);
83 }