Blender  V2.93
filter_reconstruction.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 
20  int y,
21  int storage_stride,
22  int dx,
23  int dy,
24  int t,
25  int buffer_stride,
26  int pass_stride,
27  int frame_offset,
28  bool use_time,
29  const ccl_global float *ccl_restrict buffer,
30  const ccl_global float *ccl_restrict
31  transform,
32  ccl_global int *rank,
33  float weight,
34  ccl_global float *XtWX,
35  ccl_global float3 *XtWY,
36  int localIdx)
37 {
38  if (weight < 1e-3f) {
39  return;
40  }
41 
42  int p_offset = y * buffer_stride + x;
43  int q_offset = (y + dy) * buffer_stride + (x + dx) + frame_offset;
44 
45 #ifdef __KERNEL_GPU__
46  const int stride = storage_stride;
47 #else
48  const int stride = 1;
49  (void)storage_stride;
50 #endif
51 
52 #ifdef __KERNEL_CUDA__
53  ccl_local float shared_design_row[(DENOISE_FEATURES + 1) * CCL_MAX_LOCAL_SIZE];
54  ccl_local_param float *design_row = shared_design_row + localIdx * (DENOISE_FEATURES + 1);
55 #else
56  float design_row[DENOISE_FEATURES + 1];
57 #endif
58 
59  float3 q_color = filter_get_color(buffer + q_offset, pass_stride);
60 
61  /* If the pixel was flagged as an outlier during prefiltering, skip it. */
62  if (ccl_get_feature(buffer + q_offset, 0) < 0.0f) {
63  return;
64  }
65 
67  buffer + p_offset,
68  make_int3(x + dx, y + dy, t),
69  buffer + q_offset,
70  pass_stride,
71  *rank,
72  design_row,
73  transform,
74  stride,
75  use_time);
76 
77 #ifdef __KERNEL_GPU__
78  math_trimatrix_add_gramian_strided(XtWX, (*rank) + 1, design_row, weight, stride);
79  math_vec3_add_strided(XtWY, (*rank) + 1, design_row, weight * q_color, stride);
80 #else
81  math_trimatrix_add_gramian(XtWX, (*rank) + 1, design_row, weight);
82  math_vec3_add(XtWY, (*rank) + 1, design_row, weight * q_color);
83 #endif
84 }
85 
87  int y,
88  ccl_global float *buffer,
89  ccl_global int *rank,
90  int storage_stride,
91  ccl_global float *XtWX,
92  ccl_global float3 *XtWY,
93  int4 buffer_params,
94  int sample)
95 {
96 #ifdef __KERNEL_GPU__
97  const int stride = storage_stride;
98 #else
99  const int stride = 1;
100  (void)storage_stride;
101 #endif
102 
103  if (XtWX[0] < 1e-3f) {
104  /* There is not enough information to determine a denoised result.
105  * As a fallback, keep the original value of the pixel. */
106  return;
107  }
108 
109  /* The weighted average of pixel colors (essentially, the NLM-filtered image).
110  * In case the solution of the linear model fails due to numerical issues or
111  * returns nonsensical negative values, fall back to this value. */
112  float3 mean_color = XtWY[0] / XtWX[0];
113 
114  math_trimatrix_vec3_solve(XtWX, XtWY, (*rank) + 1, stride);
115 
116  float3 final_color = XtWY[0];
117  if (!isfinite3_safe(final_color) ||
118  (final_color.x < -0.01f || final_color.y < -0.01f || final_color.z < -0.01f)) {
119  final_color = mean_color;
120  }
121 
122  /* Clamp pixel value to positive values and reverse the highlight compression transform. */
123  final_color = color_highlight_uncompress(max(final_color, make_float3(0.0f, 0.0f, 0.0f)));
124 
125  ccl_global float *combined_buffer = buffer + (y * buffer_params.y + x + buffer_params.x) *
126  buffer_params.z;
127  if (buffer_params.w >= 0) {
128  final_color *= sample;
129  if (buffer_params.w > 0) {
130  final_color.x += combined_buffer[buffer_params.w + 0];
131  final_color.y += combined_buffer[buffer_params.w + 1];
132  final_color.z += combined_buffer[buffer_params.w + 2];
133  }
134  }
135  combined_buffer[0] = final_color.x;
136  combined_buffer[1] = final_color.y;
137  combined_buffer[2] = final_color.z;
138 }
139 
_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
_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 stride
_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
SIMD_FORCE_INLINE btVector3 transform(const btVector3 &point) const
int x
Definition: btConvexHull.h:149
int w
Definition: btConvexHull.h:149
int y
Definition: btConvexHull.h:149
int z
Definition: btConvexHull.h:149
#define DENOISE_FEATURES
ccl_device_inline float3 filter_get_color(const ccl_global float *ccl_restrict buffer, int pass_stride)
ccl_device_inline void filter_get_design_row_transform(int3 p_pixel, const ccl_global float *ccl_restrict p_buffer, int3 q_pixel, const ccl_global float *ccl_restrict q_buffer, int pass_stride, int rank, float *design_row, const ccl_global float *ccl_restrict transform, int stride, bool use_time)
#define ccl_get_feature(buffer, pass)
CCL_NAMESPACE_BEGIN ccl_device_inline void kernel_filter_construct_gramian(int x, int y, int storage_stride, int dx, int dy, int t, int buffer_stride, int pass_stride, int frame_offset, bool use_time, const ccl_global float *ccl_restrict buffer, const ccl_global float *ccl_restrict transform, ccl_global int *rank, float weight, ccl_global float *XtWX, ccl_global float3 *XtWY, int localIdx)
ccl_device_inline void kernel_filter_finalize(int x, int y, ccl_global float *buffer, ccl_global int *rank, int storage_stride, ccl_global float *XtWX, ccl_global float3 *XtWY, int4 buffer_params, int sample)
#define ccl_restrict
#define ccl_local_param
#define CCL_MAX_LOCAL_SIZE
#define ccl_device_inline
#define ccl_local
#define ccl_global
#define CCL_NAMESPACE_END
#define make_int3(x, y, z)
#define make_float3(x, y, z)
__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
static void sample(SocketReader *reader, int x, int y, float color[4])
float z
Definition: sky_float3.h:35
float y
Definition: sky_float3.h:35
float x
Definition: sky_float3.h:35
float max
ccl_device float3 color_highlight_uncompress(float3 color)
Definition: util_color.h:287
ccl_device_inline bool isfinite3_safe(float3 v)
ccl_device_inline void math_vec3_add(float3 *v, int n, float *x, float3 w)
ccl_device_inline void math_vec3_add_strided(ccl_global float3 *v, int n, float *x, float3 w, int stride)
ccl_device_inline void math_trimatrix_add_gramian_strided(ccl_global float *A, int n, const float *ccl_restrict v, float weight, int stride)
ccl_device_inline void math_trimatrix_vec3_solve(ccl_global float *A, ccl_global float3 *y, int n, int stride)
ccl_device_inline void math_trimatrix_add_gramian(ccl_global float *A, int n, const float *ccl_restrict v, float weight)