intern/cycles/kernel/kernels/opencl/filter.cl

   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 /* OpenCL kernel entry points */
  18
  19 #include "kernel/kernel_compat_opencl.h"
  20
  21 #include "kernel/filter/filter_kernel.h"
  22
  23 /* kernels */
  24
  25 __kernel void kernel_ocl_filter_divide_shadow(int sample,
  26                                               ccl_global TilesInfo *tiles,
  27                                               ccl_global float *unfilteredA,
  28                                               ccl_global float *unfilteredB,
  29                                               ccl_global float *sampleVariance,
  30                                               ccl_global float *sampleVarianceV,
  31                                               ccl_global float *bufferVariance,
  32                                               int4 prefilter_rect,
  33                                               int buffer_pass_stride,
  34                                               int buffer_denoising_offset,
  35                                               char use_split_variance)
  36 {
  37         int x = prefilter_rect.x + get_global_id(0);
  38         int y = prefilter_rect.y + get_global_id(1);
  39         if(x < prefilter_rect.z && y < prefilter_rect.w) {
  40                 kernel_filter_divide_shadow(sample,
  41                                             tiles,
  42                                             x, y,
  43                                             unfilteredA,
  44                                             unfilteredB,
  45                                             sampleVariance,
  46                                             sampleVarianceV,
  47                                             bufferVariance,
  48                                             prefilter_rect,
  49                                             buffer_pass_stride,
  50                                             buffer_denoising_offset,
  51                                             use_split_variance);
  52         }
  53 }
  54
  55 __kernel void kernel_ocl_filter_get_feature(int sample,
  56                                             ccl_global TilesInfo *tiles,
  57                                             int m_offset,
  58                                             int v_offset,
  59                                             ccl_global float *mean,
  60                                             ccl_global float *variance,
  61                                             int4 prefilter_rect,
  62                                             int buffer_pass_stride,
  63                                             int buffer_denoising_offset,
  64                                             char use_split_variance)
  65 {
  66         int x = prefilter_rect.x + get_global_id(0);
  67         int y = prefilter_rect.y + get_global_id(1);
  68         if(x < prefilter_rect.z && y < prefilter_rect.w) {
  69                 kernel_filter_get_feature(sample,
  70                                           tiles,
  71                                           m_offset, v_offset,
  72                                           x, y,
  73                                           mean, variance,
  74                                           prefilter_rect,
  75                                           buffer_pass_stride,
  76                                           buffer_denoising_offset,
  77                                           use_split_variance);
  78         }
  79 }
  80
  81 __kernel void kernel_ocl_filter_detect_outliers(ccl_global float *image,
  82                                                 ccl_global float *variance,
  83                                                 ccl_global float *depth,
  84                                                 ccl_global float *output,
  85                                                 int4 prefilter_rect,
  86                                                 int pass_stride)
  87 {
  88         int x = prefilter_rect.x + get_global_id(0);
  89         int y = prefilter_rect.y + get_global_id(1);
  90         if(x < prefilter_rect.z && y < prefilter_rect.w) {
  91                 kernel_filter_detect_outliers(x, y, image, variance, depth, output, prefilter_rect, pass_stride);
  92         }
  93 }
  94
  95 __kernel void kernel_ocl_filter_combine_halves(ccl_global float *mean,
  96                                                ccl_global float *variance,
  97                                                ccl_global float *a,
  98                                                ccl_global float *b,
  99                                                int4 prefilter_rect,
 100                                                int r)
 101 {
 102         int x = prefilter_rect.x + get_global_id(0);
 103         int y = prefilter_rect.y + get_global_id(1);
 104         if(x < prefilter_rect.z && y < prefilter_rect.w) {
 105                 kernel_filter_combine_halves(x, y, mean, variance, a, b, prefilter_rect, r);
 106         }
 107 }
 108
 109 __kernel void kernel_ocl_filter_construct_transform(const ccl_global float *ccl_restrict buffer,
 110                                                     ccl_global float *transform,
 111                                                     ccl_global int *rank,
 112                                                     int4 filter_area,
 113                                                     int4 rect,
 114                                                     int pass_stride,
 115                                                     int radius,
 116                                                     float pca_threshold)
 117 {
 118         int x = get_global_id(0);
 119         int y = get_global_id(1);
 120         if(x < filter_area.z && y < filter_area.w) {
 121                 ccl_global int *l_rank = rank + y*filter_area.z + x;
 122                 ccl_global float *l_transform = transform + y*filter_area.z + x;
 123                 kernel_filter_construct_transform(buffer,
 124                                                   x + filter_area.x, y + filter_area.y,
 125                                                   rect, pass_stride,
 126                                                   l_transform, l_rank,
 127                                                   radius, pca_threshold,
 128                                                   filter_area.z*filter_area.w,
 129                                                   get_local_id(1)*get_local_size(0) + get_local_id(0));
 130         }
 131 }
 132
 133 __kernel void kernel_ocl_filter_nlm_calc_difference(int dx,
 134                                                     int dy,
 135                                                     const ccl_global float *ccl_restrict weightImage,
 136                                                     const ccl_global float *ccl_restrict varianceImage,
 137                                                     ccl_global float *differenceImage,
 138                                                     int4 rect,
 139                                                     int w,
 140                                                     int channel_offset,
 141                                                     float a,
 142                                                     float k_2) {
 143         int x = get_global_id(0) + rect.x;
 144         int y = get_global_id(1) + rect.y;
 145         if(x < rect.z && y < rect.w) {
 146                 kernel_filter_nlm_calc_difference(x, y, dx, dy, weightImage, varianceImage, differenceImage, rect, w, channel_offset, a, k_2);
 147         }
 148 }
 149
 150 __kernel void kernel_ocl_filter_nlm_blur(const ccl_global float *ccl_restrict differenceImage,
 151                                          ccl_global float *outImage,
 152                                          int4 rect,
 153                                          int w,
 154                                          int f) {
 155         int x = get_global_id(0) + rect.x;
 156         int y = get_global_id(1) + rect.y;
 157         if(x < rect.z && y < rect.w) {
 158                 kernel_filter_nlm_blur(x, y, differenceImage, outImage, rect, w, f);
 159         }
 160 }
 161
 162 __kernel void kernel_ocl_filter_nlm_calc_weight(const ccl_global float *ccl_restrict differenceImage,
 163                                                 ccl_global float *outImage,
 164                                                 int4 rect,
 165                                                 int w,
 166                                                 int f) {
 167         int x = get_global_id(0) + rect.x;
 168         int y = get_global_id(1) + rect.y;
 169         if(x < rect.z && y < rect.w) {
 170                 kernel_filter_nlm_calc_weight(x, y, differenceImage, outImage, rect, w, f);
 171         }
 172 }
 173
 174 __kernel void kernel_ocl_filter_nlm_update_output(int dx,
 175                                                   int dy,
 176                                                   const ccl_global float *ccl_restrict differenceImage,
 177                                                   const ccl_global float *ccl_restrict image,
 178                                                   ccl_global float *outImage,
 179                                                   ccl_global float *accumImage,
 180                                                   int4 rect,
 181                                                   int w,
 182                                                   int f) {
 183         int x = get_global_id(0) + rect.x;
 184         int y = get_global_id(1) + rect.y;
 185         if(x < rect.z && y < rect.w) {
 186                 kernel_filter_nlm_update_output(x, y, dx, dy, differenceImage, image, outImage, accumImage, rect, w, f);
 187         }
 188 }
 189
 190 __kernel void kernel_ocl_filter_nlm_normalize(ccl_global float *outImage,
 191                                               const ccl_global float *ccl_restrict accumImage,
 192                                               int4 rect,
 193                                               int w) {
 194         int x = get_global_id(0) + rect.x;
 195         int y = get_global_id(1) + rect.y;
 196         if(x < rect.z && y < rect.w) {
 197                 kernel_filter_nlm_normalize(x, y, outImage, accumImage, rect, w);
 198         }
 199 }
 200
 201 __kernel void kernel_ocl_filter_nlm_construct_gramian(int dx,
 202                                                       int dy,
 203                                                       const ccl_global float *ccl_restrict differenceImage,
 204                                                       const ccl_global float *ccl_restrict buffer,
 205                                                       ccl_global float *color_pass,
 206                                                       ccl_global float *variance_pass,
 207                                                       const ccl_global float *ccl_restrict transform,
 208                                                       ccl_global int *rank,
 209                                                       ccl_global float *XtWX,
 210                                                       ccl_global float3 *XtWY,
 211                                                       int4 rect,
 212                                                       int4 filter_rect,
 213                                                       int w,
 214                                                       int h,
 215                                                       int f,
 216                                                       int pass_stride) {
 217         int x = get_global_id(0) + max(0, rect.x-filter_rect.x);
 218         int y = get_global_id(1) + max(0, rect.y-filter_rect.y);
 219         if(x < min(filter_rect.z, rect.z-filter_rect.x) && y < min(filter_rect.w, rect.w-filter_rect.y)) {
 220                 kernel_filter_nlm_construct_gramian(x, y,
 221                                                     dx, dy,
 222                                                     differenceImage,
 223                                                     buffer,
 224                                                     color_pass, variance_pass,
 225                                                     transform, rank,
 226                                                     XtWX, XtWY,
 227                                                     rect, filter_rect,
 228                                                     w, h, f,
 229                                                     pass_stride,
 230                                                     get_local_id(1)*get_local_size(0) + get_local_id(0));
 231         }
 232 }
 233
 234 __kernel void kernel_ocl_filter_finalize(int w,
 235                                              int h,
 236                                          ccl_global float *buffer,
 237                                          ccl_global int *rank,
 238                                          ccl_global float *XtWX,
 239                                          ccl_global float3 *XtWY,
 240                                          int4 filter_area,
 241                                          int4 buffer_params,
 242                                          int sample) {
 243         int x = get_global_id(0);
 244         int y = get_global_id(1);
 245         if(x < filter_area.z && y < filter_area.w) {
 246                 int storage_ofs = y*filter_area.z+x;
 247                 rank += storage_ofs;
 248                 XtWX += storage_ofs;
 249                 XtWY += storage_ofs;
 250                 kernel_filter_finalize(x, y, w, h, buffer, rank, filter_area.z*filter_area.w, XtWX, XtWY, buffer_params, sample);
 251         }
 252 }
 253
 254 __kernel void kernel_ocl_filter_set_tiles(ccl_global TilesInfo* tiles,
 255                                           ccl_global float *buffer_1,
 256                                           ccl_global float *buffer_2,
 257                                           ccl_global float *buffer_3,
 258                                           ccl_global float *buffer_4,
 259                                           ccl_global float *buffer_5,
 260                                           ccl_global float *buffer_6,
 261                                           ccl_global float *buffer_7,
 262                                           ccl_global float *buffer_8,
 263                                           ccl_global float *buffer_9)
 264 {
 265         if((get_global_id(0) == 0) && (get_global_id(1) == 0)) {
 266                 tiles->buffers[0] = buffer_1;
 267                 tiles->buffers[1] = buffer_2;
 268                 tiles->buffers[2] = buffer_3;
 269                 tiles->buffers[3] = buffer_4;
 270                 tiles->buffers[4] = buffer_5;
 271                 tiles->buffers[5] = buffer_6;
 272                 tiles->buffers[6] = buffer_7;
 273                 tiles->buffers[7] = buffer_8;
 274                 tiles->buffers[8] = buffer_9;
 275         }
 276 }