diff options
author | Sergey Sharybin <sergey.vfx@gmail.com> | 2017-05-19 13:52:12 +0300 |
---|---|---|
committer | Sergey Sharybin <sergey.vfx@gmail.com> | 2017-05-19 13:52:12 +0300 |
commit | 90a62404cb74be2d0601d8dd5abbce452c1a5c87 (patch) | |
tree | 0e6cf57f6830385e40f872ede7c3a23c9d7fa0de /intern | |
parent | 3a634524e39635cfee7b6ad13d5da7020f48912b (diff) |
Cycles: Cleanup, variable names
Don't use camel case for variable names. Leave that for the structures.
Diffstat (limited to 'intern')
-rw-r--r-- | intern/cycles/kernel/filter/filter_nlm_cpu.h | 58 | ||||
-rw-r--r-- | intern/cycles/kernel/filter/filter_nlm_gpu.h | 56 | ||||
-rw-r--r-- | intern/cycles/kernel/kernels/cpu/filter_cpu.h | 24 | ||||
-rw-r--r-- | intern/cycles/kernel/kernels/cpu/filter_cpu_impl.h | 36 | ||||
-rw-r--r-- | intern/cycles/kernel/kernels/cuda/filter.cu | 30 | ||||
-rw-r--r-- | intern/cycles/kernel/kernels/opencl/filter.cl | 38 |
6 files changed, 121 insertions, 121 deletions
diff --git a/intern/cycles/kernel/filter/filter_nlm_cpu.h b/intern/cycles/kernel/filter/filter_nlm_cpu.h index 57222811992..5cb4038bc33 100644 --- a/intern/cycles/kernel/filter/filter_nlm_cpu.h +++ b/intern/cycles/kernel/filter/filter_nlm_cpu.h @@ -17,9 +17,9 @@ CCL_NAMESPACE_BEGIN ccl_device_inline void kernel_filter_nlm_calc_difference(int dx, int dy, - const float *ccl_restrict weightImage, - const float *ccl_restrict varianceImage, - float *differenceImage, + const float *ccl_restrict weight_image, + const float *ccl_restrict variance_image, + float *difference_image, int4 rect, int w, int channel_offset, @@ -31,21 +31,21 @@ ccl_device_inline void kernel_filter_nlm_calc_difference(int dx, int dy, float diff = 0.0f; int numChannels = channel_offset? 3 : 1; for(int c = 0; c < numChannels; c++) { - float cdiff = weightImage[c*channel_offset + y*w+x] - weightImage[c*channel_offset + (y+dy)*w+(x+dx)]; - float pvar = varianceImage[c*channel_offset + y*w+x]; - float qvar = varianceImage[c*channel_offset + (y+dy)*w+(x+dx)]; + float cdiff = weight_image[c*channel_offset + y*w+x] - weight_image[c*channel_offset + (y+dy)*w+(x+dx)]; + float pvar = variance_image[c*channel_offset + y*w+x]; + float qvar = variance_image[c*channel_offset + (y+dy)*w+(x+dx)]; diff += (cdiff*cdiff - a*(pvar + min(pvar, qvar))) / (1e-8f + k_2*(pvar+qvar)); } if(numChannels > 1) { diff *= 1.0f/numChannels; } - differenceImage[y*w+x] = diff; + difference_image[y*w+x] = diff; } } } -ccl_device_inline void kernel_filter_nlm_blur(const float *ccl_restrict differenceImage, - float *outImage, +ccl_device_inline void kernel_filter_nlm_blur(const float *ccl_restrict difference_image, + float *out_image, int4 rect, int w, int f) @@ -58,34 +58,34 @@ ccl_device_inline void kernel_filter_nlm_blur(const float *ccl_restrict differen const int low = max(rect.y, y-f); const int high = min(rect.w, y+f+1); for(int x = rect.x; x < rect.z; x++) { - outImage[y*w+x] = 0.0f; + out_image[y*w+x] = 0.0f; } for(int y1 = low; y1 < high; y1++) { #ifdef __KERNEL_SSE3__ for(int x = aligned_lowx; x < aligned_highx; x+=4) { - _mm_store_ps(outImage + y*w+x, _mm_add_ps(_mm_load_ps(outImage + y*w+x), _mm_load_ps(differenceImage + y1*w+x))); + _mm_store_ps(out_image + y*w+x, _mm_add_ps(_mm_load_ps(out_image + y*w+x), _mm_load_ps(difference_image + y1*w+x))); } #else for(int x = rect.x; x < rect.z; x++) { - outImage[y*w+x] += differenceImage[y1*w+x]; + out_image[y*w+x] += difference_image[y1*w+x]; } #endif } for(int x = rect.x; x < rect.z; x++) { - outImage[y*w+x] *= 1.0f/(high - low); + out_image[y*w+x] *= 1.0f/(high - low); } } } -ccl_device_inline void kernel_filter_nlm_calc_weight(const float *ccl_restrict differenceImage, - float *outImage, +ccl_device_inline void kernel_filter_nlm_calc_weight(const float *ccl_restrict difference_image, + float *out_image, int4 rect, int w, int f) { for(int y = rect.y; y < rect.w; y++) { for(int x = rect.x; x < rect.z; x++) { - outImage[y*w+x] = 0.0f; + out_image[y*w+x] = 0.0f; } } for(int dx = -f; dx <= f; dx++) { @@ -93,7 +93,7 @@ ccl_device_inline void kernel_filter_nlm_calc_weight(const float *ccl_restrict d int neg_dx = min(0, dx); for(int y = rect.y; y < rect.w; y++) { for(int x = rect.x-neg_dx; x < rect.z-pos_dx; x++) { - outImage[y*w+x] += differenceImage[y*w+dx+x]; + out_image[y*w+x] += difference_image[y*w+dx+x]; } } } @@ -101,16 +101,16 @@ ccl_device_inline void kernel_filter_nlm_calc_weight(const float *ccl_restrict d for(int x = rect.x; x < rect.z; x++) { const int low = max(rect.x, x-f); const int high = min(rect.z, x+f+1); - outImage[y*w+x] = expf(-max(outImage[y*w+x] * (1.0f/(high - low)), 0.0f)); + out_image[y*w+x] = expf(-max(out_image[y*w+x] * (1.0f/(high - low)), 0.0f)); } } } ccl_device_inline void kernel_filter_nlm_update_output(int dx, int dy, - const float *ccl_restrict differenceImage, + const float *ccl_restrict difference_image, const float *ccl_restrict image, - float *outImage, - float *accumImage, + float *out_image, + float *accum_image, int4 rect, int w, int f) @@ -121,17 +121,17 @@ ccl_device_inline void kernel_filter_nlm_update_output(int dx, int dy, const int high = min(rect.z, x+f+1); float sum = 0.0f; for(int x1 = low; x1 < high; x1++) { - sum += differenceImage[y*w+x1]; + sum += difference_image[y*w+x1]; } float weight = sum * (1.0f/(high - low)); - accumImage[y*w+x] += weight; - outImage[y*w+x] += weight*image[(y+dy)*w+(x+dx)]; + accum_image[y*w+x] += weight; + out_image[y*w+x] += weight*image[(y+dy)*w+(x+dx)]; } } } ccl_device_inline void kernel_filter_nlm_construct_gramian(int dx, int dy, - const float *ccl_restrict differenceImage, + const float *ccl_restrict difference_image, const float *ccl_restrict buffer, float *color_pass, float *variance_pass, @@ -153,7 +153,7 @@ ccl_device_inline void kernel_filter_nlm_construct_gramian(int dx, int dy, const int high = min(rect.z, x+f+1); float sum = 0.0f; for(int x1 = low; x1 < high; x1++) { - sum += differenceImage[y*w+x1]; + sum += difference_image[y*w+x1]; } float weight = sum * (1.0f/(high - low)); @@ -174,14 +174,14 @@ ccl_device_inline void kernel_filter_nlm_construct_gramian(int dx, int dy, } } -ccl_device_inline void kernel_filter_nlm_normalize(float *outImage, - const float *ccl_restrict accumImage, +ccl_device_inline void kernel_filter_nlm_normalize(float *out_image, + const float *ccl_restrict accum_image, int4 rect, int w) { for(int y = rect.y; y < rect.w; y++) { for(int x = rect.x; x < rect.z; x++) { - outImage[y*w+x] /= accumImage[y*w+x]; + out_image[y*w+x] /= accum_image[y*w+x]; } } } diff --git a/intern/cycles/kernel/filter/filter_nlm_gpu.h b/intern/cycles/kernel/filter/filter_nlm_gpu.h index fd0a88340ea..078c5f56763 100644 --- a/intern/cycles/kernel/filter/filter_nlm_gpu.h +++ b/intern/cycles/kernel/filter/filter_nlm_gpu.h @@ -18,9 +18,9 @@ CCL_NAMESPACE_BEGIN ccl_device_inline void kernel_filter_nlm_calc_difference(int x, int y, int dx, int dy, - const ccl_global float *ccl_restrict weightImage, - const ccl_global float *ccl_restrict varianceImage, - ccl_global float *differenceImage, + const ccl_global float *ccl_restrict weight_image, + const ccl_global float *ccl_restrict variance_image, + ccl_global float *difference_image, int4 rect, int w, int channel_offset, float a, float k_2) @@ -28,74 +28,74 @@ ccl_device_inline void kernel_filter_nlm_calc_difference(int x, int y, float diff = 0.0f; int numChannels = channel_offset? 3 : 1; for(int c = 0; c < numChannels; c++) { - float cdiff = weightImage[c*channel_offset + y*w+x] - weightImage[c*channel_offset + (y+dy)*w+(x+dx)]; - float pvar = varianceImage[c*channel_offset + y*w+x]; - float qvar = varianceImage[c*channel_offset + (y+dy)*w+(x+dx)]; + float cdiff = weight_image[c*channel_offset + y*w+x] - weight_image[c*channel_offset + (y+dy)*w+(x+dx)]; + float pvar = variance_image[c*channel_offset + y*w+x]; + float qvar = variance_image[c*channel_offset + (y+dy)*w+(x+dx)]; diff += (cdiff*cdiff - a*(pvar + min(pvar, qvar))) / (1e-8f + k_2*(pvar+qvar)); } if(numChannels > 1) { diff *= 1.0f/numChannels; } - differenceImage[y*w+x] = diff; + difference_image[y*w+x] = diff; } ccl_device_inline void kernel_filter_nlm_blur(int x, int y, - const ccl_global float *ccl_restrict differenceImage, - ccl_global float *outImage, + const ccl_global float *ccl_restrict difference_image, + ccl_global float *out_image, int4 rect, int w, int f) { float sum = 0.0f; const int low = max(rect.y, y-f); const int high = min(rect.w, y+f+1); for(int y1 = low; y1 < high; y1++) { - sum += differenceImage[y1*w+x]; + sum += difference_image[y1*w+x]; } sum *= 1.0f/(high-low); - outImage[y*w+x] = sum; + out_image[y*w+x] = sum; } ccl_device_inline void kernel_filter_nlm_calc_weight(int x, int y, - const ccl_global float *ccl_restrict differenceImage, - ccl_global float *outImage, + const ccl_global float *ccl_restrict difference_image, + ccl_global float *out_image, int4 rect, int w, int f) { float sum = 0.0f; const int low = max(rect.x, x-f); const int high = min(rect.z, x+f+1); for(int x1 = low; x1 < high; x1++) { - sum += differenceImage[y*w+x1]; + sum += difference_image[y*w+x1]; } sum *= 1.0f/(high-low); - outImage[y*w+x] = expf(-max(sum, 0.0f)); + out_image[y*w+x] = expf(-max(sum, 0.0f)); } ccl_device_inline void kernel_filter_nlm_update_output(int x, int y, int dx, int dy, - const ccl_global float *ccl_restrict differenceImage, + const ccl_global float *ccl_restrict difference_image, const ccl_global float *ccl_restrict image, - ccl_global float *outImage, - ccl_global float *accumImage, + ccl_global float *out_image, + ccl_global float *accum_image, int4 rect, int w, int f) { float sum = 0.0f; const int low = max(rect.x, x-f); const int high = min(rect.z, x+f+1); for(int x1 = low; x1 < high; x1++) { - sum += differenceImage[y*w+x1]; + sum += difference_image[y*w+x1]; } sum *= 1.0f/(high-low); - if(outImage) { - accumImage[y*w+x] += sum; - outImage[y*w+x] += sum*image[(y+dy)*w+(x+dx)]; + if(out_image) { + accum_image[y*w+x] += sum; + out_image[y*w+x] += sum*image[(y+dy)*w+(x+dx)]; } else { - accumImage[y*w+x] = sum; + accum_image[y*w+x] = sum; } } ccl_device_inline void kernel_filter_nlm_construct_gramian(int fx, int fy, int dx, int dy, - const ccl_global float *ccl_restrict differenceImage, + const ccl_global float *ccl_restrict difference_image, const ccl_global float *ccl_restrict buffer, ccl_global float *color_pass, ccl_global float *variance_pass, @@ -115,7 +115,7 @@ ccl_device_inline void kernel_filter_nlm_construct_gramian(int fx, int fy, const int high = min(rect.z, x+f+1); float sum = 0.0f; for(int x1 = low; x1 < high; x1++) { - sum += differenceImage[y*w+x1]; + sum += difference_image[y*w+x1]; } float weight = sum * (1.0f/(high - low)); @@ -137,11 +137,11 @@ ccl_device_inline void kernel_filter_nlm_construct_gramian(int fx, int fy, } ccl_device_inline void kernel_filter_nlm_normalize(int x, int y, - ccl_global float *outImage, - const ccl_global float *ccl_restrict accumImage, + ccl_global float *out_image, + const ccl_global float *ccl_restrict accum_image, int4 rect, int w) { - outImage[y*w+x] /= accumImage[y*w+x]; + out_image[y*w+x] /= accum_image[y*w+x]; } CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/kernels/cpu/filter_cpu.h b/intern/cycles/kernel/kernels/cpu/filter_cpu.h index 9708b4b5b58..ffd34c293fc 100644 --- a/intern/cycles/kernel/kernels/cpu/filter_cpu.h +++ b/intern/cycles/kernel/kernels/cpu/filter_cpu.h @@ -72,40 +72,40 @@ void KERNEL_FUNCTION_FULL_NAME(filter_construct_transform)(float* buffer, void KERNEL_FUNCTION_FULL_NAME(filter_nlm_calc_difference)(int dx, int dy, - float *weightImage, + float *weight_image, float *variance, - float *differenceImage, + float *difference_image, int* rect, int w, int channel_offset, float a, float k_2); -void KERNEL_FUNCTION_FULL_NAME(filter_nlm_blur)(float *differenceImage, - float *outImage, +void KERNEL_FUNCTION_FULL_NAME(filter_nlm_blur)(float *difference_image, + float *out_image, int* rect, int w, int f); -void KERNEL_FUNCTION_FULL_NAME(filter_nlm_calc_weight)(float *differenceImage, - float *outImage, +void KERNEL_FUNCTION_FULL_NAME(filter_nlm_calc_weight)(float *difference_image, + float *out_image, int* rect, int w, int f); void KERNEL_FUNCTION_FULL_NAME(filter_nlm_update_output)(int dx, int dy, - float *differenceImage, + float *difference_image, float *image, - float *outImage, - float *accumImage, + float *out_image, + float *accum_image, int* rect, int w, int f); void KERNEL_FUNCTION_FULL_NAME(filter_nlm_construct_gramian)(int dx, int dy, - float *differenceImage, + float *difference_image, float *buffer, float *color_pass, float *variance_pass, @@ -120,8 +120,8 @@ void KERNEL_FUNCTION_FULL_NAME(filter_nlm_construct_gramian)(int dx, int f, int pass_stride); -void KERNEL_FUNCTION_FULL_NAME(filter_nlm_normalize)(float *outImage, - float *accumImage, +void KERNEL_FUNCTION_FULL_NAME(filter_nlm_normalize)(float *out_image, + float *accum_image, int* rect, int w); diff --git a/intern/cycles/kernel/kernels/cpu/filter_cpu_impl.h b/intern/cycles/kernel/kernels/cpu/filter_cpu_impl.h index 15325abdccd..261176846b1 100644 --- a/intern/cycles/kernel/kernels/cpu/filter_cpu_impl.h +++ b/intern/cycles/kernel/kernels/cpu/filter_cpu_impl.h @@ -150,9 +150,9 @@ void KERNEL_FUNCTION_FULL_NAME(filter_construct_transform)(float* buffer, void KERNEL_FUNCTION_FULL_NAME(filter_nlm_calc_difference)(int dx, int dy, - float *weightImage, + float *weight_image, float *variance, - float *differenceImage, + float *difference_image, int *rect, int w, int channel_offset, @@ -162,12 +162,12 @@ void KERNEL_FUNCTION_FULL_NAME(filter_nlm_calc_difference)(int dx, #ifdef KERNEL_STUB STUB_ASSERT(KERNEL_ARCH, filter_nlm_calc_difference); #else - kernel_filter_nlm_calc_difference(dx, dy, weightImage, variance, differenceImage, load_int4(rect), w, channel_offset, a, k_2); + kernel_filter_nlm_calc_difference(dx, dy, weight_image, variance, difference_image, load_int4(rect), w, channel_offset, a, k_2); #endif } -void KERNEL_FUNCTION_FULL_NAME(filter_nlm_blur)(float *differenceImage, - float *outImage, +void KERNEL_FUNCTION_FULL_NAME(filter_nlm_blur)(float *difference_image, + float *out_image, int *rect, int w, int f) @@ -175,12 +175,12 @@ void KERNEL_FUNCTION_FULL_NAME(filter_nlm_blur)(float *differenceImage, #ifdef KERNEL_STUB STUB_ASSERT(KERNEL_ARCH, filter_nlm_blur); #else - kernel_filter_nlm_blur(differenceImage, outImage, load_int4(rect), w, f); + kernel_filter_nlm_blur(difference_image, out_image, load_int4(rect), w, f); #endif } -void KERNEL_FUNCTION_FULL_NAME(filter_nlm_calc_weight)(float *differenceImage, - float *outImage, +void KERNEL_FUNCTION_FULL_NAME(filter_nlm_calc_weight)(float *difference_image, + float *out_image, int *rect, int w, int f) @@ -188,16 +188,16 @@ void KERNEL_FUNCTION_FULL_NAME(filter_nlm_calc_weight)(float *differenceImage, #ifdef KERNEL_STUB STUB_ASSERT(KERNEL_ARCH, filter_nlm_calc_weight); #else - kernel_filter_nlm_calc_weight(differenceImage, outImage, load_int4(rect), w, f); + kernel_filter_nlm_calc_weight(difference_image, out_image, load_int4(rect), w, f); #endif } void KERNEL_FUNCTION_FULL_NAME(filter_nlm_update_output)(int dx, int dy, - float *differenceImage, + float *difference_image, float *image, - float *outImage, - float *accumImage, + float *out_image, + float *accum_image, int *rect, int w, int f) @@ -205,13 +205,13 @@ void KERNEL_FUNCTION_FULL_NAME(filter_nlm_update_output)(int dx, #ifdef KERNEL_STUB STUB_ASSERT(KERNEL_ARCH, filter_nlm_update_output); #else - kernel_filter_nlm_update_output(dx, dy, differenceImage, image, outImage, accumImage, load_int4(rect), w, f); + kernel_filter_nlm_update_output(dx, dy, difference_image, image, out_image, accum_image, load_int4(rect), w, f); #endif } void KERNEL_FUNCTION_FULL_NAME(filter_nlm_construct_gramian)(int dx, int dy, - float *differenceImage, + float *difference_image, float *buffer, float *color_pass, float *variance_pass, @@ -229,19 +229,19 @@ void KERNEL_FUNCTION_FULL_NAME(filter_nlm_construct_gramian)(int dx, #ifdef KERNEL_STUB STUB_ASSERT(KERNEL_ARCH, filter_nlm_construct_gramian); #else - kernel_filter_nlm_construct_gramian(dx, dy, differenceImage, buffer, color_pass, variance_pass, transform, rank, XtWX, XtWY, load_int4(rect), load_int4(filter_rect), w, h, f, pass_stride); + kernel_filter_nlm_construct_gramian(dx, dy, difference_image, buffer, color_pass, variance_pass, transform, rank, XtWX, XtWY, load_int4(rect), load_int4(filter_rect), w, h, f, pass_stride); #endif } -void KERNEL_FUNCTION_FULL_NAME(filter_nlm_normalize)(float *outImage, - float *accumImage, +void KERNEL_FUNCTION_FULL_NAME(filter_nlm_normalize)(float *out_image, + float *accum_image, int *rect, int w) { #ifdef KERNEL_STUB STUB_ASSERT(KERNEL_ARCH, filter_nlm_normalize); #else - kernel_filter_nlm_normalize(outImage, accumImage, load_int4(rect), w); + kernel_filter_nlm_normalize(out_image, accum_image, load_int4(rect), w); #endif } diff --git a/intern/cycles/kernel/kernels/cuda/filter.cu b/intern/cycles/kernel/kernels/cuda/filter.cu index 2f560cf8d3c..2edbff08087 100644 --- a/intern/cycles/kernel/kernels/cuda/filter.cu +++ b/intern/cycles/kernel/kernels/cuda/filter.cu @@ -139,9 +139,9 @@ kernel_cuda_filter_construct_transform(float const* __restrict__ buffer, extern "C" __global__ void CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS) kernel_cuda_filter_nlm_calc_difference(int dx, int dy, - const float *ccl_restrict weightImage, - const float *ccl_restrict varianceImage, - float *differenceImage, + const float *ccl_restrict weight_image, + const float *ccl_restrict variance_image, + float *difference_image, int4 rect, int w, int channel_offset, float a, float k_2) @@ -149,63 +149,63 @@ kernel_cuda_filter_nlm_calc_difference(int dx, int dy, int x = blockDim.x*blockIdx.x + threadIdx.x + rect.x; int y = blockDim.y*blockIdx.y + threadIdx.y + 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); + kernel_filter_nlm_calc_difference(x, y, dx, dy, weight_image, variance_image, difference_image, rect, w, channel_offset, a, k_2); } } extern "C" __global__ void CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS) -kernel_cuda_filter_nlm_blur(const float *ccl_restrict differenceImage, float *outImage, int4 rect, int w, int f) +kernel_cuda_filter_nlm_blur(const float *ccl_restrict difference_image, float *out_image, int4 rect, int w, int f) { int x = blockDim.x*blockIdx.x + threadIdx.x + rect.x; int y = blockDim.y*blockIdx.y + threadIdx.y + rect.y; if(x < rect.z && y < rect.w) { - kernel_filter_nlm_blur(x, y, differenceImage, outImage, rect, w, f); + kernel_filter_nlm_blur(x, y, difference_image, out_image, rect, w, f); } } extern "C" __global__ void CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS) -kernel_cuda_filter_nlm_calc_weight(const float *ccl_restrict differenceImage, float *outImage, int4 rect, int w, int f) +kernel_cuda_filter_nlm_calc_weight(const float *ccl_restrict difference_image, float *out_image, int4 rect, int w, int f) { int x = blockDim.x*blockIdx.x + threadIdx.x + rect.x; int y = blockDim.y*blockIdx.y + threadIdx.y + rect.y; if(x < rect.z && y < rect.w) { - kernel_filter_nlm_calc_weight(x, y, differenceImage, outImage, rect, w, f); + kernel_filter_nlm_calc_weight(x, y, difference_image, out_image, rect, w, f); } } extern "C" __global__ void CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS) kernel_cuda_filter_nlm_update_output(int dx, int dy, - const float *ccl_restrict differenceImage, + const float *ccl_restrict difference_image, const float *ccl_restrict image, - float *outImage, float *accumImage, + float *out_image, float *accum_image, int4 rect, int w, int f) { int x = blockDim.x*blockIdx.x + threadIdx.x + rect.x; int y = blockDim.y*blockIdx.y + threadIdx.y + 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); + kernel_filter_nlm_update_output(x, y, dx, dy, difference_image, image, out_image, accum_image, rect, w, f); } } extern "C" __global__ void CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS) -kernel_cuda_filter_nlm_normalize(float *outImage, const float *ccl_restrict accumImage, int4 rect, int w) +kernel_cuda_filter_nlm_normalize(float *out_image, const float *ccl_restrict accum_image, int4 rect, int w) { int x = blockDim.x*blockIdx.x + threadIdx.x + rect.x; int y = blockDim.y*blockIdx.y + threadIdx.y + rect.y; if(x < rect.z && y < rect.w) { - kernel_filter_nlm_normalize(x, y, outImage, accumImage, rect, w); + kernel_filter_nlm_normalize(x, y, out_image, accum_image, rect, w); } } extern "C" __global__ void CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS) kernel_cuda_filter_nlm_construct_gramian(int dx, int dy, - const float *ccl_restrict differenceImage, + const float *ccl_restrict difference_image, const float *ccl_restrict buffer, float *color_pass, float *variance_pass, @@ -223,7 +223,7 @@ kernel_cuda_filter_nlm_construct_gramian(int dx, int dy, 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, + difference_image, buffer, color_pass, variance_pass, transform, rank, diff --git a/intern/cycles/kernel/kernels/opencl/filter.cl b/intern/cycles/kernel/kernels/opencl/filter.cl index 4621dbaffe1..0462ca6f9bc 100644 --- a/intern/cycles/kernel/kernels/opencl/filter.cl +++ b/intern/cycles/kernel/kernels/opencl/filter.cl @@ -132,9 +132,9 @@ __kernel void kernel_ocl_filter_construct_transform(const ccl_global float *ccl_ __kernel void kernel_ocl_filter_nlm_calc_difference(int dx, int dy, - const ccl_global float *ccl_restrict weightImage, - const ccl_global float *ccl_restrict varianceImage, - ccl_global float *differenceImage, + const ccl_global float *ccl_restrict weight_image, + const ccl_global float *ccl_restrict variance_image, + ccl_global float *difference_image, int4 rect, int w, int channel_offset, @@ -144,12 +144,12 @@ __kernel void kernel_ocl_filter_nlm_calc_difference(int dx, 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); + kernel_filter_nlm_calc_difference(x, y, dx, dy, weight_image, variance_image, difference_image, rect, w, channel_offset, a, k_2); } } -__kernel void kernel_ocl_filter_nlm_blur(const ccl_global float *ccl_restrict differenceImage, - ccl_global float *outImage, +__kernel void kernel_ocl_filter_nlm_blur(const ccl_global float *ccl_restrict difference_image, + ccl_global float *out_image, int4 rect, int w, int f) @@ -157,12 +157,12 @@ __kernel void kernel_ocl_filter_nlm_blur(const ccl_global float *ccl_restrict di 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); + kernel_filter_nlm_blur(x, y, difference_image, out_image, rect, w, f); } } -__kernel void kernel_ocl_filter_nlm_calc_weight(const ccl_global float *ccl_restrict differenceImage, - ccl_global float *outImage, +__kernel void kernel_ocl_filter_nlm_calc_weight(const ccl_global float *ccl_restrict difference_image, + ccl_global float *out_image, int4 rect, int w, int f) @@ -170,16 +170,16 @@ __kernel void kernel_ocl_filter_nlm_calc_weight(const ccl_global float *ccl_rest 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); + kernel_filter_nlm_calc_weight(x, y, difference_image, out_image, rect, w, f); } } __kernel void kernel_ocl_filter_nlm_update_output(int dx, int dy, - const ccl_global float *ccl_restrict differenceImage, + const ccl_global float *ccl_restrict difference_image, const ccl_global float *ccl_restrict image, - ccl_global float *outImage, - ccl_global float *accumImage, + ccl_global float *out_image, + ccl_global float *accum_image, int4 rect, int w, int f) @@ -187,25 +187,25 @@ __kernel void kernel_ocl_filter_nlm_update_output(int dx, 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); + kernel_filter_nlm_update_output(x, y, dx, dy, difference_image, image, out_image, accum_image, rect, w, f); } } -__kernel void kernel_ocl_filter_nlm_normalize(ccl_global float *outImage, - const ccl_global float *ccl_restrict accumImage, +__kernel void kernel_ocl_filter_nlm_normalize(ccl_global float *out_image, + const ccl_global float *ccl_restrict accum_image, 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_filter_nlm_normalize(x, y, out_image, accum_image, rect, w); } } __kernel void kernel_ocl_filter_nlm_construct_gramian(int dx, int dy, - const ccl_global float *ccl_restrict differenceImage, + const ccl_global float *ccl_restrict difference_image, const ccl_global float *ccl_restrict buffer, ccl_global float *color_pass, ccl_global float *variance_pass, @@ -225,7 +225,7 @@ __kernel void kernel_ocl_filter_nlm_construct_gramian(int dx, 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, + difference_image, buffer, color_pass, variance_pass, transform, rank, |