Blender  V2.93
filter_cpu_impl.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 
17 /* Templated common implementation part of all CPU kernels.
18  *
19  * The idea is that particular .cpp files sets needed optimization flags and
20  * simply includes this file without worry of copying actual implementation over.
21  */
22 
24 
26 
27 #ifdef KERNEL_STUB
28 # define STUB_ASSERT(arch, name) \
29  assert(!(#name " kernel stub for architecture " #arch " was called!"))
30 #endif
31 
33 
34 /* Denoise filter */
35 
37  TileInfo *tile_info,
38  int x,
39  int y,
40  float *unfilteredA,
41  float *unfilteredB,
42  float *sampleVariance,
43  float *sampleVarianceV,
44  float *bufferVariance,
45  int *prefilter_rect,
46  int buffer_pass_stride,
47  int buffer_denoising_offset)
48 {
49 #ifdef KERNEL_STUB
50  STUB_ASSERT(KERNEL_ARCH, filter_divide_shadow);
51 #else
53  tile_info,
54  x,
55  y,
56  unfilteredA,
57  unfilteredB,
58  sampleVariance,
59  sampleVarianceV,
60  bufferVariance,
61  load_int4(prefilter_rect),
62  buffer_pass_stride,
63  buffer_denoising_offset);
64 #endif
65 }
66 
68  TileInfo *tile_info,
69  int m_offset,
70  int v_offset,
71  int x,
72  int y,
73  float *mean,
74  float *variance,
75  float scale,
76  int *prefilter_rect,
77  int buffer_pass_stride,
78  int buffer_denoising_offset)
79 {
80 #ifdef KERNEL_STUB
81  STUB_ASSERT(KERNEL_ARCH, filter_get_feature);
82 #else
84  tile_info,
85  m_offset,
86  v_offset,
87  x,
88  y,
89  mean,
90  variance,
91  scale,
92  load_int4(prefilter_rect),
93  buffer_pass_stride,
94  buffer_denoising_offset);
95 #endif
96 }
97 
99  int x,
100  int y,
101  int *buffer_params,
102  float *from,
103  float *buffer,
104  int out_offset,
105  int *prefilter_rect)
106 {
107 #ifdef KERNEL_STUB
108  STUB_ASSERT(KERNEL_ARCH, filter_write_feature);
109 #else
111  sample, x, y, load_int4(buffer_params), from, buffer, out_offset, load_int4(prefilter_rect));
112 #endif
113 }
114 
116  int y,
117  ccl_global float *image,
118  ccl_global float *variance,
119  ccl_global float *depth,
120  ccl_global float *output,
121  int *rect,
122  int pass_stride)
123 {
124 #ifdef KERNEL_STUB
125  STUB_ASSERT(KERNEL_ARCH, filter_detect_outliers);
126 #else
128  x, y, image, variance, depth, output, load_int4(rect), pass_stride);
129 #endif
130 }
131 
133  int x, int y, float *mean, float *variance, float *a, float *b, int *prefilter_rect, int r)
134 {
135 #ifdef KERNEL_STUB
136  STUB_ASSERT(KERNEL_ARCH, filter_combine_halves);
137 #else
138  kernel_filter_combine_halves(x, y, mean, variance, a, b, load_int4(prefilter_rect), r);
139 #endif
140 }
141 
143  TileInfo *tile_info,
144  int x,
145  int y,
146  int storage_ofs,
147  float *transform,
148  int *rank,
149  int *prefilter_rect,
150  int pass_stride,
151  int frame_stride,
152  bool use_time,
153  int radius,
154  float pca_threshold)
155 {
156 #ifdef KERNEL_STUB
158 #else
159  rank += storage_ofs;
160  transform += storage_ofs * TRANSFORM_SIZE;
162  tile_info,
163  x,
164  y,
165  load_int4(prefilter_rect),
166  pass_stride,
167  frame_stride,
168  use_time,
169  transform,
170  rank,
171  radius,
172  pca_threshold);
173 #endif
174 }
175 
177  int dy,
178  float *weight_image,
179  float *variance_image,
180  float *scale_image,
181  float *difference_image,
182  int *rect,
183  int stride,
184  int channel_offset,
185  int frame_offset,
186  float a,
187  float k_2)
188 {
189 #ifdef KERNEL_STUB
191 #else
193  dy,
194  weight_image,
195  variance_image,
196  scale_image,
197  difference_image,
198  load_int4(rect),
199  stride,
200  channel_offset,
201  frame_offset,
202  a,
203  k_2);
204 #endif
205 }
206 
208  float *difference_image, float *out_image, int *rect, int stride, int f)
209 {
210 #ifdef KERNEL_STUB
211  STUB_ASSERT(KERNEL_ARCH, filter_nlm_blur);
212 #else
213  kernel_filter_nlm_blur(difference_image, out_image, load_int4(rect), stride, f);
214 #endif
215 }
216 
218  float *difference_image, float *out_image, int *rect, int stride, int f)
219 {
220 #ifdef KERNEL_STUB
221  STUB_ASSERT(KERNEL_ARCH, filter_nlm_calc_weight);
222 #else
223  kernel_filter_nlm_calc_weight(difference_image, out_image, load_int4(rect), stride, f);
224 #endif
225 }
226 
228  int dy,
229  float *difference_image,
230  float *image,
231  float *temp_image,
232  float *out_image,
233  float *accum_image,
234  int *rect,
235  int channel_offset,
236  int stride,
237  int f)
238 {
239 #ifdef KERNEL_STUB
241 #else
243  dy,
244  difference_image,
245  image,
246  temp_image,
247  out_image,
248  accum_image,
249  load_int4(rect),
250  channel_offset,
251  stride,
252  f);
253 #endif
254 }
255 
257  int dy,
258  int t,
259  float *difference_image,
260  float *buffer,
261  float *transform,
262  int *rank,
263  float *XtWX,
264  float3 *XtWY,
265  int *rect,
266  int *filter_window,
267  int stride,
268  int f,
269  int pass_stride,
270  int frame_offset,
271  bool use_time)
272 {
273 #ifdef KERNEL_STUB
275 #else
277  dy,
278  t,
279  difference_image,
280  buffer,
281  transform,
282  rank,
283  XtWX,
284  XtWY,
285  load_int4(rect),
286  load_int4(filter_window),
287  stride,
288  f,
289  pass_stride,
290  frame_offset,
291  use_time);
292 #endif
293 }
294 
296  float *accum_image,
297  int *rect,
298  int stride)
299 {
300 #ifdef KERNEL_STUB
301  STUB_ASSERT(KERNEL_ARCH, filter_nlm_normalize);
302 #else
303  kernel_filter_nlm_normalize(out_image, accum_image, load_int4(rect), stride);
304 #endif
305 }
306 
308  int y,
309  int storage_ofs,
310  float *buffer,
311  int *rank,
312  float *XtWX,
313  float3 *XtWY,
314  int *buffer_params,
315  int sample)
316 {
317 #ifdef KERNEL_STUB
318  STUB_ASSERT(KERNEL_ARCH, filter_finalize);
319 #else
320  XtWX += storage_ofs * XTWX_SIZE;
321  XtWY += storage_ofs * XTWY_SIZE;
322  rank += storage_ofs;
323  kernel_filter_finalize(x, y, buffer, rank, 1, XtWX, XtWY, load_int4(buffer_params), sample);
324 #endif
325 }
326 
327 #undef KERNEL_STUB
328 #undef STUB_ASSERT
329 #undef KERNEL_ARCH
330 
_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 GLdouble r _GL_VOID_RET _GL_VOID GLfloat GLfloat r _GL_VOID_RET _GL_VOID GLint GLint r _GL_VOID_RET _GL_VOID GLshort GLshort r _GL_VOID_RET _GL_VOID GLdouble GLdouble r
_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
#define output
StackEntry * from
#define KERNEL_FUNCTION_FULL_NAME(name)
Definition: filter.h:30
#define KERNEL_ARCH
Definition: filter.h:47
void KERNEL_FUNCTION_FULL_NAME() filter_nlm_calc_difference(int dx, int dy, float *weight_image, float *variance_image, float *scale_image, float *difference_image, int *rect, int stride, int channel_offset, int frame_offset, float a, float k_2)
void KERNEL_FUNCTION_FULL_NAME() filter_nlm_update_output(int dx, int dy, float *difference_image, float *image, float *temp_image, float *out_image, float *accum_image, int *rect, int channel_offset, int stride, int f)
void KERNEL_FUNCTION_FULL_NAME() filter_write_feature(int sample, int x, int y, int *buffer_params, float *from, float *buffer, int out_offset, int *prefilter_rect)
void KERNEL_FUNCTION_FULL_NAME() filter_get_feature(int sample, TileInfo *tile_info, int m_offset, int v_offset, int x, int y, float *mean, float *variance, float scale, int *prefilter_rect, int buffer_pass_stride, int buffer_denoising_offset)
void KERNEL_FUNCTION_FULL_NAME() filter_construct_transform(float *buffer, TileInfo *tile_info, int x, int y, int storage_ofs, float *transform, int *rank, int *prefilter_rect, int pass_stride, int frame_stride, bool use_time, int radius, float pca_threshold)
CCL_NAMESPACE_BEGIN void KERNEL_FUNCTION_FULL_NAME() filter_divide_shadow(int sample, TileInfo *tile_info, int x, int y, float *unfilteredA, float *unfilteredB, float *sampleVariance, float *sampleVarianceV, float *bufferVariance, int *prefilter_rect, int buffer_pass_stride, int buffer_denoising_offset)
void KERNEL_FUNCTION_FULL_NAME() filter_detect_outliers(int x, int y, ccl_global float *image, ccl_global float *variance, ccl_global float *depth, ccl_global float *output, int *rect, int pass_stride)
void KERNEL_FUNCTION_FULL_NAME() filter_nlm_calc_weight(float *difference_image, float *out_image, int *rect, int stride, int f)
void KERNEL_FUNCTION_FULL_NAME() filter_finalize(int x, int y, int storage_ofs, float *buffer, int *rank, float *XtWX, float3 *XtWY, int *buffer_params, int sample)
void KERNEL_FUNCTION_FULL_NAME() filter_nlm_construct_gramian(int dx, int dy, int t, float *difference_image, float *buffer, float *transform, int *rank, float *XtWX, float3 *XtWY, int *rect, int *filter_window, int stride, int f, int pass_stride, int frame_offset, bool use_time)
void KERNEL_FUNCTION_FULL_NAME() filter_combine_halves(int x, int y, float *mean, float *variance, float *a, float *b, int *prefilter_rect, int r)
void KERNEL_FUNCTION_FULL_NAME() filter_nlm_blur(float *difference_image, float *out_image, int *rect, int stride, int f)
void KERNEL_FUNCTION_FULL_NAME() filter_nlm_normalize(float *out_image, float *accum_image, int *rect, int stride)
#define XTWY_SIZE
#define TRANSFORM_SIZE
#define XTWX_SIZE
ccl_device_inline void kernel_filter_nlm_construct_gramian(int dx, int dy, int t, const float *ccl_restrict difference_image, const float *ccl_restrict buffer, float *transform, int *rank, float *XtWX, float3 *XtWY, int4 rect, int4 filter_window, int stride, int f, int pass_stride, int frame_offset, bool use_time)
ccl_device_inline void kernel_filter_nlm_blur(const float *ccl_restrict difference_image, float *out_image, int4 rect, int stride, int f)
ccl_device_inline void kernel_filter_nlm_calc_weight(const float *ccl_restrict difference_image, float *out_image, int4 rect, int stride, int f)
ccl_device_inline void kernel_filter_nlm_update_output(int dx, int dy, const float *ccl_restrict difference_image, const float *ccl_restrict image, float *temp_image, float *out_image, float *accum_image, int4 rect, int channel_offset, int stride, int f)
ccl_device_inline void kernel_filter_nlm_calc_difference(int dx, int dy, const float *ccl_restrict weight_image, const float *ccl_restrict variance_image, const float *ccl_restrict scale_image, float *difference_image, int4 rect, int stride, int channel_offset, int frame_offset, float a, float k_2)
ccl_device_inline void kernel_filter_nlm_normalize(float *out_image, const float *ccl_restrict accum_image, int4 rect, int w)
ccl_device void kernel_filter_write_feature(int sample, int x, int y, int4 buffer_params, ccl_global float *from, ccl_global float *buffer, int out_offset, int4 rect)
ccl_device void kernel_filter_get_feature(int sample, CCL_FILTER_TILE_INFO, int m_offset, int v_offset, int x, int y, ccl_global float *mean, ccl_global float *variance, float scale, int4 rect, int buffer_pass_stride, int buffer_denoising_offset)
CCL_NAMESPACE_BEGIN ccl_device void kernel_filter_divide_shadow(int sample, CCL_FILTER_TILE_INFO, int x, int y, ccl_global float *unfilteredA, ccl_global float *unfilteredB, ccl_global float *sampleVariance, ccl_global float *sampleVarianceV, ccl_global float *bufferVariance, int4 rect, int buffer_pass_stride, int buffer_denoising_offset)
ccl_device void kernel_filter_detect_outliers(int x, int y, ccl_global float *in, ccl_global float *variance_out, ccl_global float *depth, ccl_global float *image_out, int4 rect, int pass_stride)
ccl_device void kernel_filter_combine_halves(int x, int y, ccl_global float *mean, ccl_global float *variance, ccl_global float *a, ccl_global float *b, int4 rect, int r)
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)
CCL_NAMESPACE_BEGIN ccl_device void kernel_filter_construct_transform(const float *ccl_restrict buffer, CCL_FILTER_TILE_INFO, int x, int y, int4 rect, int pass_stride, int frame_stride, bool use_time, float *transform, int *rank, int radius, float pca_threshold)
#define ccl_global
#define CCL_NAMESPACE_END
__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 unsigned a[3]
Definition: RandGen.cpp:92
static void sample(SocketReader *reader, int x, int y, float color[4])
ccl_device_inline int4 load_int4(const int *v)