Blender  V2.93
kernel_subsurface_scatter.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 #if defined(__BRANCHED_PATH__) && defined(__SUBSURFACE__)
20 
21 ccl_device_inline void kernel_split_branched_path_subsurface_indirect_light_init(KernelGlobals *kg,
22  int ray_index)
23 {
24  kernel_split_branched_path_indirect_loop_init(kg, ray_index);
25 
26  SplitBranchedState *branched_state = &kernel_split_state.branched_state[ray_index];
27 
28  branched_state->ss_next_closure = 0;
29  branched_state->ss_next_sample = 0;
30 
31  branched_state->num_hits = 0;
32  branched_state->next_hit = 0;
33 
35 }
36 
37 ccl_device_noinline bool kernel_split_branched_path_subsurface_indirect_light_iter(
38  KernelGlobals *kg, int ray_index)
39 {
40  SplitBranchedState *branched_state = &kernel_split_state.branched_state[ray_index];
41 
42  ShaderData *sd = kernel_split_sd(branched_state_sd, ray_index);
43  PathRadiance *L = &kernel_split_state.path_radiance[ray_index];
44  ShaderData *emission_sd = AS_SHADER_DATA(&kernel_split_state.sd_DL_shadow[ray_index]);
45 
46  for (int i = branched_state->ss_next_closure; i < sd->num_closure; i++) {
47  ShaderClosure *sc = &sd->closure[i];
48 
49  if (!CLOSURE_IS_BSSRDF(sc->type))
50  continue;
51 
52  /* Closure memory will be overwritten, so read required variables now. */
53  Bssrdf *bssrdf = (Bssrdf *)sc;
54  ClosureType bssrdf_type = sc->type;
55  float bssrdf_roughness = bssrdf->roughness;
56 
57  /* set up random number generator */
58  if (branched_state->ss_next_sample == 0 && branched_state->next_hit == 0 &&
59  branched_state->next_closure == 0 && branched_state->next_sample == 0) {
60  branched_state->lcg_state = lcg_state_init_addrspace(&branched_state->path_state,
61  0x68bc21eb);
62  }
63  int num_samples = kernel_data.integrator.subsurface_samples * 3;
64  float num_samples_inv = 1.0f / num_samples;
65  uint bssrdf_rng_hash = cmj_hash(branched_state->path_state.rng_hash, i);
66 
67  /* do subsurface scatter step with copy of shader data, this will
68  * replace the BSSRDF with a diffuse BSDF closure */
69  for (int j = branched_state->ss_next_sample; j < num_samples; j++) {
70  ccl_global PathState *hit_state = &kernel_split_state.path_state[ray_index];
71  *hit_state = branched_state->path_state;
72  hit_state->rng_hash = bssrdf_rng_hash;
73  path_state_branch(hit_state, j, num_samples);
74 
75  ccl_global LocalIntersection *ss_isect = &branched_state->ss_isect;
76  float bssrdf_u, bssrdf_v;
78  kg, bssrdf_rng_hash, hit_state, j, num_samples, PRNG_BSDF_U, &bssrdf_u, &bssrdf_v);
79 
80  /* intersection is expensive so avoid doing multiple times for the same input */
81  if (branched_state->next_hit == 0 && branched_state->next_closure == 0 &&
82  branched_state->next_sample == 0) {
83  uint lcg_state = branched_state->lcg_state;
84  LocalIntersection ss_isect_private;
85 
87  kg, &ss_isect_private, sd, hit_state, sc, &lcg_state, bssrdf_u, bssrdf_v, true);
88 
89  branched_state->lcg_state = lcg_state;
90  *ss_isect = ss_isect_private;
91  }
92 
93  hit_state->rng_offset += PRNG_BOUNCE_NUM;
94 
95 # ifdef __VOLUME__
96  Ray volume_ray = branched_state->ray;
97  bool need_update_volume_stack = kernel_data.integrator.use_volumes &&
98  sd->object_flag & SD_OBJECT_INTERSECTS_VOLUME;
99 # endif /* __VOLUME__ */
100 
101  /* compute lighting with the BSDF closure */
102  for (int hit = branched_state->next_hit; hit < branched_state->num_hits; hit++) {
103  ShaderData *bssrdf_sd = kernel_split_sd(sd, ray_index);
104  *bssrdf_sd = *sd; /* note: copy happens each iteration of inner loop, this is
105  * important as the indirect path will write into bssrdf_sd */
106 
107  LocalIntersection ss_isect_private = *ss_isect;
109  kg, &ss_isect_private, hit, bssrdf_sd, hit_state, bssrdf_type, bssrdf_roughness);
110  *ss_isect = ss_isect_private;
111 
112 # ifdef __VOLUME__
113  if (need_update_volume_stack) {
114  /* Setup ray from previous surface point to the new one. */
115  float3 P = ray_offset(bssrdf_sd->P, -bssrdf_sd->Ng);
116  volume_ray.D = normalize_len(P - volume_ray.P, &volume_ray.t);
117 
118  for (int k = 0; k < VOLUME_STACK_SIZE; k++) {
119  hit_state->volume_stack[k] = branched_state->path_state.volume_stack[k];
120  }
121 
122  kernel_volume_stack_update_for_subsurface(
123  kg, emission_sd, &volume_ray, hit_state->volume_stack);
124  }
125 # endif /* __VOLUME__ */
126 
127 # ifdef __EMISSION__
128  if (branched_state->next_closure == 0 && branched_state->next_sample == 0) {
129  /* direct light */
130  if (kernel_data.integrator.use_direct_light) {
131  int all = (kernel_data.integrator.sample_all_lights_direct) ||
132  (hit_state->flag & PATH_RAY_SHADOW_CATCHER);
133  kernel_branched_path_surface_connect_light(kg,
134  bssrdf_sd,
135  emission_sd,
136  hit_state,
137  branched_state->throughput,
138  num_samples_inv,
139  L,
140  all);
141  }
142  }
143 # endif /* __EMISSION__ */
144 
145  /* indirect light */
146  if (kernel_split_branched_path_surface_indirect_light_iter(
147  kg, ray_index, num_samples_inv, bssrdf_sd, false, false)) {
148  branched_state->ss_next_closure = i;
149  branched_state->ss_next_sample = j;
150  branched_state->next_hit = hit;
151 
152  return true;
153  }
154 
155  branched_state->next_closure = 0;
156  }
157 
158  branched_state->next_hit = 0;
159  }
160 
161  branched_state->ss_next_sample = 0;
162  }
163 
164  branched_state->ss_next_closure = sd->num_closure;
165 
166  branched_state->waiting_on_shared_samples = (branched_state->shared_sample_count > 0);
167  if (branched_state->waiting_on_shared_samples) {
168  return true;
169  }
170 
171  kernel_split_branched_path_indirect_loop_end(kg, ray_index);
172 
173  return false;
174 }
175 
176 #endif /* __BRANCHED_PATH__ && __SUBSURFACE__ */
177 
179 {
180  int thread_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0);
181  if (thread_index == 0) {
182  /* We will empty both queues in this kernel. */
185  }
186 
187  int ray_index = ccl_global_id(1) * ccl_global_size(0) + ccl_global_id(0);
188  ray_index = get_ray_index(kg,
189  ray_index,
191  kernel_split_state.queue_data,
192  kernel_split_params.queue_size,
193  1);
195  thread_index,
197  kernel_split_state.queue_data,
198  kernel_split_params.queue_size,
199  1);
200 
201 #ifdef __SUBSURFACE__
202  ccl_global char *ray_state = kernel_split_state.ray_state;
203 
204  if (IS_STATE(ray_state, ray_index, RAY_ACTIVE)) {
205  ccl_global PathState *state = &kernel_split_state.path_state[ray_index];
206  PathRadiance *L = &kernel_split_state.path_radiance[ray_index];
207  ccl_global Ray *ray = &kernel_split_state.ray[ray_index];
208  ccl_global float3 *throughput = &kernel_split_state.throughput[ray_index];
209  ccl_global SubsurfaceIndirectRays *ss_indirect = &kernel_split_state.ss_rays[ray_index];
210  ShaderData *sd = kernel_split_sd(sd, ray_index);
211  ShaderData *emission_sd = AS_SHADER_DATA(&kernel_split_state.sd_DL_shadow[ray_index]);
212 
213  if (sd->flag & SD_BSSRDF) {
214 
215 # ifdef __BRANCHED_PATH__
216  if (!kernel_data.integrator.branched ||
217  IS_FLAG(ray_state, ray_index, RAY_BRANCHED_INDIRECT)) {
218 # endif
219  if (kernel_path_subsurface_scatter(
220  kg, sd, emission_sd, L, state, ray, throughput, ss_indirect)) {
221  kernel_split_path_end(kg, ray_index);
222  }
223 # ifdef __BRANCHED_PATH__
224  }
225  else {
226  kernel_split_branched_path_subsurface_indirect_light_init(kg, ray_index);
227 
228  if (kernel_split_branched_path_subsurface_indirect_light_iter(kg, ray_index)) {
230  }
231  }
232 # endif
233  }
234  }
235 
236 # ifdef __BRANCHED_PATH__
237  if (ccl_global_id(0) == 0 && ccl_global_id(1) == 0) {
239  }
240 
241  /* iter loop */
242  ray_index = get_ray_index(kg,
245  kernel_split_state.queue_data,
246  kernel_split_params.queue_size,
247  1);
248 
250  /* for render passes, sum and reset indirect light pass variables
251  * for the next samples */
252  path_radiance_sum_indirect(&kernel_split_state.path_radiance[ray_index]);
253  path_radiance_reset_indirect(&kernel_split_state.path_radiance[ray_index]);
254 
255  if (kernel_split_branched_path_subsurface_indirect_light_iter(kg, ray_index)) {
257  }
258  }
259 # endif /* __BRANCHED_PATH__ */
260 
261 #endif /* __SUBSURFACE__ */
262 }
263 
unsigned int uint
Definition: BLI_sys_types.h:83
ccl_device_inline float3 ray_offset(float3 P, float3 Ng)
ccl_device_inline void path_radiance_sum_indirect(PathRadiance *L)
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_device
#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_inline void path_state_branch(ccl_addr_space PathState *state, int branch, int num_branches)
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_inline uint lcg_state_init_addrspace(ccl_addr_space PathState *state, uint scramble)
ccl_device_inline void path_branched_rng_2D(KernelGlobals *kg, uint rng_hash, const ccl_addr_space PathState *state, int branch, int num_branches, int dimension, float *fx, float *fy)
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
ccl_device_inline int subsurface_scatter_multi_intersect(KernelGlobals *kg, LocalIntersection *ss_isect, ShaderData *sd, ccl_addr_space PathState *state, const ShaderClosure *sc, uint *lcg_state, float bssrdf_u, float bssrdf_v, bool all)
ccl_device_noinline void subsurface_scatter_multi_setup(KernelGlobals *kg, LocalIntersection *ss_isect, int hit, ShaderData *sd, ccl_addr_space PathState *state, ClosureType type, float roughness)
CCL_NAMESPACE_BEGIN ccl_device void kernel_subsurface_scatter(KernelGlobals *kg)
#define IS_FLAG(ray_state, ray_index, flag)
@ SD_BSSRDF
Definition: kernel_types.h:851
#define AS_SHADER_DATA(shader_data_tiny_storage)
#define IS_STATE(ray_state, ray_index, state)
@ PRNG_BOUNCE_NUM
Definition: kernel_types.h:248
@ PRNG_BSDF_U
Definition: kernel_types.h:240
@ PATH_RAY_SHADOW_CATCHER
Definition: kernel_types.h:306
ShaderData
@ SD_OBJECT_INTERSECTS_VOLUME
Definition: kernel_types.h:914
#define ADD_RAY_FLAG(ray_state, ray_index, flag)
#define VOLUME_STACK_SIZE
Definition: kernel_types.h:64
#define ASSIGN_RAY_STATE(ray_state, ray_index, state)
ShaderClosure
Definition: kernel_types.h:831
@ RAY_BRANCHED_INDIRECT
@ RAY_ACTIVE
@ RAY_BRANCHED_SUBSURFACE_INDIRECT
@ RAY_SUBSURFACE_INDIRECT_NEXT_ITER
@ RAY_REGENERATED
@ QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS
@ QUEUE_SUBSURFACE_INDIRECT_ITER
@ QUEUE_ACTIVE_AND_REGENERATED_RAYS
static float P(float k)
Definition: math_interp.c:41
static ulong state[N]
#define L
closure color bssrdf(string method, normal N, vector radius, color albedo) BUILTIN
Definition: bssrdf.h:22
float t
Definition: kernel_types.h:649
float3 P
Definition: kernel_types.h:647
float3 D
Definition: kernel_types.h:648
ClosureType
Definition: svm_types.h:527
#define CLOSURE_IS_BSSRDF(type)
Definition: svm_types.h:623
__forceinline bool all(const avxb &b)
Definition: util_avxb.h:214
ccl_device_inline float2 normalize_len(const float2 &a, float *t)