ccl_global float *bufferVariance,
int4 prefilter_rect,
int buffer_pass_stride,
- int buffer_denoising_offset,
- char use_split_variance)
+ int buffer_denoising_offset)
{
int x = prefilter_rect.x + get_global_id(0);
int y = prefilter_rect.y + get_global_id(1);
bufferVariance,
prefilter_rect,
buffer_pass_stride,
- buffer_denoising_offset,
- use_split_variance);
+ buffer_denoising_offset);
}
}
ccl_global float *variance,
int4 prefilter_rect,
int buffer_pass_stride,
- int buffer_denoising_offset,
- char use_split_variance)
+ int buffer_denoising_offset)
{
int x = prefilter_rect.x + get_global_id(0);
int y = prefilter_rect.y + get_global_id(1);
mean, variance,
prefilter_rect,
buffer_pass_stride,
- buffer_denoising_offset,
- use_split_variance);
+ buffer_denoising_offset);
+ }
+}
+
+__kernel void kernel_ocl_filter_detect_outliers(ccl_global float *image,
+ ccl_global float *variance,
+ ccl_global float *depth,
+ ccl_global float *output,
+ int4 prefilter_rect,
+ int pass_stride)
+{
+ int x = prefilter_rect.x + get_global_id(0);
+ int y = prefilter_rect.y + get_global_id(1);
+ if(x < prefilter_rect.z && y < prefilter_rect.w) {
+ kernel_filter_detect_outliers(x, y, image, variance, depth, output, prefilter_rect, pass_stride);
}
}
}
}
-__kernel void kernel_ocl_filter_construct_transform(ccl_global float ccl_restrict_ptr buffer,
+__kernel void kernel_ocl_filter_construct_transform(const ccl_global float *ccl_restrict buffer,
ccl_global float *transform,
ccl_global int *rank,
int4 filter_area,
}
}
-__kernel void kernel_ocl_filter_nlm_calc_difference(int dx,
- int dy,
- ccl_global float ccl_restrict_ptr weightImage,
- ccl_global float ccl_restrict_ptr varianceImage,
- ccl_global float *differenceImage,
- int4 rect,
+__kernel void kernel_ocl_filter_nlm_calc_difference(const ccl_global float *ccl_restrict weight_image,
+ const ccl_global float *ccl_restrict variance_image,
+ ccl_global float *difference_image,
int w,
+ int h,
+ int stride,
+ int shift_stride,
+ int r,
int channel_offset,
float a,
- float k_2) {
- int x = get_global_id(0) + rect.x;
- int y = get_global_id(1) + rect.y;
- if(x < rect.z && y < rect.w) {
- kernel_filter_nlm_calc_difference(x, y, dx, dy, weightImage, varianceImage, differenceImage, rect, w, channel_offset, a, k_2);
+ float k_2)
+{
+ int4 co, rect;
+ int ofs;
+ if(get_nlm_coords(w, h, r, shift_stride, &rect, &co, &ofs)) {
+ kernel_filter_nlm_calc_difference(co.x, co.y, co.z, co.w,
+ weight_image,
+ variance_image,
+ difference_image + ofs,
+ rect, stride,
+ channel_offset, a, k_2);
}
}
-__kernel void kernel_ocl_filter_nlm_blur(ccl_global float ccl_restrict_ptr differenceImage,
- ccl_global float *outImage,
- int4 rect,
+__kernel void kernel_ocl_filter_nlm_blur(const ccl_global float *ccl_restrict difference_image,
+ ccl_global float *out_image,
int w,
- int f) {
- int x = get_global_id(0) + rect.x;
- int y = get_global_id(1) + rect.y;
- if(x < rect.z && y < rect.w) {
- kernel_filter_nlm_blur(x, y, differenceImage, outImage, rect, w, f);
+ int h,
+ int stride,
+ int shift_stride,
+ int r,
+ int f)
+{
+ int4 co, rect;
+ int ofs;
+ if(get_nlm_coords(w, h, r, shift_stride, &rect, &co, &ofs)) {
+ kernel_filter_nlm_blur(co.x, co.y,
+ difference_image + ofs,
+ out_image + ofs,
+ rect, stride, f);
}
}
-__kernel void kernel_ocl_filter_nlm_calc_weight(ccl_global float ccl_restrict_ptr differenceImage,
- ccl_global float *outImage,
- int4 rect,
+__kernel void kernel_ocl_filter_nlm_calc_weight(const ccl_global float *ccl_restrict difference_image,
+ ccl_global float *out_image,
int w,
- int f) {
- int x = get_global_id(0) + rect.x;
- int y = get_global_id(1) + rect.y;
- if(x < rect.z && y < rect.w) {
- kernel_filter_nlm_calc_weight(x, y, differenceImage, outImage, rect, w, f);
+ int h,
+ int stride,
+ int shift_stride,
+ int r,
+ int f)
+{
+ int4 co, rect;
+ int ofs;
+ if(get_nlm_coords(w, h, r, shift_stride, &rect, &co, &ofs)) {
+ kernel_filter_nlm_calc_weight(co.x, co.y,
+ difference_image + ofs,
+ out_image + ofs,
+ rect, stride, f);
}
}
-__kernel void kernel_ocl_filter_nlm_update_output(int dx,
- int dy,
- ccl_global float ccl_restrict_ptr differenceImage,
- ccl_global float ccl_restrict_ptr image,
- ccl_global float *outImage,
- ccl_global float *accumImage,
- int4 rect,
+__kernel void kernel_ocl_filter_nlm_update_output(const ccl_global float *ccl_restrict difference_image,
+ const ccl_global float *ccl_restrict image,
+ ccl_global float *out_image,
+ ccl_global float *accum_image,
int w,
- int f) {
- int x = get_global_id(0) + rect.x;
- int y = get_global_id(1) + rect.y;
- if(x < rect.z && y < rect.w) {
- kernel_filter_nlm_update_output(x, y, dx, dy, differenceImage, image, outImage, accumImage, rect, w, f);
+ int h,
+ int stride,
+ int shift_stride,
+ int r,
+ int f)
+{
+ int4 co, rect;
+ int ofs;
+ if(get_nlm_coords(w, h, r, shift_stride, &rect, &co, &ofs)) {
+ kernel_filter_nlm_update_output(co.x, co.y, co.z, co.w,
+ difference_image + ofs,
+ image,
+ out_image,
+ accum_image,
+ rect, stride, f);
}
}
-__kernel void kernel_ocl_filter_nlm_normalize(ccl_global float *outImage,
- ccl_global float ccl_restrict_ptr accumImage,
- int4 rect,
- int w) {
- int x = get_global_id(0) + rect.x;
- int y = get_global_id(1) + rect.y;
- if(x < rect.z && y < rect.w) {
- kernel_filter_nlm_normalize(x, y, outImage, accumImage, rect, w);
+__kernel void kernel_ocl_filter_nlm_normalize(ccl_global float *out_image,
+ const ccl_global float *ccl_restrict accum_image,
+ int w,
+ int h,
+ int stride)
+{
+ int x = get_global_id(0);
+ int y = get_global_id(1);
+ if(x < w && y < h) {
+ kernel_filter_nlm_normalize(x, y, out_image, accum_image, stride);
}
}
-__kernel void kernel_ocl_filter_nlm_construct_gramian(int dx,
- int dy,
- ccl_global float ccl_restrict_ptr differenceImage,
- ccl_global float ccl_restrict_ptr buffer,
- ccl_global float *color_pass,
- ccl_global float *variance_pass,
- ccl_global float ccl_restrict_ptr transform,
+__kernel void kernel_ocl_filter_nlm_construct_gramian(const ccl_global float *ccl_restrict difference_image,
+ const ccl_global float *ccl_restrict buffer,
+ const ccl_global float *ccl_restrict transform,
ccl_global int *rank,
ccl_global float *XtWX,
ccl_global float3 *XtWY,
- int4 rect,
- int4 filter_rect,
+ int4 filter_window,
int w,
int h,
+ int stride,
+ int shift_stride,
+ int r,
int f,
- int pass_stride) {
- int x = get_global_id(0) + max(0, rect.x-filter_rect.x);
- int y = get_global_id(1) + max(0, rect.y-filter_rect.y);
- if(x < min(filter_rect.z, rect.z-filter_rect.x) && y < min(filter_rect.w, rect.w-filter_rect.y)) {
- kernel_filter_nlm_construct_gramian(x, y,
- dx, dy,
- differenceImage,
+ int pass_stride)
+{
+ int4 co, rect;
+ int ofs;
+ if(get_nlm_coords_window(w, h, r, shift_stride, &rect, &co, &ofs, filter_window)) {
+ kernel_filter_nlm_construct_gramian(co.x, co.y,
+ co.z, co.w,
+ difference_image + ofs,
buffer,
- color_pass, variance_pass,
transform, rank,
XtWX, XtWY,
- rect, filter_rect,
- w, h, f,
+ rect, filter_window,
+ stride, f,
pass_stride,
get_local_id(1)*get_local_size(0) + get_local_id(0));
}
}
-__kernel void kernel_ocl_filter_finalize(int w,
- int h,
- ccl_global float *buffer,
+__kernel void kernel_ocl_filter_finalize(ccl_global float *buffer,
ccl_global int *rank,
ccl_global float *XtWX,
ccl_global float3 *XtWY,
int4 filter_area,
int4 buffer_params,
- int sample) {
+ int sample)
+{
int x = get_global_id(0);
int y = get_global_id(1);
if(x < filter_area.z && y < filter_area.w) {
rank += storage_ofs;
XtWX += storage_ofs;
XtWY += storage_ofs;
- kernel_filter_finalize(x, y, w, h, buffer, rank, filter_area.z*filter_area.w, XtWX, XtWY, buffer_params, sample);
+ kernel_filter_finalize(x, y, buffer, rank,
+ filter_area.z*filter_area.w,
+ XtWX, XtWY,
+ buffer_params, sample);
}
}
projects / blender.git / blobdiff