Blender  V2.93
kernel_emission.h
Go to the documentation of this file.
1 /*
2  * Copyright 2011-2013 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 /* Direction Emission */
21  ShaderData *emission_sd,
22  LightSample *ls,
24  float3 I,
25  differential3 dI,
26  float t,
27  float time)
28 {
29  /* setup shading at emitter */
30  float3 eval = zero_float3();
31 
32  if (shader_constant_emission_eval(kg, ls->shader, &eval)) {
33  if ((ls->prim != PRIM_NONE) && dot(ls->Ng, I) < 0.0f) {
34  ls->Ng = -ls->Ng;
35  }
36  }
37  else {
38  /* Setup shader data and call shader_eval_surface once, better
39  * for GPU coherence and compile times. */
40 #ifdef __BACKGROUND_MIS__
41  if (ls->type == LIGHT_BACKGROUND) {
42  Ray ray;
43  ray.D = ls->D;
44  ray.P = ls->P;
45  ray.t = 1.0f;
46  ray.time = time;
47  ray.dP = differential3_zero();
48  ray.dD = dI;
49 
50  shader_setup_from_background(kg, emission_sd, &ray);
51  }
52  else
53 #endif
54  {
56  emission_sd,
57  ls->P,
58  ls->Ng,
59  I,
60  ls->shader,
61  ls->object,
62  ls->prim,
63  ls->u,
64  ls->v,
65  t,
66  time,
67  false,
68  ls->lamp);
69 
70  ls->Ng = emission_sd->Ng;
71  }
72 
73  /* No proper path flag, we're evaluating this for all closures. that's
74  * weak but we'd have to do multiple evaluations otherwise. */
78 
79  /* Evaluate closures. */
80 #ifdef __BACKGROUND_MIS__
81  if (ls->type == LIGHT_BACKGROUND) {
82  eval = shader_background_eval(emission_sd);
83  }
84  else
85 #endif
86  {
87  eval = shader_emissive_eval(emission_sd);
88  }
89  }
90 
91  eval *= ls->eval_fac;
92 
93  if (ls->lamp != LAMP_NONE) {
94  const ccl_global KernelLight *klight = &kernel_tex_fetch(__lights, ls->lamp);
95  eval *= make_float3(klight->strength[0], klight->strength[1], klight->strength[2]);
96  }
97 
98  return eval;
99 }
100 
102  ShaderData *sd,
103  ShaderData *emission_sd,
104  LightSample *ls,
106  Ray *ray,
107  BsdfEval *eval,
108  bool *is_lamp,
109  float rand_terminate)
110 {
111  if (ls->pdf == 0.0f)
112  return false;
113 
114  /* todo: implement */
116 
117  /* evaluate closure */
118 
119  float3 light_eval = direct_emissive_eval(
120  kg, emission_sd, ls, state, -ls->D, dD, ls->t, sd->time);
121 
122  if (is_zero(light_eval))
123  return false;
124 
125  /* evaluate BSDF at shading point */
126 
127 #ifdef __VOLUME__
128  if (sd->prim != PRIM_NONE)
129  shader_bsdf_eval(kg, sd, ls->D, eval, ls->pdf, ls->shader & SHADER_USE_MIS);
130  else {
131  float bsdf_pdf;
132  shader_volume_phase_eval(kg, sd, ls->D, eval, &bsdf_pdf);
133  if (ls->shader & SHADER_USE_MIS) {
134  /* Multiple importance sampling. */
135  float mis_weight = power_heuristic(ls->pdf, bsdf_pdf);
136  light_eval *= mis_weight;
137  }
138  }
139 #else
140  shader_bsdf_eval(kg, sd, ls->D, eval, ls->pdf, ls->shader & SHADER_USE_MIS);
141 #endif
142 
143  bsdf_eval_mul3(eval, light_eval / ls->pdf);
144 
145 #ifdef __PASSES__
146  /* use visibility flag to skip lights */
147  if (ls->shader & SHADER_EXCLUDE_ANY) {
148  if (ls->shader & SHADER_EXCLUDE_DIFFUSE)
149  eval->diffuse = zero_float3();
150  if (ls->shader & SHADER_EXCLUDE_GLOSSY)
151  eval->glossy = zero_float3();
153  eval->transmission = zero_float3();
154  if (ls->shader & SHADER_EXCLUDE_SCATTER)
155  eval->volume = zero_float3();
156  }
157 #endif
158 
159  if (bsdf_eval_is_zero(eval))
160  return false;
161 
162  if (kernel_data.integrator.light_inv_rr_threshold > 0.0f
163 #ifdef __SHADOW_TRICKS__
164  && (state->flag & PATH_RAY_SHADOW_CATCHER) == 0
165 #endif
166  ) {
167  float probability = max3(fabs(bsdf_eval_sum(eval))) *
168  kernel_data.integrator.light_inv_rr_threshold;
169  if (probability < 1.0f) {
170  if (rand_terminate >= probability) {
171  return false;
172  }
173  bsdf_eval_mul(eval, 1.0f / probability);
174  }
175  }
176 
177  if (ls->shader & SHADER_CAST_SHADOW) {
178  /* setup ray */
179  bool transmit = (dot(sd->Ng, ls->D) < 0.0f);
180  ray->P = ray_offset(sd->P, (transmit) ? -sd->Ng : sd->Ng);
181 
182  if (ls->t == FLT_MAX) {
183  /* distant light */
184  ray->D = ls->D;
185  ray->t = ls->t;
186  }
187  else {
188  /* other lights, avoid self-intersection */
189  ray->D = ray_offset(ls->P, ls->Ng) - ray->P;
190  ray->D = normalize_len(ray->D, &ray->t);
191  }
192 
193  ray->dP = sd->dP;
194  ray->dD = differential3_zero();
195  }
196  else {
197  /* signal to not cast shadow ray */
198  ray->t = 0.0f;
199  }
200 
201  /* return if it's a lamp for shadow pass */
202  *is_lamp = (ls->prim == PRIM_NONE && ls->type != LIGHT_BACKGROUND);
203 
204  return true;
205 }
206 
207 /* Indirect Primitive Emission */
208 
210  KernelGlobals *kg, ShaderData *sd, float t, int path_flag, float bsdf_pdf)
211 {
212  /* evaluate emissive closure */
214 
215 #ifdef __HAIR__
216  if (!(path_flag & PATH_RAY_MIS_SKIP) && (sd->flag & SD_USE_MIS) &&
217  (sd->type & PRIMITIVE_ALL_TRIANGLE))
218 #else
219  if (!(path_flag & PATH_RAY_MIS_SKIP) && (sd->flag & SD_USE_MIS))
220 #endif
221  {
222  /* multiple importance sampling, get triangle light pdf,
223  * and compute weight with respect to BSDF pdf */
224  float pdf = triangle_light_pdf(kg, sd, t);
225  float mis_weight = power_heuristic(bsdf_pdf, pdf);
226 
227  return L * mis_weight;
228  }
229 
230  return L;
231 }
232 
233 /* Indirect Lamp Emission */
234 
236  ShaderData *emission_sd,
238  PathRadiance *L,
239  Ray *ray,
240  float3 throughput)
241 {
242  for (int lamp = 0; lamp < kernel_data.integrator.num_all_lights; lamp++) {
244 
245  if (!lamp_light_eval(kg, lamp, ray->P, ray->D, ray->t, &ls))
246  continue;
247 
248 #ifdef __PASSES__
249  /* use visibility flag to skip lights */
250  if (ls.shader & SHADER_EXCLUDE_ANY) {
251  if (((ls.shader & SHADER_EXCLUDE_DIFFUSE) && (state->flag & PATH_RAY_DIFFUSE)) ||
252  ((ls.shader & SHADER_EXCLUDE_GLOSSY) &&
253  ((state->flag & (PATH_RAY_GLOSSY | PATH_RAY_REFLECT)) ==
255  ((ls.shader & SHADER_EXCLUDE_TRANSMIT) && (state->flag & PATH_RAY_TRANSMIT)) ||
256  ((ls.shader & SHADER_EXCLUDE_SCATTER) && (state->flag & PATH_RAY_VOLUME_SCATTER)))
257  continue;
258  }
259 #endif
260 
261  float3 lamp_L = direct_emissive_eval(
262  kg, emission_sd, &ls, state, -ray->D, ray->dD, ls.t, ray->time);
263 
264 #ifdef __VOLUME__
265  if (state->volume_stack[0].shader != SHADER_NONE) {
266  /* shadow attenuation */
267  Ray volume_ray = *ray;
268  volume_ray.t = ls.t;
269  float3 volume_tp = one_float3();
270  kernel_volume_shadow(kg, emission_sd, state, &volume_ray, &volume_tp);
271  lamp_L *= volume_tp;
272  }
273 #endif
274 
275  if (!(state->flag & PATH_RAY_MIS_SKIP)) {
276  /* multiple importance sampling, get regular light pdf,
277  * and compute weight with respect to BSDF pdf */
278  float mis_weight = power_heuristic(state->ray_pdf, ls.pdf);
279  lamp_L *= mis_weight;
280  }
281 
282  path_radiance_accum_emission(kg, L, state, throughput, lamp_L);
283  }
284 }
285 
286 /* Indirect Background */
287 
289  ShaderData *emission_sd,
291  ccl_global float *buffer,
292  ccl_addr_space Ray *ray)
293 {
294 #ifdef __BACKGROUND__
295  int shader = kernel_data.background.surface_shader;
296 
297  /* Use visibility flag to skip lights. */
298  if (shader & SHADER_EXCLUDE_ANY) {
299  if (((shader & SHADER_EXCLUDE_DIFFUSE) && (state->flag & PATH_RAY_DIFFUSE)) ||
301  ((state->flag & (PATH_RAY_GLOSSY | PATH_RAY_REFLECT)) ==
304  ((shader & SHADER_EXCLUDE_CAMERA) && (state->flag & PATH_RAY_CAMERA)) ||
306  return zero_float3();
307  }
308 
309  /* Evaluate background shader. */
310  float3 L = zero_float3();
312 # ifdef __SPLIT_KERNEL__
313  Ray priv_ray = *ray;
314  shader_setup_from_background(kg, emission_sd, &priv_ray);
315 # else
316  shader_setup_from_background(kg, emission_sd, ray);
317 # endif
318 
320  shader_eval_surface(kg, emission_sd, state, buffer, state->flag | PATH_RAY_EMISSION);
322 
323  L = shader_background_eval(emission_sd);
324  }
325 
326  /* Background MIS weights. */
327 # ifdef __BACKGROUND_MIS__
328  /* Check if background light exists or if we should skip pdf. */
329  if (!(state->flag & PATH_RAY_MIS_SKIP) && kernel_data.background.use_mis) {
330  /* multiple importance sampling, get background light pdf for ray
331  * direction, and compute weight with respect to BSDF pdf */
332  float pdf = background_light_pdf(kg, ray->P, ray->D);
333  float mis_weight = power_heuristic(state->ray_pdf, pdf);
334 
335  return L * mis_weight;
336  }
337 # endif
338 
339  return L;
340 #else
341  return make_float3(0.8f, 0.8f, 0.8f);
342 #endif
343 }
344 
_GL_VOID GLfloat value _GL_VOID_RET _GL_VOID const GLuint GLboolean *residences _GL_BOOL_RET _GL_VOID GLsizei GLfloat GLfloat GLfloat GLfloat const GLubyte *bitmap _GL_VOID_RET _GL_VOID GLenum const void *lists _GL_VOID_RET _GL_VOID const GLdouble *equation _GL_VOID_RET _GL_VOID GLdouble GLdouble blue _GL_VOID_RET _GL_VOID GLfloat GLfloat blue _GL_VOID_RET _GL_VOID GLint GLint blue _GL_VOID_RET _GL_VOID GLshort GLshort blue _GL_VOID_RET _GL_VOID GLubyte GLubyte blue _GL_VOID_RET _GL_VOID GLuint GLuint blue _GL_VOID_RET _GL_VOID GLushort GLushort blue _GL_VOID_RET _GL_VOID GLbyte GLbyte GLbyte alpha _GL_VOID_RET _GL_VOID GLdouble GLdouble GLdouble alpha _GL_VOID_RET _GL_VOID GLfloat GLfloat GLfloat alpha _GL_VOID_RET _GL_VOID GLint GLint GLint alpha _GL_VOID_RET _GL_VOID GLshort GLshort GLshort alpha _GL_VOID_RET _GL_VOID GLubyte GLubyte GLubyte alpha _GL_VOID_RET _GL_VOID GLuint GLuint GLuint alpha _GL_VOID_RET _GL_VOID GLushort GLushort GLushort alpha _GL_VOID_RET _GL_VOID GLenum mode _GL_VOID_RET _GL_VOID GLint GLsizei GLsizei GLenum type _GL_VOID_RET _GL_VOID GLsizei GLenum GLenum const void *pixels _GL_VOID_RET _GL_VOID const void *pointer _GL_VOID_RET _GL_VOID GLdouble v _GL_VOID_RET _GL_VOID GLfloat v _GL_VOID_RET _GL_VOID GLint GLint i2 _GL_VOID_RET _GL_VOID GLint j _GL_VOID_RET _GL_VOID GLfloat param _GL_VOID_RET _GL_VOID GLint param _GL_VOID_RET _GL_VOID GLdouble GLdouble GLdouble GLdouble GLdouble zFar _GL_VOID_RET _GL_UINT GLdouble *equation _GL_VOID_RET _GL_VOID GLenum GLint *params _GL_VOID_RET _GL_VOID GLenum GLfloat *v _GL_VOID_RET _GL_VOID GLenum GLfloat *params _GL_VOID_RET _GL_VOID GLfloat *values _GL_VOID_RET _GL_VOID GLushort *values _GL_VOID_RET _GL_VOID GLenum GLfloat *params _GL_VOID_RET _GL_VOID GLenum GLdouble *params _GL_VOID_RET _GL_VOID GLenum GLint *params _GL_VOID_RET _GL_VOID GLsizei const void *pointer _GL_VOID_RET _GL_VOID GLsizei const void *pointer _GL_VOID_RET _GL_BOOL GLfloat param _GL_VOID_RET _GL_VOID GLint param _GL_VOID_RET _GL_VOID GLenum GLfloat param _GL_VOID_RET _GL_VOID GLenum GLint param _GL_VOID_RET _GL_VOID GLushort pattern _GL_VOID_RET _GL_VOID GLdouble GLdouble GLint GLint const GLdouble *points _GL_VOID_RET _GL_VOID GLdouble GLdouble GLint GLint GLdouble GLdouble GLint GLint const GLdouble *points _GL_VOID_RET _GL_VOID GLdouble GLdouble u2 _GL_VOID_RET _GL_VOID GLdouble GLdouble GLint GLdouble GLdouble v2 _GL_VOID_RET _GL_VOID GLenum GLfloat param _GL_VOID_RET _GL_VOID GLenum GLint param _GL_VOID_RET _GL_VOID GLenum mode _GL_VOID_RET _GL_VOID GLdouble GLdouble nz _GL_VOID_RET _GL_VOID GLfloat GLfloat nz _GL_VOID_RET _GL_VOID GLint GLint nz _GL_VOID_RET _GL_VOID GLshort GLshort nz _GL_VOID_RET _GL_VOID GLsizei const void *pointer _GL_VOID_RET _GL_VOID GLsizei const GLfloat *values _GL_VOID_RET _GL_VOID GLsizei const GLushort *values _GL_VOID_RET _GL_VOID GLint param _GL_VOID_RET _GL_VOID const GLuint const GLclampf *priorities _GL_VOID_RET _GL_VOID GLdouble y _GL_VOID_RET _GL_VOID GLfloat y _GL_VOID_RET _GL_VOID GLint y _GL_VOID_RET _GL_VOID GLshort y _GL_VOID_RET _GL_VOID GLdouble GLdouble z _GL_VOID_RET _GL_VOID GLfloat GLfloat z _GL_VOID_RET _GL_VOID GLint GLint z _GL_VOID_RET _GL_VOID GLshort GLshort z _GL_VOID_RET _GL_VOID GLdouble GLdouble GLdouble w _GL_VOID_RET _GL_VOID GLfloat GLfloat GLfloat w _GL_VOID_RET _GL_VOID GLint GLint GLint w _GL_VOID_RET _GL_VOID GLshort GLshort GLshort w _GL_VOID_RET _GL_VOID GLdouble GLdouble GLdouble y2 _GL_VOID_RET _GL_VOID GLfloat GLfloat GLfloat y2 _GL_VOID_RET _GL_VOID GLint GLint GLint y2 _GL_VOID_RET _GL_VOID GLshort GLshort GLshort y2 _GL_VOID_RET _GL_VOID GLdouble GLdouble GLdouble z _GL_VOID_RET _GL_VOID GLdouble GLdouble z _GL_VOID_RET _GL_VOID GLuint *buffer _GL_VOID_RET _GL_VOID GLdouble t _GL_VOID_RET _GL_VOID GLfloat t _GL_VOID_RET _GL_VOID GLint t _GL_VOID_RET _GL_VOID GLshort t _GL_VOID_RET _GL_VOID GLdouble t
double time
Light lamp
ccl_device_inline float3 ray_offset(float3 P, float3 Ng)
ccl_device_inline bool bsdf_eval_is_zero(BsdfEval *eval)
ccl_device_inline float3 bsdf_eval_sum(const BsdfEval *eval)
ccl_device_inline void bsdf_eval_mul(BsdfEval *eval, float value)
ccl_device_inline void path_radiance_accum_emission(KernelGlobals *kg, PathRadiance *L, ccl_addr_space PathState *state, float3 throughput, float3 value)
ccl_device_inline void bsdf_eval_mul3(BsdfEval *eval, float3 value)
#define kernel_data
#define ccl_addr_space
#define kernel_tex_fetch(tex, index)
#define ccl_optional_struct_init
#define ccl_device_noinline_cpu
#define ccl_global
#define CCL_NAMESPACE_END
#define make_float3(x, y, z)
void KERNEL_FUNCTION_FULL_NAME() shader(KernelGlobals *kg, uint4 *input, float4 *output, int type, int filter, int i, int offset, int sample)
ccl_device differential3 differential3_zero()
ccl_device_noinline_cpu float3 indirect_primitive_emission(KernelGlobals *kg, ShaderData *sd, float t, int path_flag, float bsdf_pdf)
ccl_device_noinline_cpu void indirect_lamp_emission(KernelGlobals *kg, ShaderData *emission_sd, ccl_addr_space PathState *state, PathRadiance *L, Ray *ray, float3 throughput)
ccl_device_noinline_cpu float3 indirect_background(KernelGlobals *kg, ShaderData *emission_sd, ccl_addr_space PathState *state, ccl_global float *buffer, ccl_addr_space Ray *ray)
CCL_NAMESPACE_BEGIN ccl_device_noinline_cpu float3 direct_emissive_eval(KernelGlobals *kg, ShaderData *emission_sd, LightSample *ls, ccl_addr_space PathState *state, float3 I, differential3 dI, float t, float time)
ccl_device_noinline_cpu bool direct_emission(KernelGlobals *kg, ShaderData *sd, ShaderData *emission_sd, LightSample *ls, ccl_addr_space PathState *state, Ray *ray, BsdfEval *eval, bool *is_lamp, float rand_terminate)
ccl_device bool lamp_light_eval(KernelGlobals *kg, int lamp, float3 P, float3 D, float t, LightSample *ls)
Definition: kernel_light.h:179
ccl_device_forceinline float triangle_light_pdf(KernelGlobals *kg, ShaderData *sd, float t)
Definition: kernel_light.h:382
ccl_device float power_heuristic(float a, float b)
ccl_device_inline void path_state_modify_bounce(ccl_addr_space PathState *state, bool increase)
ccl_device void shader_eval_surface(KernelGlobals *kg, ShaderData *sd, ccl_addr_space PathState *state, ccl_global float *buffer, int path_flag)
ccl_device_inline void shader_setup_from_sample(KernelGlobals *kg, ShaderData *sd, const float3 P, const float3 Ng, const float3 I, int shader, int object, int prim, float u, float v, float t, float time, bool object_space, int lamp)
ccl_device void shader_bsdf_eval(KernelGlobals *kg, ShaderData *sd, const float3 omega_in, BsdfEval *eval, float light_pdf, bool use_mis)
ccl_device float3 shader_background_eval(ShaderData *sd)
ccl_device_inline void shader_setup_from_background(KernelGlobals *kg, ShaderData *sd, const Ray *ray)
ccl_device bool shader_constant_emission_eval(KernelGlobals *kg, int shader, float3 *eval)
ccl_device float3 shader_emissive_eval(ShaderData *sd)
__kernel void ccl_constant KernelData ccl_global void ccl_global char ccl_global int ccl_global char ccl_global unsigned int ccl_global float * buffer
@ SD_USE_MIS
Definition: kernel_types.h:869
@ PRIMITIVE_ALL_TRIANGLE
Definition: kernel_types.h:697
#define __SHADOW_TRICKS__
Definition: kernel_types.h:96
#define SHADER_NONE
Definition: kernel_types.h:58
#define PRIM_NONE
Definition: kernel_types.h:60
@ PATH_RAY_REFLECT
Definition: kernel_types.h:267
@ PATH_RAY_SHADOW_CATCHER
Definition: kernel_types.h:306
@ PATH_RAY_TRANSMIT
Definition: kernel_types.h:268
@ PATH_RAY_EMISSION
Definition: kernel_types.h:321
@ PATH_RAY_VOLUME_SCATTER
Definition: kernel_types.h:290
@ PATH_RAY_MIS_SKIP
Definition: kernel_types.h:299
@ PATH_RAY_GLOSSY
Definition: kernel_types.h:270
@ PATH_RAY_DIFFUSE
Definition: kernel_types.h:269
@ PATH_RAY_CAMERA
Definition: kernel_types.h:266
ShaderData
@ SHADER_EXCLUDE_CAMERA
Definition: kernel_types.h:588
@ SHADER_EXCLUDE_DIFFUSE
Definition: kernel_types.h:585
@ SHADER_USE_MIS
Definition: kernel_types.h:584
@ SHADER_EXCLUDE_ANY
Definition: kernel_types.h:590
@ SHADER_EXCLUDE_SCATTER
Definition: kernel_types.h:589
@ SHADER_EXCLUDE_GLOSSY
Definition: kernel_types.h:586
@ SHADER_EXCLUDE_TRANSMIT
Definition: kernel_types.h:587
@ SHADER_CAST_SHADOW
Definition: kernel_types.h:582
@ LIGHT_BACKGROUND
Definition: kernel_types.h:602
#define LAMP_NONE
Definition: kernel_types.h:61
static ulong state[N]
#define L
#define I
float3 transmission
Definition: kernel_types.h:569
float3 diffuse
Definition: kernel_types.h:566
float3 glossy
Definition: kernel_types.h:568
float3 volume
Definition: kernel_types.h:571
float eval_fac
Definition: kernel_light.h:30
LightType type
Definition: kernel_light.h:35
float t
Definition: kernel_types.h:649
differential3 dD
Definition: kernel_types.h:660
float3 P
Definition: kernel_types.h:647
float time
Definition: kernel_types.h:650
differential3 dP
Definition: kernel_types.h:659
float3 D
Definition: kernel_types.h:648
ccl_device_inline float dot(const float2 &a, const float2 &b)
ccl_device_inline float2 normalize_len(const float2 &a, float *t)
ccl_device_inline bool is_zero(const float2 &a)
ccl_device_inline float2 fabs(const float2 &a)
ccl_device_inline float3 one_float3()
ccl_device_inline float3 zero_float3()
ccl_device_inline float max3(float3 a)