Blender  V2.93
kernel_adaptive_sampling.h
Go to the documentation of this file.
1 /*
2  * Copyright 2019 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 #ifndef __KERNEL_ADAPTIVE_SAMPLING_H__
18 #define __KERNEL_ADAPTIVE_SAMPLING_H__
19 
21 
22 /* Determines whether to continue sampling a given pixel or if it has sufficiently converged. */
23 
25  ccl_global float *buffer,
26  int sample)
27 {
28  /* TODO Stefan: Is this better in linear, sRGB or something else? */
29  float4 I = *((ccl_global float4 *)buffer);
30  float4 A = *(ccl_global float4 *)(buffer + kernel_data.film.pass_adaptive_aux_buffer);
31  /* The per pixel error as seen in section 2.1 of
32  * "A hierarchical automatic stopping condition for Monte Carlo global illumination"
33  * A small epsilon is added to the divisor to prevent division by zero. */
34  float error = (fabsf(I.x - A.x) + fabsf(I.y - A.y) + fabsf(I.z - A.z)) /
35  (sample * 0.0001f + sqrtf(I.x + I.y + I.z));
36  if (error < kernel_data.integrator.adaptive_threshold * (float)sample) {
37  /* Set the fourth component to non-zero value to indicate that this pixel has converged. */
38  buffer[kernel_data.film.pass_adaptive_aux_buffer + 3] += 1.0f;
39  }
40 }
41 
42 /* Adjust the values of an adaptively sampled pixel. */
43 
45  ccl_global float *buffer,
46  float sample_multiplier)
47 {
48  *(ccl_global float4 *)(buffer) *= sample_multiplier;
49 
50  /* Scale the aux pass too, this is necessary for progressive rendering to work properly. */
51  kernel_assert(kernel_data.film.pass_adaptive_aux_buffer);
52  *(ccl_global float4 *)(buffer + kernel_data.film.pass_adaptive_aux_buffer) *= sample_multiplier;
53 
54 #ifdef __PASSES__
55  int flag = kernel_data.film.pass_flag;
56 
57  if (flag & PASSMASK(NORMAL))
58  *(ccl_global float3 *)(buffer + kernel_data.film.pass_normal) *= sample_multiplier;
59 
60  if (flag & PASSMASK(UV))
61  *(ccl_global float3 *)(buffer + kernel_data.film.pass_uv) *= sample_multiplier;
62 
63  if (flag & PASSMASK(MOTION)) {
64  *(ccl_global float4 *)(buffer + kernel_data.film.pass_motion) *= sample_multiplier;
65  *(ccl_global float *)(buffer + kernel_data.film.pass_motion_weight) *= sample_multiplier;
66  }
67 
68  if (kernel_data.film.use_light_pass) {
69  int light_flag = kernel_data.film.light_pass_flag;
70 
71  if (light_flag & PASSMASK(MIST))
72  *(ccl_global float *)(buffer + kernel_data.film.pass_mist) *= sample_multiplier;
73 
74  /* Shadow pass omitted on purpose. It has its own scale parameter. */
75 
76  if (light_flag & PASSMASK(DIFFUSE_INDIRECT))
77  *(ccl_global float3 *)(buffer + kernel_data.film.pass_diffuse_indirect) *= sample_multiplier;
78  if (light_flag & PASSMASK(GLOSSY_INDIRECT))
79  *(ccl_global float3 *)(buffer + kernel_data.film.pass_glossy_indirect) *= sample_multiplier;
80  if (light_flag & PASSMASK(TRANSMISSION_INDIRECT))
81  *(ccl_global float3 *)(buffer +
82  kernel_data.film.pass_transmission_indirect) *= sample_multiplier;
83  if (light_flag & PASSMASK(VOLUME_INDIRECT))
84  *(ccl_global float3 *)(buffer + kernel_data.film.pass_volume_indirect) *= sample_multiplier;
85  if (light_flag & PASSMASK(DIFFUSE_DIRECT))
86  *(ccl_global float3 *)(buffer + kernel_data.film.pass_diffuse_direct) *= sample_multiplier;
87  if (light_flag & PASSMASK(GLOSSY_DIRECT))
88  *(ccl_global float3 *)(buffer + kernel_data.film.pass_glossy_direct) *= sample_multiplier;
89  if (light_flag & PASSMASK(TRANSMISSION_DIRECT))
90  *(ccl_global float3 *)(buffer +
91  kernel_data.film.pass_transmission_direct) *= sample_multiplier;
92  if (light_flag & PASSMASK(VOLUME_DIRECT))
93  *(ccl_global float3 *)(buffer + kernel_data.film.pass_volume_direct) *= sample_multiplier;
94 
95  if (light_flag & PASSMASK(EMISSION))
96  *(ccl_global float3 *)(buffer + kernel_data.film.pass_emission) *= sample_multiplier;
97  if (light_flag & PASSMASK(BACKGROUND))
98  *(ccl_global float3 *)(buffer + kernel_data.film.pass_background) *= sample_multiplier;
99  if (light_flag & PASSMASK(AO))
100  *(ccl_global float3 *)(buffer + kernel_data.film.pass_ao) *= sample_multiplier;
101 
102  if (light_flag & PASSMASK(DIFFUSE_COLOR))
103  *(ccl_global float3 *)(buffer + kernel_data.film.pass_diffuse_color) *= sample_multiplier;
104  if (light_flag & PASSMASK(GLOSSY_COLOR))
105  *(ccl_global float3 *)(buffer + kernel_data.film.pass_glossy_color) *= sample_multiplier;
106  if (light_flag & PASSMASK(TRANSMISSION_COLOR))
107  *(ccl_global float3 *)(buffer +
108  kernel_data.film.pass_transmission_color) *= sample_multiplier;
109  }
110 #endif
111 
112 #ifdef __DENOISING_FEATURES__
113 
114 # define scale_float3_variance(buffer, offset, scale) \
115  *(buffer + offset) *= scale; \
116  *(buffer + offset + 1) *= scale; \
117  *(buffer + offset + 2) *= scale; \
118  *(buffer + offset + 3) *= scale * scale; \
119  *(buffer + offset + 4) *= scale * scale; \
120  *(buffer + offset + 5) *= scale * scale;
121 
122 # define scale_shadow_variance(buffer, offset, scale) \
123  *(buffer + offset) *= scale; \
124  *(buffer + offset + 1) *= scale; \
125  *(buffer + offset + 2) *= scale * scale;
126 
127  if (kernel_data.film.pass_denoising_data) {
128  scale_shadow_variance(
129  buffer, kernel_data.film.pass_denoising_data + DENOISING_PASS_SHADOW_A, sample_multiplier);
130  scale_shadow_variance(
131  buffer, kernel_data.film.pass_denoising_data + DENOISING_PASS_SHADOW_B, sample_multiplier);
132  if (kernel_data.film.pass_denoising_clean) {
133  scale_float3_variance(
134  buffer, kernel_data.film.pass_denoising_data + DENOISING_PASS_COLOR, sample_multiplier);
135  *(buffer + kernel_data.film.pass_denoising_clean) *= sample_multiplier;
136  *(buffer + kernel_data.film.pass_denoising_clean + 1) *= sample_multiplier;
137  *(buffer + kernel_data.film.pass_denoising_clean + 2) *= sample_multiplier;
138  }
139  else {
140  scale_float3_variance(
141  buffer, kernel_data.film.pass_denoising_data + DENOISING_PASS_COLOR, sample_multiplier);
142  }
143  scale_float3_variance(
144  buffer, kernel_data.film.pass_denoising_data + DENOISING_PASS_NORMAL, sample_multiplier);
145  scale_float3_variance(
146  buffer, kernel_data.film.pass_denoising_data + DENOISING_PASS_ALBEDO, sample_multiplier);
147  *(buffer + kernel_data.film.pass_denoising_data + DENOISING_PASS_DEPTH) *= sample_multiplier;
148  *(buffer + kernel_data.film.pass_denoising_data + DENOISING_PASS_DEPTH +
149  1) *= sample_multiplier * sample_multiplier;
150  }
151 #endif /* __DENOISING_FEATURES__ */
152 
153  /* Cryptomatte. */
154  if (kernel_data.film.cryptomatte_passes) {
155  int num_slots = 0;
156  num_slots += (kernel_data.film.cryptomatte_passes & CRYPT_OBJECT) ? 1 : 0;
157  num_slots += (kernel_data.film.cryptomatte_passes & CRYPT_MATERIAL) ? 1 : 0;
158  num_slots += (kernel_data.film.cryptomatte_passes & CRYPT_ASSET) ? 1 : 0;
159  num_slots = num_slots * 2 * kernel_data.film.cryptomatte_depth;
160  ccl_global float2 *id_buffer = (ccl_global float2 *)(buffer +
161  kernel_data.film.pass_cryptomatte);
162  for (int slot = 0; slot < num_slots; slot++) {
163  id_buffer[slot].y *= sample_multiplier;
164  }
165  }
166 
167  /* AOVs. */
168  for (int i = 0; i < kernel_data.film.pass_aov_value_num; i++) {
169  *(buffer + kernel_data.film.pass_aov_value + i) *= sample_multiplier;
170  }
171  for (int i = 0; i < kernel_data.film.pass_aov_color_num; i++) {
172  *((ccl_global float4 *)(buffer + kernel_data.film.pass_aov_color) + i) *= sample_multiplier;
173  }
174 }
175 
176 /* This is a simple box filter in two passes.
177  * When a pixel demands more adaptive samples, let its neighboring pixels draw more samples too. */
178 
180 {
181  bool any = false;
182  bool prev = false;
183  for (int x = tile->x; x < tile->x + tile->w; ++x) {
184  int index = tile->offset + x + y * tile->stride;
185  ccl_global float *buffer = tile->buffer + index * kernel_data.film.pass_stride;
186  ccl_global float4 *aux = (ccl_global float4 *)(buffer +
187  kernel_data.film.pass_adaptive_aux_buffer);
188  if ((*aux).w == 0.0f) {
189  any = true;
190  if (x > tile->x && !prev) {
191  index = index - 1;
192  buffer = tile->buffer + index * kernel_data.film.pass_stride;
193  aux = (ccl_global float4 *)(buffer + kernel_data.film.pass_adaptive_aux_buffer);
194  (*aux).w = 0.0f;
195  }
196  prev = true;
197  }
198  else {
199  if (prev) {
200  (*aux).w = 0.0f;
201  }
202  prev = false;
203  }
204  }
205  return any;
206 }
207 
209 {
210  bool prev = false;
211  bool any = false;
212  for (int y = tile->y; y < tile->y + tile->h; ++y) {
213  int index = tile->offset + x + y * tile->stride;
214  ccl_global float *buffer = tile->buffer + index * kernel_data.film.pass_stride;
215  ccl_global float4 *aux = (ccl_global float4 *)(buffer +
216  kernel_data.film.pass_adaptive_aux_buffer);
217  if ((*aux).w == 0.0f) {
218  any = true;
219  if (y > tile->y && !prev) {
220  index = index - tile->stride;
221  buffer = tile->buffer + index * kernel_data.film.pass_stride;
222  aux = (ccl_global float4 *)(buffer + kernel_data.film.pass_adaptive_aux_buffer);
223  (*aux).w = 0.0f;
224  }
225  prev = true;
226  }
227  else {
228  if (prev) {
229  (*aux).w = 0.0f;
230  }
231  prev = false;
232  }
233  }
234  return any;
235 }
236 
238 
239 #endif /* __KERNEL_ADAPTIVE_SAMPLING_H__ */
_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 y
Group RGB to NORMAL
#define A
CCL_NAMESPACE_BEGIN ccl_device void kernel_do_adaptive_stopping(KernelGlobals *kg, ccl_global float *buffer, int sample)
ccl_device bool kernel_do_adaptive_filter_x(KernelGlobals *kg, int y, ccl_global WorkTile *tile)
ccl_device void kernel_adaptive_post_adjust(KernelGlobals *kg, ccl_global float *buffer, float sample_multiplier)
ccl_device bool kernel_do_adaptive_filter_y(KernelGlobals *kg, int x, ccl_global WorkTile *tile)
#define kernel_data
#define kernel_assert(cond)
#define ccl_device
#define ccl_global
#define CCL_NAMESPACE_END
#define fabsf(x)
#define sqrtf(x)
__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
@ CRYPT_ASSET
Definition: kernel_types.h:404
@ CRYPT_OBJECT
Definition: kernel_types.h:402
@ CRYPT_MATERIAL
Definition: kernel_types.h:403
#define PASSMASK(pass)
Definition: kernel_types.h:341
@ DENOISING_PASS_ALBEDO
Definition: kernel_types.h:411
@ DENOISING_PASS_DEPTH
Definition: kernel_types.h:413
@ DENOISING_PASS_SHADOW_B
Definition: kernel_types.h:416
@ DENOISING_PASS_COLOR
Definition: kernel_types.h:417
@ DENOISING_PASS_NORMAL
Definition: kernel_types.h:409
@ DENOISING_PASS_SHADOW_A
Definition: kernel_types.h:415
static void error(const char *str)
Definition: meshlaplacian.c:65
static void sample(SocketReader *reader, int x, int y, float color[4])
#define I
__forceinline bool any(const avxb &b)
Definition: util_avxb.h:218