diff options
author | Sergey Sharybin <sergey.vfx@gmail.com> | 2017-05-19 13:33:28 +0300 |
---|---|---|
committer | Sergey Sharybin <sergey.vfx@gmail.com> | 2017-05-19 13:41:03 +0300 |
commit | 803337f3f64fed240e9adc6f286d5f9d13a5026a (patch) | |
tree | a262b427bd53873be9d18e952f193d66801edac8 /intern/cycles/kernel/kernels | |
parent | 8e655446d1ec667a08a6d351d1e452fc51f1428a (diff) |
\0;115;0cCycles: Cleanup, use ccl_restrict instead of ccl_restrict_ptr
There were following issues with ccl_restrict_ptr:
- We already had ccl_restrict for all platforms.
- It was secretly adding `const` qualifier to the declaration,
which is quite weird since non-const pointer can also be
declared as restricted.
- We never in Blender are using foo_ptr or FooPtr type definitions,
so not sure why we should introduce such a thing here.
- It is absolutely wrong from semantic point of view to put pointer
into the restrict macro -- const is a part of type, not part of
hint for compiler that some pointer is never aliased.
Diffstat (limited to 'intern/cycles/kernel/kernels')
-rw-r--r-- | intern/cycles/kernel/kernels/cuda/filter.cu | 18 | ||||
-rw-r--r-- | intern/cycles/kernel/kernels/opencl/filter.cl | 22 |
2 files changed, 20 insertions, 20 deletions
diff --git a/intern/cycles/kernel/kernels/cuda/filter.cu b/intern/cycles/kernel/kernels/cuda/filter.cu index f812a6601c6..30e1414f1e9 100644 --- a/intern/cycles/kernel/kernels/cuda/filter.cu +++ b/intern/cycles/kernel/kernels/cuda/filter.cu @@ -139,8 +139,8 @@ 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, - float ccl_restrict_ptr weightImage, - float ccl_restrict_ptr varianceImage, + const float *ccl_restrict weightImage, + const float *ccl_restrict varianceImage, float *differenceImage, int4 rect, int w, int channel_offset, @@ -154,7 +154,7 @@ kernel_cuda_filter_nlm_calc_difference(int dx, int dy, extern "C" __global__ void CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS) -kernel_cuda_filter_nlm_blur(float ccl_restrict_ptr differenceImage, float *outImage, int4 rect, int w, int f) { +kernel_cuda_filter_nlm_blur(const float *ccl_restrict differenceImage, float *outImage, 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) { @@ -164,7 +164,7 @@ kernel_cuda_filter_nlm_blur(float ccl_restrict_ptr differenceImage, float *outIm extern "C" __global__ void CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS) -kernel_cuda_filter_nlm_calc_weight(float ccl_restrict_ptr differenceImage, float *outImage, int4 rect, int w, int f) { +kernel_cuda_filter_nlm_calc_weight(const float *ccl_restrict differenceImage, float *outImage, 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) { @@ -175,8 +175,8 @@ kernel_cuda_filter_nlm_calc_weight(float ccl_restrict_ptr differenceImage, float 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, - float ccl_restrict_ptr differenceImage, - float ccl_restrict_ptr image, + const float *ccl_restrict differenceImage, + const float *ccl_restrict image, float *outImage, float *accumImage, int4 rect, int w, int f) { @@ -189,7 +189,7 @@ kernel_cuda_filter_nlm_update_output(int dx, int dy, extern "C" __global__ void CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS) -kernel_cuda_filter_nlm_normalize(float *outImage, float ccl_restrict_ptr accumImage, int4 rect, int w) { +kernel_cuda_filter_nlm_normalize(float *outImage, const float *ccl_restrict accumImage, 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) { @@ -200,8 +200,8 @@ kernel_cuda_filter_nlm_normalize(float *outImage, float ccl_restrict_ptr accumIm 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, - float ccl_restrict_ptr differenceImage, - float ccl_restrict_ptr buffer, + const float *ccl_restrict differenceImage, + const float *ccl_restrict buffer, float *color_pass, float *variance_pass, float const* __restrict__ transform, diff --git a/intern/cycles/kernel/kernels/opencl/filter.cl b/intern/cycles/kernel/kernels/opencl/filter.cl index fbc3daa62b9..f7d177b45b0 100644 --- a/intern/cycles/kernel/kernels/opencl/filter.cl +++ b/intern/cycles/kernel/kernels/opencl/filter.cl @@ -106,7 +106,7 @@ __kernel void kernel_ocl_filter_combine_halves(ccl_global float *mean, } } -__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, @@ -132,8 +132,8 @@ __kernel void kernel_ocl_filter_construct_transform(ccl_global float ccl_restric __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, + const ccl_global float *ccl_restrict weightImage, + const ccl_global float *ccl_restrict varianceImage, ccl_global float *differenceImage, int4 rect, int w, @@ -147,7 +147,7 @@ __kernel void kernel_ocl_filter_nlm_calc_difference(int dx, } } -__kernel void kernel_ocl_filter_nlm_blur(ccl_global float ccl_restrict_ptr differenceImage, +__kernel void kernel_ocl_filter_nlm_blur(const ccl_global float *ccl_restrict differenceImage, ccl_global float *outImage, int4 rect, int w, @@ -159,7 +159,7 @@ __kernel void kernel_ocl_filter_nlm_blur(ccl_global float ccl_restrict_ptr diffe } } -__kernel void kernel_ocl_filter_nlm_calc_weight(ccl_global float ccl_restrict_ptr differenceImage, +__kernel void kernel_ocl_filter_nlm_calc_weight(const ccl_global float *ccl_restrict differenceImage, ccl_global float *outImage, int4 rect, int w, @@ -173,8 +173,8 @@ __kernel void kernel_ocl_filter_nlm_calc_weight(ccl_global float ccl_restrict_pt __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, + const ccl_global float *ccl_restrict differenceImage, + const ccl_global float *ccl_restrict image, ccl_global float *outImage, ccl_global float *accumImage, int4 rect, @@ -188,7 +188,7 @@ __kernel void kernel_ocl_filter_nlm_update_output(int dx, } __kernel void kernel_ocl_filter_nlm_normalize(ccl_global float *outImage, - ccl_global float ccl_restrict_ptr accumImage, + const ccl_global float *ccl_restrict accumImage, int4 rect, int w) { int x = get_global_id(0) + rect.x; @@ -200,11 +200,11 @@ __kernel void kernel_ocl_filter_nlm_normalize(ccl_global float *outImage, __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, + const ccl_global float *ccl_restrict differenceImage, + const ccl_global float *ccl_restrict buffer, ccl_global float *color_pass, ccl_global float *variance_pass, - ccl_global float ccl_restrict_ptr transform, + const ccl_global float *ccl_restrict transform, ccl_global int *rank, ccl_global float *XtWX, ccl_global float3 *XtWY, |