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/opencl/filter.cl | |
| 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/opencl/filter.cl')
| -rw-r--r-- | intern/cycles/kernel/kernels/opencl/filter.cl | 22 |
1 files changed, 11 insertions, 11 deletions
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, |