Welcome to mirror list, hosted at ThFree Co, Russian Federation.

git.blender.org/blender.git - Unnamed repository; edit this file 'description' to name the repository.
summaryrefslogtreecommitdiff
path: root/intern
diff options
context:
space:
mode:
authorSergey Sharybin <sergey.vfx@gmail.com>2017-05-19 13:52:12 +0300
committerSergey Sharybin <sergey.vfx@gmail.com>2017-05-19 13:52:12 +0300
commit90a62404cb74be2d0601d8dd5abbce452c1a5c87 (patch)
tree0e6cf57f6830385e40f872ede7c3a23c9d7fa0de /intern
parent3a634524e39635cfee7b6ad13d5da7020f48912b (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.h58
-rw-r--r--intern/cycles/kernel/filter/filter_nlm_gpu.h56
-rw-r--r--intern/cycles/kernel/kernels/cpu/filter_cpu.h24
-rw-r--r--intern/cycles/kernel/kernels/cpu/filter_cpu_impl.h36
-rw-r--r--intern/cycles/kernel/kernels/cuda/filter.cu30
-rw-r--r--intern/cycles/kernel/kernels/opencl/filter.cl38
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,