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
diff options
context:
space:
mode:
authorLukas Stockner <lukas.stockner@freenet.de>2018-07-04 15:26:42 +0300
committerLukas Stockner <lukas.stockner@freenet.de>2018-07-04 15:37:24 +0300
commit9db8bdbc653eb783707a748a271797510144a8eb (patch)
tree806f8c5d27790244de203ae62829b04db0be1610 /intern/cycles/kernel
parent97a0d6fcc736e604113487196ff3c3578306fc6c (diff)
Cycles Denoising: Cleanup: Rename tiles to tile_info
Diffstat (limited to 'intern/cycles/kernel')
-rw-r--r--intern/cycles/kernel/filter/filter_defines.h4
-rw-r--r--intern/cycles/kernel/filter/filter_prefilter.h20
-rw-r--r--intern/cycles/kernel/kernels/cpu/filter_cpu.h4
-rw-r--r--intern/cycles/kernel/kernels/cpu/filter_cpu_impl.h8
-rw-r--r--intern/cycles/kernel/kernels/cuda/filter.cu8
-rw-r--r--intern/cycles/kernel/kernels/opencl/filter.cl42
6 files changed, 43 insertions, 43 deletions
diff --git a/intern/cycles/kernel/filter/filter_defines.h b/intern/cycles/kernel/filter/filter_defines.h
index ce96f733aff..57d3d90594f 100644
--- a/intern/cycles/kernel/filter/filter_defines.h
+++ b/intern/cycles/kernel/filter/filter_defines.h
@@ -22,7 +22,7 @@
#define XTWX_SIZE (((DENOISE_FEATURES+1)*(DENOISE_FEATURES+2))/2)
#define XTWY_SIZE (DENOISE_FEATURES+1)
-typedef struct TilesInfo {
+typedef struct TileInfo {
int offsets[9];
int strides[9];
int x[4];
@@ -33,6 +33,6 @@ typedef struct TilesInfo {
#else
long long int buffers[9];
#endif
-} TilesInfo;
+} TileInfo;
#endif /* __FILTER_DEFINES_H__*/
diff --git a/intern/cycles/kernel/filter/filter_prefilter.h b/intern/cycles/kernel/filter/filter_prefilter.h
index 4af209341f6..9513bf46bd7 100644
--- a/intern/cycles/kernel/filter/filter_prefilter.h
+++ b/intern/cycles/kernel/filter/filter_prefilter.h
@@ -26,7 +26,7 @@ CCL_NAMESPACE_BEGIN
* bufferVariance: The buffer-based variance of the shadow feature. Unbiased, but quite noisy.
*/
ccl_device void kernel_filter_divide_shadow(int sample,
- ccl_global TilesInfo *tiles,
+ ccl_global TileInfo *tile_info,
int x, int y,
ccl_global float *unfilteredA,
ccl_global float *unfilteredB,
@@ -37,13 +37,13 @@ ccl_device void kernel_filter_divide_shadow(int sample,
int buffer_pass_stride,
int buffer_denoising_offset)
{
- int xtile = (x < tiles->x[1])? 0: ((x < tiles->x[2])? 1: 2);
- int ytile = (y < tiles->y[1])? 0: ((y < tiles->y[2])? 1: 2);
+ int xtile = (x < tile_info->x[1])? 0: ((x < tile_info->x[2])? 1: 2);
+ int ytile = (y < tile_info->y[1])? 0: ((y < tile_info->y[2])? 1: 2);
int tile = ytile*3+xtile;
- int offset = tiles->offsets[tile];
- int stride = tiles->strides[tile];
- const ccl_global float *ccl_restrict center_buffer = (ccl_global float*) tiles->buffers[tile];
+ int offset = tile_info->offsets[tile];
+ int stride = tile_info->strides[tile];
+ const ccl_global float *ccl_restrict center_buffer = (ccl_global float*) tile_info->buffers[tile];
center_buffer += (y*stride + x + offset)*buffer_pass_stride;
center_buffer += buffer_denoising_offset + 14;
@@ -79,7 +79,7 @@ ccl_device void kernel_filter_divide_shadow(int sample,
* - rect: The prefilter area (lower pixels inclusive, upper pixels exclusive).
*/
ccl_device void kernel_filter_get_feature(int sample,
- ccl_global TilesInfo *tiles,
+ ccl_global TileInfo *tile_info,
int m_offset, int v_offset,
int x, int y,
ccl_global float *mean,
@@ -87,10 +87,10 @@ ccl_device void kernel_filter_get_feature(int sample,
int4 rect, int buffer_pass_stride,
int buffer_denoising_offset)
{
- int xtile = (x < tiles->x[1])? 0: ((x < tiles->x[2])? 1: 2);
- int ytile = (y < tiles->y[1])? 0: ((y < tiles->y[2])? 1: 2);
+ int xtile = (x < tile_info->x[1])? 0: ((x < tile_info->x[2])? 1: 2);
+ int ytile = (y < tile_info->y[1])? 0: ((y < tile_info->y[2])? 1: 2);
int tile = ytile*3+xtile;
- ccl_global float *center_buffer = ((ccl_global float*) tiles->buffers[tile]) + (tiles->offsets[tile] + y*tiles->strides[tile] + x)*buffer_pass_stride + buffer_denoising_offset;
+ ccl_global float *center_buffer = ((ccl_global float*) tile_info->buffers[tile]) + (tile_info->offsets[tile] + y*tile_info->strides[tile] + x)*buffer_pass_stride + buffer_denoising_offset;
int buffer_w = align_up(rect.z - rect.x, 4);
int idx = (y-rect.y)*buffer_w + (x - rect.x);
diff --git a/intern/cycles/kernel/kernels/cpu/filter_cpu.h b/intern/cycles/kernel/kernels/cpu/filter_cpu.h
index 4231aba88d7..b62aa9663ec 100644
--- a/intern/cycles/kernel/kernels/cpu/filter_cpu.h
+++ b/intern/cycles/kernel/kernels/cpu/filter_cpu.h
@@ -17,7 +17,7 @@
/* Templated common declaration part of all CPU kernels. */
void KERNEL_FUNCTION_FULL_NAME(filter_divide_shadow)(int sample,
- TilesInfo *tiles,
+ TileInfo *tile_info,
int x,
int y,
float *unfilteredA,
@@ -30,7 +30,7 @@ void KERNEL_FUNCTION_FULL_NAME(filter_divide_shadow)(int sample,
int buffer_denoising_offset);
void KERNEL_FUNCTION_FULL_NAME(filter_get_feature)(int sample,
- TilesInfo *tiles,
+ TileInfo *tile_info,
int m_offset,
int v_offset,
int x,
diff --git a/intern/cycles/kernel/kernels/cpu/filter_cpu_impl.h b/intern/cycles/kernel/kernels/cpu/filter_cpu_impl.h
index 504622ecfd9..26777fdabb2 100644
--- a/intern/cycles/kernel/kernels/cpu/filter_cpu_impl.h
+++ b/intern/cycles/kernel/kernels/cpu/filter_cpu_impl.h
@@ -34,7 +34,7 @@ CCL_NAMESPACE_BEGIN
/* Denoise filter */
void KERNEL_FUNCTION_FULL_NAME(filter_divide_shadow)(int sample,
- TilesInfo *tiles,
+ TileInfo *tile_info,
int x,
int y,
float *unfilteredA,
@@ -49,7 +49,7 @@ void KERNEL_FUNCTION_FULL_NAME(filter_divide_shadow)(int sample,
#ifdef KERNEL_STUB
STUB_ASSERT(KERNEL_ARCH, filter_divide_shadow);
#else
- kernel_filter_divide_shadow(sample, tiles,
+ kernel_filter_divide_shadow(sample, tile_info,
x, y,
unfilteredA,
unfilteredB,
@@ -63,7 +63,7 @@ void KERNEL_FUNCTION_FULL_NAME(filter_divide_shadow)(int sample,
}
void KERNEL_FUNCTION_FULL_NAME(filter_get_feature)(int sample,
- TilesInfo *tiles,
+ TileInfo *tile_info,
int m_offset,
int v_offset,
int x,
@@ -76,7 +76,7 @@ void KERNEL_FUNCTION_FULL_NAME(filter_get_feature)(int sample,
#ifdef KERNEL_STUB
STUB_ASSERT(KERNEL_ARCH, filter_get_feature);
#else
- kernel_filter_get_feature(sample, tiles,
+ kernel_filter_get_feature(sample, tile_info,
m_offset, v_offset,
x, y,
mean, variance,
diff --git a/intern/cycles/kernel/kernels/cuda/filter.cu b/intern/cycles/kernel/kernels/cuda/filter.cu
index 035f0484488..0561c40e6b1 100644
--- a/intern/cycles/kernel/kernels/cuda/filter.cu
+++ b/intern/cycles/kernel/kernels/cuda/filter.cu
@@ -29,7 +29,7 @@
extern "C" __global__ void
CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS)
kernel_cuda_filter_divide_shadow(int sample,
- TilesInfo *tiles,
+ TileInfo *tile_info,
float *unfilteredA,
float *unfilteredB,
float *sampleVariance,
@@ -43,7 +43,7 @@ kernel_cuda_filter_divide_shadow(int sample,
int y = prefilter_rect.y + blockDim.y*blockIdx.y + threadIdx.y;
if(x < prefilter_rect.z && y < prefilter_rect.w) {
kernel_filter_divide_shadow(sample,
- tiles,
+ tile_info,
x, y,
unfilteredA,
unfilteredB,
@@ -59,7 +59,7 @@ kernel_cuda_filter_divide_shadow(int sample,
extern "C" __global__ void
CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS)
kernel_cuda_filter_get_feature(int sample,
- TilesInfo *tiles,
+ TileInfo *tile_info,
int m_offset,
int v_offset,
float *mean,
@@ -72,7 +72,7 @@ kernel_cuda_filter_get_feature(int sample,
int y = prefilter_rect.y + blockDim.y*blockIdx.y + threadIdx.y;
if(x < prefilter_rect.z && y < prefilter_rect.w) {
kernel_filter_get_feature(sample,
- tiles,
+ tile_info,
m_offset, v_offset,
x, y,
mean, variance,
diff --git a/intern/cycles/kernel/kernels/opencl/filter.cl b/intern/cycles/kernel/kernels/opencl/filter.cl
index 2b77807c38b..d553ee6833c 100644
--- a/intern/cycles/kernel/kernels/opencl/filter.cl
+++ b/intern/cycles/kernel/kernels/opencl/filter.cl
@@ -23,7 +23,7 @@
/* kernels */
__kernel void kernel_ocl_filter_divide_shadow(int sample,
- ccl_global TilesInfo *tiles,
+ ccl_global TileInfo *tile_info,
ccl_global float *unfilteredA,
ccl_global float *unfilteredB,
ccl_global float *sampleVariance,
@@ -51,7 +51,7 @@ __kernel void kernel_ocl_filter_divide_shadow(int sample,
}
__kernel void kernel_ocl_filter_get_feature(int sample,
- ccl_global TilesInfo *tiles,
+ ccl_global TileInfo *tile_info,
int m_offset,
int v_offset,
ccl_global float *mean,
@@ -277,26 +277,26 @@ __kernel void kernel_ocl_filter_finalize(ccl_global float *buffer,
}
}
-__kernel void kernel_ocl_filter_set_tiles(ccl_global TilesInfo* tiles,
- ccl_global float *buffer_1,
- ccl_global float *buffer_2,
- ccl_global float *buffer_3,
- ccl_global float *buffer_4,
- ccl_global float *buffer_5,
- ccl_global float *buffer_6,
- ccl_global float *buffer_7,
- ccl_global float *buffer_8,
- ccl_global float *buffer_9)
+__kernel void kernel_ocl_filter_set_tile_info(ccl_global TileInfo* tile_info,
+ ccl_global float *buffer_1,
+ ccl_global float *buffer_2,
+ ccl_global float *buffer_3,
+ ccl_global float *buffer_4,
+ ccl_global float *buffer_5,
+ ccl_global float *buffer_6,
+ ccl_global float *buffer_7,
+ ccl_global float *buffer_8,
+ ccl_global float *buffer_9)
{
if((get_global_id(0) == 0) && (get_global_id(1) == 0)) {
- tiles->buffers[0] = buffer_1;
- tiles->buffers[1] = buffer_2;
- tiles->buffers[2] = buffer_3;
- tiles->buffers[3] = buffer_4;
- tiles->buffers[4] = buffer_5;
- tiles->buffers[5] = buffer_6;
- tiles->buffers[6] = buffer_7;
- tiles->buffers[7] = buffer_8;
- tiles->buffers[8] = buffer_9;
+ tile_info->buffers[0] = buffer_1;
+ tile_info->buffers[1] = buffer_2;
+ tile_info->buffers[2] = buffer_3;
+ tile_info->buffers[3] = buffer_4;
+ tile_info->buffers[4] = buffer_5;
+ tile_info->buffers[5] = buffer_6;
+ tile_info->buffers[6] = buffer_7;
+ tile_info->buffers[7] = buffer_8;
+ tile_info->buffers[8] = buffer_9;
}
}