Blender  V2.93
kernel_branched.h
Go to the documentation of this file.
1 /*
2  * Copyright 2011-2017 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 #ifdef __BRANCHED_PATH__
20 
21 /* sets up the various state needed to do an indirect loop */
22 ccl_device_inline void kernel_split_branched_path_indirect_loop_init(KernelGlobals *kg,
23  int ray_index)
24 {
25  SplitBranchedState *branched_state = &kernel_split_state.branched_state[ray_index];
26 
27  /* save a copy of the state to restore later */
28 # define BRANCHED_STORE(name) branched_state->name = kernel_split_state.name[ray_index];
29 
30  BRANCHED_STORE(path_state);
31  BRANCHED_STORE(throughput);
32  BRANCHED_STORE(ray);
33  BRANCHED_STORE(isect);
34  BRANCHED_STORE(ray_state);
35 
36  *kernel_split_sd(branched_state_sd, ray_index) = *kernel_split_sd(sd, ray_index);
37  for (int i = 0; i < kernel_split_sd(branched_state_sd, ray_index)->num_closure; i++) {
38  kernel_split_sd(branched_state_sd, ray_index)->closure[i] =
39  kernel_split_sd(sd, ray_index)->closure[i];
40  }
41 
42 # undef BRANCHED_STORE
43 
44  /* Set loop counters to initial position. */
45  branched_state->next_closure = 0;
46  branched_state->next_sample = 0;
47 }
48 
49 /* ends an indirect loop and restores the previous state */
50 ccl_device_inline void kernel_split_branched_path_indirect_loop_end(KernelGlobals *kg,
51  int ray_index)
52 {
53  SplitBranchedState *branched_state = &kernel_split_state.branched_state[ray_index];
54 
55  /* restore state */
56 # define BRANCHED_RESTORE(name) kernel_split_state.name[ray_index] = branched_state->name;
57 
58  BRANCHED_RESTORE(path_state);
59  BRANCHED_RESTORE(throughput);
60  BRANCHED_RESTORE(ray);
61  BRANCHED_RESTORE(isect);
62  BRANCHED_RESTORE(ray_state);
63 
64  *kernel_split_sd(sd, ray_index) = *kernel_split_sd(branched_state_sd, ray_index);
65  for (int i = 0; i < kernel_split_sd(branched_state_sd, ray_index)->num_closure; i++) {
66  kernel_split_sd(sd, ray_index)->closure[i] =
67  kernel_split_sd(branched_state_sd, ray_index)->closure[i];
68  }
69 
70 # undef BRANCHED_RESTORE
71 
72  /* leave indirect loop */
74 }
75 
76 ccl_device_inline bool kernel_split_branched_indirect_start_shared(KernelGlobals *kg,
77  int ray_index)
78 {
79  ccl_global char *ray_state = kernel_split_state.ray_state;
80 
81  int inactive_ray = dequeue_ray_index(QUEUE_INACTIVE_RAYS,
82  kernel_split_state.queue_data,
83  kernel_split_params.queue_size,
84  kernel_split_params.queue_index);
85 
86  if (!IS_STATE(ray_state, inactive_ray, RAY_INACTIVE)) {
87  return false;
88  }
89 
90 # define SPLIT_DATA_ENTRY(type, name, num) \
91  if (num) { \
92  kernel_split_state.name[inactive_ray] = kernel_split_state.name[ray_index]; \
93  }
95 # undef SPLIT_DATA_ENTRY
96 
97  *kernel_split_sd(sd, inactive_ray) = *kernel_split_sd(sd, ray_index);
98  for (int i = 0; i < kernel_split_sd(sd, ray_index)->num_closure; i++) {
99  kernel_split_sd(sd, inactive_ray)->closure[i] = kernel_split_sd(sd, ray_index)->closure[i];
100  }
101 
102  kernel_split_state.branched_state[inactive_ray].shared_sample_count = 0;
103  kernel_split_state.branched_state[inactive_ray].original_ray = ray_index;
104  kernel_split_state.branched_state[inactive_ray].waiting_on_shared_samples = false;
105 
106  PathRadiance *L = &kernel_split_state.path_radiance[ray_index];
107  PathRadiance *inactive_L = &kernel_split_state.path_radiance[inactive_ray];
108 
109  path_radiance_init(kg, inactive_L);
110  path_radiance_copy_indirect(inactive_L, L);
111 
112  ray_state[inactive_ray] = RAY_REGENERATED;
114  ADD_RAY_FLAG(ray_state, inactive_ray, IS_FLAG(ray_state, ray_index, RAY_BRANCHED_INDIRECT));
115 
117  (ccl_global uint *)&kernel_split_state.branched_state[ray_index].shared_sample_count);
118 
119  return true;
120 }
121 
122 /* bounce off surface and integrate indirect light */
123 ccl_device_noinline bool kernel_split_branched_path_surface_indirect_light_iter(
124  KernelGlobals *kg,
125  int ray_index,
126  float num_samples_adjust,
127  ShaderData *saved_sd,
128  bool reset_path_state,
129  bool wait_for_shared)
130 {
131  SplitBranchedState *branched_state = &kernel_split_state.branched_state[ray_index];
132 
133  ShaderData *sd = saved_sd;
134  PathRadiance *L = &kernel_split_state.path_radiance[ray_index];
135  float3 throughput = branched_state->throughput;
136  ccl_global PathState *ps = &kernel_split_state.path_state[ray_index];
137 
138  float sum_sample_weight = 0.0f;
139 # ifdef __DENOISING_FEATURES__
140  if (ps->denoising_feature_weight > 0.0f) {
141  for (int i = 0; i < sd->num_closure; i++) {
142  const ShaderClosure *sc = &sd->closure[i];
143 
144  /* transparency is not handled here, but in outer loop */
145  if (!CLOSURE_IS_BSDF(sc->type) || CLOSURE_IS_BSDF_TRANSPARENT(sc->type)) {
146  continue;
147  }
148 
149  sum_sample_weight += sc->sample_weight;
150  }
151  }
152  else {
153  sum_sample_weight = 1.0f;
154  }
155 # endif /* __DENOISING_FEATURES__ */
156 
157  for (int i = branched_state->next_closure; i < sd->num_closure; i++) {
158  const ShaderClosure *sc = &sd->closure[i];
159 
160  if (!CLOSURE_IS_BSDF(sc->type))
161  continue;
162  /* transparency is not handled here, but in outer loop */
163  if (sc->type == CLOSURE_BSDF_TRANSPARENT_ID)
164  continue;
165 
166  int num_samples;
167 
168  if (CLOSURE_IS_BSDF_DIFFUSE(sc->type))
169  num_samples = kernel_data.integrator.diffuse_samples;
170  else if (CLOSURE_IS_BSDF_BSSRDF(sc->type))
171  num_samples = 1;
172  else if (CLOSURE_IS_BSDF_GLOSSY(sc->type))
173  num_samples = kernel_data.integrator.glossy_samples;
174  else
175  num_samples = kernel_data.integrator.transmission_samples;
176 
177  num_samples = ceil_to_int(num_samples_adjust * num_samples);
178 
179  float num_samples_inv = num_samples_adjust / num_samples;
180 
181  for (int j = branched_state->next_sample; j < num_samples; j++) {
182  if (reset_path_state) {
183  *ps = branched_state->path_state;
184  }
185 
186  ps->rng_hash = cmj_hash(branched_state->path_state.rng_hash, i);
187 
188  ccl_global float3 *tp = &kernel_split_state.throughput[ray_index];
189  *tp = throughput;
190 
191  ccl_global Ray *bsdf_ray = &kernel_split_state.ray[ray_index];
192 
193  if (!kernel_branched_path_surface_bounce(
194  kg, sd, sc, j, num_samples, tp, ps, &L->state, bsdf_ray, sum_sample_weight)) {
195  continue;
196  }
197 
198  ps->rng_hash = branched_state->path_state.rng_hash;
199 
200  /* update state for next iteration */
201  branched_state->next_closure = i;
202  branched_state->next_sample = j + 1;
203 
204  /* start the indirect path */
205  *tp *= num_samples_inv;
206 
207  if (kernel_split_branched_indirect_start_shared(kg, ray_index)) {
208  continue;
209  }
210 
211  return true;
212  }
213 
214  branched_state->next_sample = 0;
215  }
216 
217  branched_state->next_closure = sd->num_closure;
218 
219  if (wait_for_shared) {
220  branched_state->waiting_on_shared_samples = (branched_state->shared_sample_count > 0);
221  if (branched_state->waiting_on_shared_samples) {
222  return true;
223  }
224  }
225 
226  return false;
227 }
228 
229 #endif /* __BRANCHED_PATH__ */
230 
unsigned int uint
Definition: BLI_sys_types.h:83
ccl_device_inline void path_radiance_init(KernelGlobals *kg, PathRadiance *L)
ccl_device_inline void path_radiance_copy_indirect(PathRadiance *L, const PathRadiance *L_src)
#define kernel_data
#define ccl_device_inline
#define ccl_global
#define ccl_device_noinline
#define CCL_NAMESPACE_END
ccl_device_inline uint cmj_hash(uint i, uint p)
ccl_device int dequeue_ray_index(int queue_number, ccl_global int *queues, int queue_size, ccl_global int *queue_index)
#define kernel_split_params
#define kernel_split_sd(sd, ray_index)
#define SPLIT_DATA_ENTRIES_BRANCHED_SHARED
#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)
#define IS_STATE(ray_state, ray_index, state)
#define REMOVE_RAY_FLAG(ray_state, ray_index, flag)
ShaderData
#define ADD_RAY_FLAG(ray_state, ray_index, flag)
ShaderClosure
Definition: kernel_types.h:831
@ RAY_BRANCHED_INDIRECT
@ RAY_BRANCHED_INDIRECT_SHARED
@ RAY_REGENERATED
@ RAY_INACTIVE
@ QUEUE_INACTIVE_RAYS
#define L
#define CLOSURE_IS_BSDF_TRANSPARENT(type)
Definition: svm_types.h:608
#define CLOSURE_IS_BSDF(type)
Definition: svm_types.h:595
#define CLOSURE_IS_BSDF_GLOSSY(type)
Definition: svm_types.h:598
#define CLOSURE_IS_BSDF_BSSRDF(type)
Definition: svm_types.h:603
#define CLOSURE_IS_BSDF_DIFFUSE(type)
Definition: svm_types.h:596
@ CLOSURE_BSDF_TRANSPARENT_ID
Definition: svm_types.h:571
#define atomic_fetch_and_inc_uint32(p)
Definition: util_atomic.h:29
ccl_device_inline int ceil_to_int(float f)
Definition: util_math.h:342