Fix T53410: 3D Text always recalculated
[blender.git] / intern / cycles / kernel / kernels / cuda / filter.cu
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 /* CUDA kernel entry points */
18
19 #ifdef __CUDA_ARCH__
20
21 #include "kernel_config.h"
22
23 #include "kernel/kernel_compat_cuda.h"
24
25 #include "kernel/filter/filter_kernel.h"
26
27 /* kernels */
28
29 extern "C" __global__ void
30 CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS)
31 kernel_cuda_filter_divide_shadow(int sample,
32                                  TilesInfo *tiles,
33                                  float *unfilteredA,
34                                  float *unfilteredB,
35                                  float *sampleVariance,
36                                  float *sampleVarianceV,
37                                  float *bufferVariance,
38                                  int4 prefilter_rect,
39                                  int buffer_pass_stride,
40                                  int buffer_denoising_offset)
41 {
42         int x = prefilter_rect.x + blockDim.x*blockIdx.x + threadIdx.x;
43         int y = prefilter_rect.y + blockDim.y*blockIdx.y + threadIdx.y;
44         if(x < prefilter_rect.z && y < prefilter_rect.w) {
45                 kernel_filter_divide_shadow(sample,
46                                             tiles,
47                                             x, y,
48                                             unfilteredA,
49                                             unfilteredB,
50                                             sampleVariance,
51                                             sampleVarianceV,
52                                             bufferVariance,
53                                             prefilter_rect,
54                                             buffer_pass_stride,
55                                             buffer_denoising_offset);
56         }
57 }
58
59 extern "C" __global__ void
60 CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS)
61 kernel_cuda_filter_get_feature(int sample,
62                                TilesInfo *tiles,
63                                int m_offset,
64                                int v_offset,
65                                float *mean,
66                                float *variance,
67                                int4 prefilter_rect,
68                                int buffer_pass_stride,
69                                int buffer_denoising_offset)
70 {
71         int x = prefilter_rect.x + blockDim.x*blockIdx.x + threadIdx.x;
72         int y = prefilter_rect.y + blockDim.y*blockIdx.y + threadIdx.y;
73         if(x < prefilter_rect.z && y < prefilter_rect.w) {
74                 kernel_filter_get_feature(sample,
75                                           tiles,
76                                           m_offset, v_offset,
77                                           x, y,
78                                           mean, variance,
79                                           prefilter_rect,
80                                           buffer_pass_stride,
81                                           buffer_denoising_offset);
82         }
83 }
84
85 extern "C" __global__ void
86 CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS)
87 kernel_cuda_filter_detect_outliers(float *image,
88                                    float *variance,
89                                    float *depth,
90                                    float *output,
91                                    int4 prefilter_rect,
92                                    int pass_stride)
93 {
94         int x = prefilter_rect.x + blockDim.x*blockIdx.x + threadIdx.x;
95         int y = prefilter_rect.y + blockDim.y*blockIdx.y + threadIdx.y;
96         if(x < prefilter_rect.z && y < prefilter_rect.w) {
97                 kernel_filter_detect_outliers(x, y, image, variance, depth, output, prefilter_rect, pass_stride);
98         }
99 }
100
101 extern "C" __global__ void
102 CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS)
103 kernel_cuda_filter_combine_halves(float *mean, float *variance, float *a, float *b, int4 prefilter_rect, int r)
104 {
105         int x = prefilter_rect.x + blockDim.x*blockIdx.x + threadIdx.x;
106         int y = prefilter_rect.y + blockDim.y*blockIdx.y + threadIdx.y;
107         if(x < prefilter_rect.z && y < prefilter_rect.w) {
108                 kernel_filter_combine_halves(x, y, mean, variance, a, b, prefilter_rect, r);
109         }
110 }
111
112 extern "C" __global__ void
113 CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS)
114 kernel_cuda_filter_construct_transform(float const* __restrict__ buffer,
115                                        float *transform, int *rank,
116                                        int4 filter_area, int4 rect,
117                                        int radius, float pca_threshold,
118                                        int pass_stride)
119 {
120         int x = blockDim.x*blockIdx.x + threadIdx.x;
121         int y = blockDim.y*blockIdx.y + threadIdx.y;
122         if(x < filter_area.z && y < filter_area.w) {
123                 int *l_rank = rank + y*filter_area.z + x;
124                 float *l_transform = transform + y*filter_area.z + x;
125                 kernel_filter_construct_transform(buffer,
126                                                   x + filter_area.x, y + filter_area.y,
127                                                   rect, pass_stride,
128                                                   l_transform, l_rank,
129                                                   radius, pca_threshold,
130                                                   filter_area.z*filter_area.w,
131                                                   threadIdx.y*blockDim.x + threadIdx.x);
132         }
133 }
134
135 extern "C" __global__ void
136 CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS)
137 kernel_cuda_filter_nlm_calc_difference(int dx, int dy,
138                                        const float *ccl_restrict weight_image,
139                                        const float *ccl_restrict variance_image,
140                                        float *difference_image,
141                                        int4 rect, int w,
142                                        int channel_offset,
143                                        float a, float k_2)
144 {
145         int x = blockDim.x*blockIdx.x + threadIdx.x + rect.x;
146         int y = blockDim.y*blockIdx.y + threadIdx.y + rect.y;
147         if(x < rect.z && y < rect.w) {
148                 kernel_filter_nlm_calc_difference(x, y, dx, dy, weight_image, variance_image, difference_image, rect, w, channel_offset, a, k_2);
149         }
150 }
151
152 extern "C" __global__ void
153 CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS)
154 kernel_cuda_filter_nlm_blur(const float *ccl_restrict difference_image, float *out_image, int4 rect, int w, int f)
155 {
156         int x = blockDim.x*blockIdx.x + threadIdx.x + rect.x;
157         int y = blockDim.y*blockIdx.y + threadIdx.y + rect.y;
158         if(x < rect.z && y < rect.w) {
159                 kernel_filter_nlm_blur(x, y, difference_image, out_image, rect, w, f);
160         }
161 }
162
163 extern "C" __global__ void
164 CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS)
165 kernel_cuda_filter_nlm_calc_weight(const float *ccl_restrict difference_image, float *out_image, int4 rect, int w, int f)
166 {
167         int x = blockDim.x*blockIdx.x + threadIdx.x + rect.x;
168         int y = blockDim.y*blockIdx.y + threadIdx.y + rect.y;
169         if(x < rect.z && y < rect.w) {
170                 kernel_filter_nlm_calc_weight(x, y, difference_image, out_image, rect, w, f);
171         }
172 }
173
174 extern "C" __global__ void
175 CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS)
176 kernel_cuda_filter_nlm_update_output(int dx, int dy,
177                                      const float *ccl_restrict difference_image,
178                                      const float *ccl_restrict image,
179                                      float *out_image, float *accum_image,
180                                      int4 rect, int w,
181                                      int f)
182 {
183         int x = blockDim.x*blockIdx.x + threadIdx.x + rect.x;
184         int y = blockDim.y*blockIdx.y + threadIdx.y + rect.y;
185         if(x < rect.z && y < rect.w) {
186                 kernel_filter_nlm_update_output(x, y, dx, dy, difference_image, image, out_image, accum_image, rect, w, f);
187         }
188 }
189
190 extern "C" __global__ void
191 CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS)
192 kernel_cuda_filter_nlm_normalize(float *out_image, const float *ccl_restrict accum_image, int4 rect, int w)
193 {
194         int x = blockDim.x*blockIdx.x + threadIdx.x + rect.x;
195         int y = blockDim.y*blockIdx.y + threadIdx.y + rect.y;
196         if(x < rect.z && y < rect.w) {
197                 kernel_filter_nlm_normalize(x, y, out_image, accum_image, rect, w);
198         }
199 }
200
201 extern "C" __global__ void
202 CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS)
203 kernel_cuda_filter_nlm_construct_gramian(int dx, int dy,
204                                          const float *ccl_restrict difference_image,
205                                          const float *ccl_restrict buffer,
206                                          float const* __restrict__ transform,
207                                          int *rank,
208                                          float *XtWX,
209                                          float3 *XtWY,
210                                          int4 rect,
211                                          int4 filter_rect,
212                                          int w, int h, int f,
213                                          int pass_stride)
214 {
215         int x = blockDim.x*blockIdx.x + threadIdx.x + max(0, rect.x-filter_rect.x);
216         int y = blockDim.y*blockIdx.y + threadIdx.y + max(0, rect.y-filter_rect.y);
217         if(x < min(filter_rect.z, rect.z-filter_rect.x) && y < min(filter_rect.w, rect.w-filter_rect.y)) {
218                 kernel_filter_nlm_construct_gramian(x, y,
219                                                     dx, dy,
220                                                     difference_image,
221                                                     buffer,
222                                                     transform, rank,
223                                                     XtWX, XtWY,
224                                                     rect, filter_rect,
225                                                     w, h, f,
226                                                     pass_stride,
227                                                     threadIdx.y*blockDim.x + threadIdx.x);
228         }
229 }
230
231 extern "C" __global__ void
232 CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS)
233 kernel_cuda_filter_finalize(int w, int h,
234                             float *buffer, int *rank,
235                             float *XtWX, float3 *XtWY,
236                             int4 filter_area, int4 buffer_params,
237                             int sample)
238 {
239         int x = blockDim.x*blockIdx.x + threadIdx.x;
240         int y = blockDim.y*blockIdx.y + threadIdx.y;
241         if(x < filter_area.z && y < filter_area.w) {
242                 int storage_ofs = y*filter_area.z+x;
243                 rank += storage_ofs;
244                 XtWX += storage_ofs;
245                 XtWY += storage_ofs;
246                 kernel_filter_finalize(x, y, w, h, buffer, rank, filter_area.z*filter_area.w, XtWX, XtWY, buffer_params, sample);
247         }
248 }
249
250 #endif
251