Blender  V2.93
kernel_next_iteration_setup.h
Go to the documentation of this file.
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 
18 
19 /*This kernel takes care of setting up ray for the next iteration of
20  * path-iteration and accumulating radiance corresponding to AO and
21  * direct-lighting
22  *
23  * Ray state of rays that are terminated in this kernel are changed
24  * to RAY_UPDATE_BUFFER.
25  *
26  * Note on queues:
27  * This kernel fetches rays from the queue QUEUE_ACTIVE_AND_REGENERATED_RAYS
28  * and processes only the rays of state RAY_ACTIVE.
29  * There are different points in this kernel where a ray may terminate and
30  * reach RAY_UPDATE_BUFF state. These rays are enqueued into
31  * QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS queue. These rays will still be present
32  * in QUEUE_ACTIVE_AND_REGENERATED_RAYS queue, but since their ray-state has
33  * been changed to RAY_UPDATE_BUFF, there is no problem.
34  *
35  * State of queues when this kernel is called:
36  * At entry,
37  * - QUEUE_ACTIVE_AND_REGENERATED_RAYS will be filled with RAY_ACTIVE,
38  * RAY_REGENERATED, RAY_UPDATE_BUFFER rays.
39  * - QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS will be filled with
40  * RAY_TO_REGENERATE and RAY_UPDATE_BUFFER rays.
41  * At exit,
42  * - QUEUE_ACTIVE_AND_REGENERATED_RAYS will be filled with RAY_ACTIVE,
43  * RAY_REGENERATED and more RAY_UPDATE_BUFFER rays.
44  * - QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS will be filled with
45  * RAY_TO_REGENERATE and more RAY_UPDATE_BUFFER rays.
46  */
47 
48 #ifdef __BRANCHED_PATH__
49 ccl_device_inline void kernel_split_branched_indirect_light_init(KernelGlobals *kg, int ray_index)
50 {
51  kernel_split_branched_path_indirect_loop_init(kg, ray_index);
52 
54 }
55 
56 ccl_device void kernel_split_branched_transparent_bounce(KernelGlobals *kg, int ray_index)
57 {
58  ccl_global float3 *throughput = &kernel_split_state.throughput[ray_index];
59  ShaderData *sd = kernel_split_sd(sd, ray_index);
60  ccl_global PathState *state = &kernel_split_state.path_state[ray_index];
61  ccl_global Ray *ray = &kernel_split_state.ray[ray_index];
62 
63 # ifdef __VOLUME__
64  if (!(sd->flag & SD_HAS_ONLY_VOLUME)) {
65 # endif
66  /* continue in case of transparency */
67  *throughput *= shader_bsdf_transparency(kg, sd);
68 
69  if (is_zero(*throughput)) {
70  kernel_split_path_end(kg, ray_index);
71  return;
72  }
73 
74  /* Update Path State */
76 # ifdef __VOLUME__
77  }
78  else {
79  if (!path_state_volume_next(kg, state)) {
80  kernel_split_path_end(kg, ray_index);
81  return;
82  }
83  }
84 # endif
85 
86  ray->P = ray_offset(sd->P, -sd->Ng);
87  ray->t -= sd->ray_length; /* clipping works through transparent */
88 
89 # ifdef __RAY_DIFFERENTIALS__
90  ray->dP = sd->dP;
91  ray->dD.dx = -sd->dI.dx;
92  ray->dD.dy = -sd->dI.dy;
93 # endif /* __RAY_DIFFERENTIALS__ */
94 
95 # ifdef __VOLUME__
96  /* enter/exit volume */
97  kernel_volume_stack_enter_exit(kg, sd, state->volume_stack);
98 # endif /* __VOLUME__ */
99 }
100 #endif /* __BRANCHED_PATH__ */
101 
103  ccl_local_param unsigned int *local_queue_atomics)
104 {
105  if (ccl_local_id(0) == 0 && ccl_local_id(1) == 0) {
106  *local_queue_atomics = 0;
107  }
109 
110  if (ccl_global_id(0) == 0 && ccl_global_id(1) == 0) {
111  /* If we are here, then it means that scene-intersect kernel
112  * has already been executed at least once. From the next time,
113  * scene-intersect kernel may operate on queues to fetch ray index
114  */
115  *kernel_split_params.use_queues_flag = 1;
116 
117  /* Mark queue indices of QUEUE_SHADOW_RAY_CAST_AO_RAYS and
118  * QUEUE_SHADOW_RAY_CAST_DL_RAYS queues that were made empty during the
119  * previous kernel.
120  */
123  }
124 
125  int ray_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0);
126  ray_index = get_ray_index(kg,
127  ray_index,
129  kernel_split_state.queue_data,
130  kernel_split_params.queue_size,
131  0);
132 
133  ccl_global char *ray_state = kernel_split_state.ray_state;
134 
135 #ifdef __VOLUME__
136  /* Reactivate only volume rays here, most surface work was skipped. */
137  if (IS_STATE(ray_state, ray_index, RAY_HAS_ONLY_VOLUME)) {
139  }
140 #endif
141 
142  bool active = IS_STATE(ray_state, ray_index, RAY_ACTIVE);
143  if (active) {
144  ccl_global float3 *throughput = &kernel_split_state.throughput[ray_index];
145  ccl_global Ray *ray = &kernel_split_state.ray[ray_index];
146  ShaderData *sd = kernel_split_sd(sd, ray_index);
147  ccl_global PathState *state = &kernel_split_state.path_state[ray_index];
148  PathRadiance *L = &kernel_split_state.path_radiance[ray_index];
149 
150 #ifdef __BRANCHED_PATH__
151  if (!kernel_data.integrator.branched || IS_FLAG(ray_state, ray_index, RAY_BRANCHED_INDIRECT)) {
152 #endif
153  /* Compute direct lighting and next bounce. */
154  if (!kernel_path_surface_bounce(kg, sd, throughput, state, &L->state, ray)) {
155  kernel_split_path_end(kg, ray_index);
156  }
157 #ifdef __BRANCHED_PATH__
158  }
159  else if (sd->flag & SD_HAS_ONLY_VOLUME) {
160  kernel_split_branched_transparent_bounce(kg, ray_index);
161  }
162  else {
163  kernel_split_branched_indirect_light_init(kg, ray_index);
164 
165  if (kernel_split_branched_path_surface_indirect_light_iter(
166  kg, ray_index, 1.0f, kernel_split_sd(branched_state_sd, ray_index), true, true)) {
168  }
169  else {
170  kernel_split_branched_path_indirect_loop_end(kg, ray_index);
171  kernel_split_branched_transparent_bounce(kg, ray_index);
172  }
173  }
174 #endif /* __BRANCHED_PATH__ */
175  }
176 
177  /* Enqueue RAY_UPDATE_BUFFER rays. */
178  enqueue_ray_index_local(ray_index,
180  IS_STATE(ray_state, ray_index, RAY_UPDATE_BUFFER) && active,
181  kernel_split_params.queue_size,
182  local_queue_atomics,
183  kernel_split_state.queue_data,
184  kernel_split_params.queue_index);
185 
186 #ifdef __BRANCHED_PATH__
187  /* iter loop */
188  if (ccl_global_id(0) == 0 && ccl_global_id(1) == 0) {
190  }
191 
192  ray_index = get_ray_index(kg,
195  kernel_split_state.queue_data,
196  kernel_split_params.queue_size,
197  1);
198 
200  /* for render passes, sum and reset indirect light pass variables
201  * for the next samples */
202  PathRadiance *L = &kernel_split_state.path_radiance[ray_index];
203 
206 
207  if (kernel_split_branched_path_surface_indirect_light_iter(
208  kg, ray_index, 1.0f, kernel_split_sd(branched_state_sd, ray_index), true, true)) {
210  }
211  else {
212  kernel_split_branched_path_indirect_loop_end(kg, ray_index);
213  kernel_split_branched_transparent_bounce(kg, ray_index);
214  }
215  }
216 
217 # ifdef __VOLUME__
218  /* Enqueue RAY_VOLUME_INDIRECT_NEXT_ITER rays */
220  if (ccl_local_id(0) == 0 && ccl_local_id(1) == 0) {
221  *local_queue_atomics = 0;
222  }
224 
225  ray_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0);
227  ray_index,
230  kernel_split_params.queue_size,
231  local_queue_atomics,
232  kernel_split_state.queue_data,
233  kernel_split_params.queue_index);
234 
235 # endif /* __VOLUME__ */
236 
237 # ifdef __SUBSURFACE__
238  /* Enqueue RAY_SUBSURFACE_INDIRECT_NEXT_ITER rays */
240  if (ccl_local_id(0) == 0 && ccl_local_id(1) == 0) {
241  *local_queue_atomics = 0;
242  }
244 
245  ray_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0);
247  ray_index,
250  kernel_split_params.queue_size,
251  local_queue_atomics,
252  kernel_split_state.queue_data,
253  kernel_split_params.queue_index);
254 # endif /* __SUBSURFACE__ */
255 #endif /* __BRANCHED_PATH__ */
256 }
257 
ccl_device_inline float3 ray_offset(float3 P, float3 Ng)
ccl_device_inline void path_radiance_sum_indirect(PathRadiance *L)
CCL_NAMESPACE_BEGIN ccl_device float3 shader_bsdf_transparency(KernelGlobals *kg, const ShaderData *sd)
ccl_device_inline void path_radiance_reset_indirect(PathRadiance *L)
#define kernel_data
#define ccl_global_id(d)
#define ccl_global_size(d)
#define ccl_local_param
#define ccl_device
ccl_device_inline uint ccl_local_id(uint d)
#define ccl_device_inline
#define ccl_global
#define CCL_NAMESPACE_END
CCL_NAMESPACE_BEGIN ccl_device void kernel_next_iteration_setup(KernelGlobals *kg, ccl_local_param unsigned int *local_queue_atomics)
ccl_device_inline void path_state_next(KernelGlobals *kg, ccl_addr_space PathState *state, int label)
ccl_device bool kernel_path_surface_bounce(KernelGlobals *kg, ShaderData *sd, ccl_addr_space float3 *throughput, ccl_addr_space PathState *state, PathRadianceState *L_state, ccl_addr_space Ray *ray)
ccl_device int get_ray_index(KernelGlobals *kg, int thread_index, int queue_number, ccl_global int *queues, int queuesize, int empty_queue)
Definition: kernel_queues.h:53
ccl_device void enqueue_ray_index_local(int ray_index, int queue_number, char enqueue_flag, int queuesize, ccl_local_param unsigned int *local_queue_atomics, ccl_global int *Queue_data, ccl_global int *Queue_index)
Definition: kernel_queues.h:71
CCL_NAMESPACE_BEGIN ccl_device_inline void kernel_split_path_end(KernelGlobals *kg, int ray_index)
#define kernel_split_params
#define kernel_split_sd(sd, ray_index)
#define kernel_split_state
__kernel void ccl_constant KernelData ccl_global void ccl_global char * ray_state
#define IS_FLAG(ray_state, ray_index, flag)
@ SD_HAS_ONLY_VOLUME
Definition: kernel_types.h:875
#define IS_STATE(ray_state, ray_index, state)
ShaderData
#define ADD_RAY_FLAG(ray_state, ray_index, flag)
@ LABEL_TRANSPARENT
Definition: kernel_types.h:333
#define ASSIGN_RAY_STATE(ray_state, ray_index, state)
@ RAY_BRANCHED_INDIRECT
@ RAY_UPDATE_BUFFER
@ RAY_ACTIVE
@ RAY_LIGHT_INDIRECT_NEXT_ITER
@ RAY_BRANCHED_LIGHT_INDIRECT
@ RAY_HAS_ONLY_VOLUME
@ RAY_SUBSURFACE_INDIRECT_NEXT_ITER
@ RAY_VOLUME_INDIRECT_NEXT_ITER
@ RAY_REGENERATED
@ QUEUE_LIGHT_INDIRECT_ITER
@ QUEUE_SHADOW_RAY_CAST_DL_RAYS
@ QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS
@ QUEUE_SHADOW_RAY_CAST_AO_RAYS
@ QUEUE_VOLUME_INDIRECT_ITER
@ QUEUE_SUBSURFACE_INDIRECT_ITER
@ QUEUE_ACTIVE_AND_REGENERATED_RAYS
static ulong state[N]
#define L
bool active
all scheduled work for the GPU.
#define CCL_LOCAL_MEM_FENCE
Definition: util_atomic.h:32
#define ccl_barrier(flags)
Definition: util_atomic.h:33
ccl_device_inline bool is_zero(const float2 &a)