Cycles: OpenCL split kernel cleanup, move casts from .h files to .cl files
[blender.git] / intern / cycles / kernel / kernels / opencl / kernel_next_iteration_setup.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_next_iteration_setup.h"
18
19 __kernel void kernel_ocl_path_trace_next_iteration_setup(
20         ccl_global char *kg,
21         ccl_constant KernelData *data,
22         ccl_global char *sd,                  /* Required for setting up ray for next iteration */
23         ccl_global uint *rng_coop,            /* Required for setting up ray for next iteration */
24         ccl_global float3 *throughput_coop,   /* Required for setting up ray for next iteration */
25         PathRadiance *PathRadiance_coop,      /* Required for setting up ray for next iteration */
26         ccl_global Ray *Ray_coop,             /* Required for setting up ray for next iteration */
27         ccl_global PathState *PathState_coop, /* Required for setting up ray for next iteration */
28         ccl_global Ray *LightRay_dl_coop,     /* Required for radiance update - direct lighting */
29         ccl_global int *ISLamp_coop,          /* Required for radiance update - direct lighting */
30         ccl_global BsdfEval *BSDFEval_coop,   /* Required for radiance update - direct lighting */
31         ccl_global Ray *LightRay_ao_coop,     /* Required for radiance update - AO */
32         ccl_global float3 *AOBSDF_coop,       /* Required for radiance update - AO */
33         ccl_global float3 *AOAlpha_coop,      /* Required for radiance update - AO */
34         ccl_global char *ray_state,           /* Denotes the state of each ray */
35         ccl_global int *Queue_data,           /* Queue memory */
36         ccl_global int *Queue_index,          /* Tracks the number of elements in each queue */
37         int queuesize,                        /* Size (capacity) of each queue */
38         ccl_global char *use_queues_flag)     /* flag to decide if scene_intersect kernel should
39                                                * use queues to fetch ray index */
40 {
41         ccl_local unsigned int local_queue_atomics;
42         if(get_local_id(0) == 0 && get_local_id(1) == 0) {
43                 local_queue_atomics = 0;
44         }
45         barrier(CLK_LOCAL_MEM_FENCE);
46
47         if(get_global_id(0) == 0 && get_global_id(1) == 0) {
48                 /* If we are here, then it means that scene-intersect kernel
49                 * has already been executed atleast once. From the next time,
50                 * scene-intersect kernel may operate on queues to fetch ray index
51                 */
52                 use_queues_flag[0] = 1;
53
54                 /* Mark queue indices of QUEUE_SHADOW_RAY_CAST_AO_RAYS and
55                  * QUEUE_SHADOW_RAY_CAST_DL_RAYS queues that were made empty during the
56                  * previous kernel.
57                  */
58                 Queue_index[QUEUE_SHADOW_RAY_CAST_AO_RAYS] = 0;
59                 Queue_index[QUEUE_SHADOW_RAY_CAST_DL_RAYS] = 0;
60         }
61
62         char enqueue_flag = 0;
63         int ray_index = get_global_id(1) * get_global_size(0) + get_global_id(0);
64         ray_index = get_ray_index(ray_index,
65                                   QUEUE_ACTIVE_AND_REGENERATED_RAYS,
66                                   Queue_data,
67                                   queuesize,
68                                   0);
69
70 #ifdef __COMPUTE_DEVICE_GPU__
71         /* If we are executing on a GPU device, we exit all threads that are not
72          * required.
73          *
74          * If we are executing on a CPU device, then we need to keep all threads
75          * active since we have barrier() calls later in the kernel. CPU devices,
76          * expect all threads to execute barrier statement.
77          */
78         if(ray_index == QUEUE_EMPTY_SLOT) {
79                 return;
80         }
81 #endif
82
83 #ifndef __COMPUTE_DEVICE_GPU__
84         if(ray_index != QUEUE_EMPTY_SLOT) {
85 #endif
86                 enqueue_flag = kernel_next_iteration_setup((KernelGlobals *)kg,
87                                                            data,
88                                                            (ShaderData *)sd,
89                                                            rng_coop,
90                                                            throughput_coop,
91                                                            PathRadiance_coop,
92                                                            Ray_coop,
93                                                            PathState_coop,
94                                                            LightRay_dl_coop,
95                                                            ISLamp_coop,
96                                                            BSDFEval_coop,
97                                                            LightRay_ao_coop,
98                                                            AOBSDF_coop,
99                                                            AOAlpha_coop,
100                                                            ray_state,
101                                                            use_queues_flag,
102                                                            ray_index);
103 #ifndef __COMPUTE_DEVICE_GPU__
104         }
105 #endif
106
107         /* Enqueue RAY_UPDATE_BUFFER rays. */
108         enqueue_ray_index_local(ray_index,
109                                 QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS,
110                                 enqueue_flag,
111                                 queuesize,
112                                 &local_queue_atomics,
113                                 Queue_data,
114                                 Queue_index);
115 }