From 934394c5e2b1e142b4baeb7dc32a1b8cefaa93bc Mon Sep 17 00:00:00 2001 From: Tom Stellard Date: Tue, 4 Mar 2014 12:48:07 -0800 Subject: [PATCH] ocl: Don't use 'inline' attribute on functions In C99 'inline' is not a hint to the compiler to inline the function, it is an attribute that affects the linkage of the function. 'inline' functions are required to have a definition in a different compiliation unit, so compilers are free to delete 'inline' functions if they want to. This issue can be seen in Clang when compiling at -O0. Clang will sometimes delete 'inline' functions which creates an invalid program. Issue 3746: http://code.opencv.org/issues/3746 --- modules/ocl/src/opencl/bgfg_mog.cl | 22 ++++++++-------- modules/ocl/src/opencl/brute_force_match.cl | 10 +++---- modules/ocl/src/opencl/filtering_boxFilter.cl | 4 +-- modules/ocl/src/opencl/filtering_filter2D.cl | 4 +-- modules/ocl/src/opencl/imgproc_canny.cl | 10 +++---- modules/ocl/src/opencl/imgproc_clahe.cl | 8 +++--- modules/ocl/src/opencl/imgproc_warpAffine.cl | 4 +-- .../ocl/src/opencl/imgproc_warpPerspective.cl | 4 +-- modules/ocl/src/opencl/interpolate_frames.cl | 4 +-- .../src/opencl/kernel_radix_sort_by_key.cl | 4 +-- .../src/opencl/kernel_stablesort_by_key.cl | 6 ++--- modules/ocl/src/opencl/kmeans_kernel.cl | 4 +-- modules/ocl/src/opencl/match_template.cl | 6 ++--- modules/ocl/src/opencl/meanShift.cl | 4 +-- modules/ocl/src/opencl/objdetect_hog.cl | 4 +-- .../ocl/src/opencl/optical_flow_farneback.cl | 14 +++++----- modules/ocl/src/opencl/pyr_down.cl | 14 +++++----- modules/ocl/src/opencl/pyrlk.cl | 26 +++++++++---------- modules/ocl/src/opencl/stereobm.cl | 14 +++++----- modules/ocl/src/opencl/stereobp.cl | 14 +++++----- modules/ocl/src/opencl/stereocsbp.cl | 14 +++++----- modules/ocl/src/opencl/tvl1flow.cl | 8 +++--- 22 files changed, 101 insertions(+), 101 deletions(-) diff --git a/modules/ocl/src/opencl/bgfg_mog.cl b/modules/ocl/src/opencl/bgfg_mog.cl index a7479b929c..a3c2619abc 100644 --- a/modules/ocl/src/opencl/bgfg_mog.cl +++ b/modules/ocl/src/opencl/bgfg_mog.cl @@ -11,7 +11,7 @@ // For Open Source Computer Vision Library // // Copyright (C) 2010-2013, Multicoreware, Inc., all rights reserved. -// Copyright (C) 2010-2013, Advanced Micro Devices, Inc., all rights reserved. +// Copyright (C) 2010,2014, Advanced Micro Devices, Inc., all rights reserved. // Third party copyrights are property of their respective owners. // // @Authors @@ -48,22 +48,22 @@ #define T_MEAN_VAR float #define CONVERT_TYPE convert_uchar_sat #define F_ZERO (0.0f) -inline float cvt(uchar val) +float cvt(uchar val) { return val; } -inline float sqr(float val) +float sqr(float val) { return val * val; } -inline float sum(float val) +float sum(float val) { return val; } -inline float clamp1(float var, float learningRate, float diff, float minVar) +float clamp1(float var, float learningRate, float diff, float minVar) { return fmax(var + learningRate * (diff * diff - var), minVar); } @@ -75,7 +75,7 @@ inline float clamp1(float var, float learningRate, float diff, float minVar) #define CONVERT_TYPE convert_uchar4_sat #define F_ZERO (0.0f, 0.0f, 0.0f, 0.0f) -inline float4 cvt(const uchar4 val) +float4 cvt(const uchar4 val) { float4 result; result.x = val.x; @@ -86,17 +86,17 @@ inline float4 cvt(const uchar4 val) return result; } -inline float sqr(const float4 val) +float sqr(const float4 val) { return val.x * val.x + val.y * val.y + val.z * val.z; } -inline float sum(const float4 val) +float sum(const float4 val) { return (val.x + val.y + val.z); } -inline void swap4(__global float4* ptr, int x, int y, int k, int rows, int ptr_step) +void swap4(__global float4* ptr, int x, int y, int k, int rows, int ptr_step) { float4 val = ptr[(k * rows + y) * ptr_step + x]; ptr[(k * rows + y) * ptr_step + x] = ptr[((k + 1) * rows + y) * ptr_step + x]; @@ -104,7 +104,7 @@ inline void swap4(__global float4* ptr, int x, int y, int k, int rows, int ptr_s } -inline float4 clamp1(const float4 var, float learningRate, const float4 diff, float minVar) +float4 clamp1(const float4 var, float learningRate, const float4 diff, float minVar) { float4 result; result.x = fmax(var.x + learningRate * (diff.x * diff.x - var.x), minVar); @@ -128,7 +128,7 @@ typedef struct uchar c_shadowVal; } con_srtuct_t; -inline void swap(__global float* ptr, int x, int y, int k, int rows, int ptr_step) +void swap(__global float* ptr, int x, int y, int k, int rows, int ptr_step) { float val = ptr[(k * rows + y) * ptr_step + x]; ptr[(k * rows + y) * ptr_step + x] = ptr[((k + 1) * rows + y) * ptr_step + x]; diff --git a/modules/ocl/src/opencl/brute_force_match.cl b/modules/ocl/src/opencl/brute_force_match.cl index 2bf206a08b..5a54aa7ceb 100644 --- a/modules/ocl/src/opencl/brute_force_match.cl +++ b/modules/ocl/src/opencl/brute_force_match.cl @@ -11,7 +11,7 @@ // For Open Source Computer Vision Library // // Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. -// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Copyright (C) 2010,2014, Advanced Micro Devices, Inc., all rights reserved. // Third party copyrights are property of their respective owners. // // @Authors @@ -82,7 +82,7 @@ typedef float result_type; #define DIST_RES(x) sqrt(x) #elif (DIST_TYPE == 2) // Hamming //http://graphics.stanford.edu/~seander/bithacks.html#CountBitsSetParallel -inline int bit1Count(int v) +int bit1Count(int v) { v = v - ((v >> 1) & 0x55555555); // reuse input as temporary v = (v & 0x33333333) + ((v >> 2) & 0x33333333); // temp @@ -94,7 +94,7 @@ typedef int result_type; #define DIST_RES(x) (x) #endif -inline result_type reduce_block( +result_type reduce_block( __local value_type *s_query, __local value_type *s_train, int lidx, @@ -112,7 +112,7 @@ inline result_type reduce_block( return DIST_RES(result); } -inline result_type reduce_block_match( +result_type reduce_block_match( __local value_type *s_query, __local value_type *s_train, int lidx, @@ -130,7 +130,7 @@ inline result_type reduce_block_match( return (result); } -inline result_type reduce_multi_block( +result_type reduce_multi_block( __local value_type *s_query, __local value_type *s_train, int block_index, diff --git a/modules/ocl/src/opencl/filtering_boxFilter.cl b/modules/ocl/src/opencl/filtering_boxFilter.cl index 96091ce6e2..4145b050ae 100644 --- a/modules/ocl/src/opencl/filtering_boxFilter.cl +++ b/modules/ocl/src/opencl/filtering_boxFilter.cl @@ -10,7 +10,7 @@ // License Agreement // For Open Source Computer Vision Library // -// Copyright (C) 2010-2013, Advanced Micro Devices, Inc., all rights reserved. +// Copyright (C) 2010,2014, Advanced Micro Devices, Inc., all rights reserved. // Third party copyrights are property of their respective owners. // // Redistribution and use in source and binary forms, with or without modification, @@ -225,7 +225,7 @@ struct RectCoords #endif -inline INTERMEDIATE_TYPE readSrcPixel(int2 pos, __global TYPE *src, const unsigned int srcStepBytes, const struct RectCoords srcCoords +INTERMEDIATE_TYPE readSrcPixel(int2 pos, __global TYPE *src, const unsigned int srcStepBytes, const struct RectCoords srcCoords #ifdef BORDER_CONSTANT , SCALAR_TYPE borderValue #endif diff --git a/modules/ocl/src/opencl/filtering_filter2D.cl b/modules/ocl/src/opencl/filtering_filter2D.cl index fb7dca509e..684be0c4b6 100644 --- a/modules/ocl/src/opencl/filtering_filter2D.cl +++ b/modules/ocl/src/opencl/filtering_filter2D.cl @@ -10,7 +10,7 @@ // License Agreement // For Open Source Computer Vision Library // -// Copyright (C) 2010-2013, Advanced Micro Devices, Inc., all rights reserved. +// Copyright (C) 2010,2014, Advanced Micro Devices, Inc., all rights reserved. // Third party copyrights are property of their respective owners. // // Redistribution and use in source and binary forms, with or without modification, @@ -222,7 +222,7 @@ struct RectCoords #endif -inline INTERMEDIATE_TYPE readSrcPixel(int2 pos, __global TYPE *src, const unsigned int srcStepBytes, const struct RectCoords srcCoords +INTERMEDIATE_TYPE readSrcPixel(int2 pos, __global TYPE *src, const unsigned int srcStepBytes, const struct RectCoords srcCoords #ifdef BORDER_CONSTANT , SCALAR_TYPE borderValue #endif diff --git a/modules/ocl/src/opencl/imgproc_canny.cl b/modules/ocl/src/opencl/imgproc_canny.cl index 2ddfdae5f9..8f8d4d9d8c 100644 --- a/modules/ocl/src/opencl/imgproc_canny.cl +++ b/modules/ocl/src/opencl/imgproc_canny.cl @@ -11,7 +11,7 @@ // For Open Source Computer Vision Library // // Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. -// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Copyright (C) 2010,2014, Advanced Micro Devices, Inc., all rights reserved. // Third party copyrights are property of their respective owners. // // @Authors @@ -44,12 +44,12 @@ //M*/ #ifdef L2GRAD -inline float calc(int x, int y) +float calc(int x, int y) { return sqrt((float)(x * x + y * y)); } #else -inline float calc(int x, int y) +float calc(int x, int y) { return (float)abs(x) + abs(y); } @@ -381,8 +381,8 @@ struct PtrStepSz { int step; int rows, cols; }; -inline int get(struct PtrStepSz data, int y, int x) { return *((__global int *)((__global char*)data.ptr + data.step * (y + 1) + sizeof(int) * (x + 1))); } -inline void set(struct PtrStepSz data, int y, int x, int value) { *((__global int *)((__global char*)data.ptr + data.step * (y + 1) + sizeof(int) * (x + 1))) = value; } +int get(struct PtrStepSz data, int y, int x) { return *((__global int *)((__global char*)data.ptr + data.step * (y + 1) + sizeof(int) * (x + 1))); } +void set(struct PtrStepSz data, int y, int x, int value) { *((__global int *)((__global char*)data.ptr + data.step * (y + 1) + sizeof(int) * (x + 1))) = value; } ////////////////////////////////////////////////////////////////////////////////////////// // do Hysteresis for pixel whose edge type is 1 diff --git a/modules/ocl/src/opencl/imgproc_clahe.cl b/modules/ocl/src/opencl/imgproc_clahe.cl index 71a6f895d1..bb647ecf50 100644 --- a/modules/ocl/src/opencl/imgproc_clahe.cl +++ b/modules/ocl/src/opencl/imgproc_clahe.cl @@ -11,7 +11,7 @@ // For Open Source Computer Vision Library // // Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. -// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Copyright (C) 2010,2014, Advanced Micro Devices, Inc., all rights reserved. // Third party copyrights are property of their respective owners. // // @Authors @@ -47,7 +47,7 @@ #define WAVE_SIZE 1 #endif -inline int calc_lut(__local int* smem, int val, int tid) +int calc_lut(__local int* smem, int val, int tid) { smem[tid] = val; barrier(CLK_LOCAL_MEM_FENCE); @@ -61,7 +61,7 @@ inline int calc_lut(__local int* smem, int val, int tid) } #ifdef CPU -inline void reduce(volatile __local int* smem, int val, int tid) +void reduce(volatile __local int* smem, int val, int tid) { smem[tid] = val; barrier(CLK_LOCAL_MEM_FENCE); @@ -101,7 +101,7 @@ inline void reduce(volatile __local int* smem, int val, int tid) #else -inline void reduce(__local volatile int* smem, int val, int tid) +void reduce(__local volatile int* smem, int val, int tid) { smem[tid] = val; barrier(CLK_LOCAL_MEM_FENCE); diff --git a/modules/ocl/src/opencl/imgproc_warpAffine.cl b/modules/ocl/src/opencl/imgproc_warpAffine.cl index 27f99e005e..2ff911dedb 100644 --- a/modules/ocl/src/opencl/imgproc_warpAffine.cl +++ b/modules/ocl/src/opencl/imgproc_warpAffine.cl @@ -11,7 +11,7 @@ // For Open Source Computer Vision Library // // Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved. -// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Copyright (C) 2010,2014, Advanced Micro Devices, Inc., all rights reserved. // Third party copyrights are property of their respective owners. // // @Authors @@ -70,7 +70,7 @@ typedef float4 F4; #define INTER_REMAP_COEF_BITS 15 #define INTER_REMAP_COEF_SCALE (1 << INTER_REMAP_COEF_BITS) -inline void interpolateCubic( float x, float* coeffs ) +void interpolateCubic( float x, float* coeffs ) { const float A = -0.75f; diff --git a/modules/ocl/src/opencl/imgproc_warpPerspective.cl b/modules/ocl/src/opencl/imgproc_warpPerspective.cl index 97f86640be..009993bb7c 100644 --- a/modules/ocl/src/opencl/imgproc_warpPerspective.cl +++ b/modules/ocl/src/opencl/imgproc_warpPerspective.cl @@ -11,7 +11,7 @@ // For Open Source Computer Vision Library // // Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved. -// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Copyright (C) 2010,2014, Advanced Micro Devices, Inc., all rights reserved. // Third party copyrights are property of their respective owners. // // @Authors @@ -71,7 +71,7 @@ typedef float4 F4; #define INTER_REMAP_COEF_BITS 15 #define INTER_REMAP_COEF_SCALE (1 << INTER_REMAP_COEF_BITS) -inline void interpolateCubic( float x, float* coeffs ) +void interpolateCubic( float x, float* coeffs ) { const float A = -0.75f; diff --git a/modules/ocl/src/opencl/interpolate_frames.cl b/modules/ocl/src/opencl/interpolate_frames.cl index eb0b55f33b..a0e24df295 100644 --- a/modules/ocl/src/opencl/interpolate_frames.cl +++ b/modules/ocl/src/opencl/interpolate_frames.cl @@ -11,7 +11,7 @@ // For Open Source Computer Vision Library // // Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. -// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Copyright (C) 2010,2014, Advanced Micro Devices, Inc., all rights reserved. // Third party copyrights are property of their respective owners. // // @Authors @@ -50,7 +50,7 @@ __constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR; // atomic add for 32bit floating point -inline void atomic_addf(volatile __global float *source, const float operand) { +void atomic_addf(volatile __global float *source, const float operand) { union { unsigned int intVal; float floatVal; diff --git a/modules/ocl/src/opencl/kernel_radix_sort_by_key.cl b/modules/ocl/src/opencl/kernel_radix_sort_by_key.cl index 7e09f3fc57..0348d6fa30 100644 --- a/modules/ocl/src/opencl/kernel_radix_sort_by_key.cl +++ b/modules/ocl/src/opencl/kernel_radix_sort_by_key.cl @@ -11,7 +11,7 @@ // For Open Source Computer Vision Library // // Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. -// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Copyright (C) 2010,2014, Advanced Micro Devices, Inc., all rights reserved. // Third party copyrights are property of their respective owners. // // @Authors @@ -64,7 +64,7 @@ // from Thrust::b40c, link: // https://github.com/thrust/thrust/blob/master/thrust/system/cuda/detail/detail/b40c/radixsort_key_conversion.h -__inline uint convertKey(uint converted_key) +uint convertKey(uint converted_key) { #ifdef K_FLT unsigned int mask = (converted_key & 0x80000000) ? 0xffffffff : 0x80000000; diff --git a/modules/ocl/src/opencl/kernel_stablesort_by_key.cl b/modules/ocl/src/opencl/kernel_stablesort_by_key.cl index c573e3ebb3..2f4ae295b7 100644 --- a/modules/ocl/src/opencl/kernel_stablesort_by_key.cl +++ b/modules/ocl/src/opencl/kernel_stablesort_by_key.cl @@ -11,7 +11,7 @@ // For Open Source Computer Vision Library // // Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. -// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Copyright (C) 2010,2014, Advanced Micro Devices, Inc., all rights reserved. // Third party copyrights are property of their respective owners. // // @Authors @@ -65,7 +65,7 @@ // by a base pointer and left and right index for a particular candidate value. The comparison operator is // passed as a functor parameter my_comp // This function returns an index that is the first index whos value would be equal to the searched value -inline uint lowerBoundBinary( global K_T* data, uint left, uint right, K_T searchVal) +uint lowerBoundBinary( global K_T* data, uint left, uint right, K_T searchVal) { // The values firstIndex and lastIndex get modified within the loop, narrowing down the potential sequence uint firstIndex = left; @@ -101,7 +101,7 @@ inline uint lowerBoundBinary( global K_T* data, uint left, uint right, K_T searc // passed as a functor parameter my_comp // This function returns an index that is the first index whos value would be greater than the searched value // If the search value is not found in the sequence, upperbound returns the same result as lowerbound -inline uint upperBoundBinary( global K_T* data, uint left, uint right, K_T searchVal) +uint upperBoundBinary( global K_T* data, uint left, uint right, K_T searchVal) { uint upperBound = lowerBoundBinary( data, left, right, searchVal ); diff --git a/modules/ocl/src/opencl/kmeans_kernel.cl b/modules/ocl/src/opencl/kmeans_kernel.cl index bb0e9c9a41..5e6fa8a706 100644 --- a/modules/ocl/src/opencl/kmeans_kernel.cl +++ b/modules/ocl/src/opencl/kmeans_kernel.cl @@ -11,7 +11,7 @@ // For Open Source Computer Vision Library // // Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. -// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Copyright (C) 2010,2014, Advanced Micro Devices, Inc., all rights reserved. // Third party copyrights are property of their respective owners. // // @Authors @@ -44,7 +44,7 @@ // //M*/ -inline float distance_(__global const float * center, __global const float * src, int feature_length) +float distance_(__global const float * center, __global const float * src, int feature_length) { float res = 0; float4 v0, v1, v2; diff --git a/modules/ocl/src/opencl/match_template.cl b/modules/ocl/src/opencl/match_template.cl index 4d46d0084c..a68be6802a 100644 --- a/modules/ocl/src/opencl/match_template.cl +++ b/modules/ocl/src/opencl/match_template.cl @@ -11,7 +11,7 @@ // For Open Source Computer Vision Library // // Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. -// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Copyright (C) 2010,2014, Advanced Micro Devices, Inc., all rights reserved. // Third party copyrights are property of their respective owners. // // @Authors @@ -66,7 +66,7 @@ #define SUMS_PTR(ox, oy) mad24(gidy + oy, img_sums_step, gidx + img_sums_offset + ox) // normAcc* are accurate normalization routines which make GPU matchTemplate // consistent with CPU one -inline float normAcc(float num, float denum) +float normAcc(float num, float denum) { if(fabs(num) < denum) { @@ -79,7 +79,7 @@ inline float normAcc(float num, float denum) return 0; } -inline float normAcc_SQDIFF(float num, float denum) +float normAcc_SQDIFF(float num, float denum) { if(fabs(num) < denum) { diff --git a/modules/ocl/src/opencl/meanShift.cl b/modules/ocl/src/opencl/meanShift.cl index 3fff473a83..a1986670c0 100644 --- a/modules/ocl/src/opencl/meanShift.cl +++ b/modules/ocl/src/opencl/meanShift.cl @@ -11,7 +11,7 @@ // For Open Source Computer Vision Library // // Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved. -// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Copyright (C) 2010,2014, Advanced Micro Devices, Inc., all rights reserved. // Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. // Third party copyrights are property of their respective owners. // @@ -46,7 +46,7 @@ // //M*/ -inline short2 do_mean_shift(int x0, int y0, __global uchar4* out,int out_step, +short2 do_mean_shift(int x0, int y0, __global uchar4* out,int out_step, __global uchar4* in, int in_step, int dst_off, int src_off, int cols, int rows, int sp, int sr, int maxIter, float eps) { diff --git a/modules/ocl/src/opencl/objdetect_hog.cl b/modules/ocl/src/opencl/objdetect_hog.cl index e931e82b57..1a0ceb21ba 100644 --- a/modules/ocl/src/opencl/objdetect_hog.cl +++ b/modules/ocl/src/opencl/objdetect_hog.cl @@ -11,7 +11,7 @@ // For Open Source Computer Vision Library // // Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. -// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Copyright (C) 2010,2014, Advanced Micro Devices, Inc., all rights reserved. // Third party copyrights are property of their respective owners. // // @Authors @@ -208,7 +208,7 @@ __kernel void normalize_hists_36_kernel(__global float* block_hists, //------------------------------------------------------------- // Normalization of histograms via L2Hys_norm // -inline float reduce_smem(volatile __local float* smem, int size) +float reduce_smem(volatile __local float* smem, int size) { unsigned int tid = get_local_id(0); float sum = smem[tid]; diff --git a/modules/ocl/src/opencl/optical_flow_farneback.cl b/modules/ocl/src/opencl/optical_flow_farneback.cl index 4725662c60..8967b2f3b7 100644 --- a/modules/ocl/src/opencl/optical_flow_farneback.cl +++ b/modules/ocl/src/opencl/optical_flow_farneback.cl @@ -11,7 +11,7 @@ // For Open Source Computer Vision Library // // Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. -// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Copyright (C) 2010,2014, Advanced Micro Devices, Inc., all rights reserved. // Third party copyrights are property of their respective owners. // // @Authors @@ -123,32 +123,32 @@ __kernel void polynomialExpansion(__global float * dst, } } -inline int idx_row_low(const int y, const int last_row) +int idx_row_low(const int y, const int last_row) { return abs(y) % (last_row + 1); } -inline int idx_row_high(const int y, const int last_row) +int idx_row_high(const int y, const int last_row) { return abs(last_row - abs(last_row - y)) % (last_row + 1); } -inline int idx_row(const int y, const int last_row) +int idx_row(const int y, const int last_row) { return idx_row_low(idx_row_high(y, last_row), last_row); } -inline int idx_col_low(const int x, const int last_col) +int idx_col_low(const int x, const int last_col) { return abs(x) % (last_col + 1); } -inline int idx_col_high(const int x, const int last_col) +int idx_col_high(const int x, const int last_col) { return abs(last_col - abs(last_col - x)) % (last_col + 1); } -inline int idx_col(const int x, const int last_col) +int idx_col(const int x, const int last_col) { return idx_col_low(idx_col_high(x, last_col), last_col); } diff --git a/modules/ocl/src/opencl/pyr_down.cl b/modules/ocl/src/opencl/pyr_down.cl index 6f10067e9f..0025f7c76a 100644 --- a/modules/ocl/src/opencl/pyr_down.cl +++ b/modules/ocl/src/opencl/pyr_down.cl @@ -11,7 +11,7 @@ // For Open Source Computer Vision Library // // Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. -// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Copyright (C) 2010,2014, Advanced Micro Devices, Inc., all rights reserved. // Third party copyrights are property of their respective owners. // // @Authors @@ -43,32 +43,32 @@ // //M*/ -inline int idx_row_low(int y, int last_row) +int idx_row_low(int y, int last_row) { return abs(y) % (last_row + 1); } -inline int idx_row_high(int y, int last_row) +int idx_row_high(int y, int last_row) { return abs(last_row - (int)abs(last_row - y)) % (last_row + 1); } -inline int idx_row(int y, int last_row) +int idx_row(int y, int last_row) { return idx_row_low(idx_row_high(y, last_row), last_row); } -inline int idx_col_low(int x, int last_col) +int idx_col_low(int x, int last_col) { return abs(x) % (last_col + 1); } -inline int idx_col_high(int x, int last_col) +int idx_col_high(int x, int last_col) { return abs(last_col - (int)abs(last_col - x)) % (last_col + 1); } -inline int idx_col(int x, int last_col) +int idx_col(int x, int last_col) { return idx_col_low(idx_col_high(x, last_col), last_col); } diff --git a/modules/ocl/src/opencl/pyrlk.cl b/modules/ocl/src/opencl/pyrlk.cl index f34aee9009..8afdc8708e 100644 --- a/modules/ocl/src/opencl/pyrlk.cl +++ b/modules/ocl/src/opencl/pyrlk.cl @@ -11,7 +11,7 @@ // For Open Source Computer Vision Library // // Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. -// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Copyright (C) 2010,2014, Advanced Micro Devices, Inc., all rights reserved. // Third party copyrights are property of their respective owners. // // @Authors @@ -52,7 +52,7 @@ #endif #ifdef CPU -inline void reduce3(float val1, float val2, float val3, __local float* smem1, __local float* smem2, __local float* smem3, int tid) +void reduce3(float val1, float val2, float val3, __local float* smem1, __local float* smem2, __local float* smem3, int tid) { smem1[tid] = val1; smem2[tid] = val2; @@ -71,7 +71,7 @@ inline void reduce3(float val1, float val2, float val3, __local float* smem1, } } -inline void reduce2(float val1, float val2, volatile __local float* smem1, volatile __local float* smem2, int tid) +void reduce2(float val1, float val2, volatile __local float* smem1, volatile __local float* smem2, int tid) { smem1[tid] = val1; smem2[tid] = val2; @@ -88,7 +88,7 @@ inline void reduce2(float val1, float val2, volatile __local float* smem1, volat } } -inline void reduce1(float val1, volatile __local float* smem1, int tid) +void reduce1(float val1, volatile __local float* smem1, int tid) { smem1[tid] = val1; barrier(CLK_LOCAL_MEM_FENCE); @@ -103,7 +103,7 @@ inline void reduce1(float val1, volatile __local float* smem1, int tid) } } #else -inline void reduce3(float val1, float val2, float val3, +void reduce3(float val1, float val2, float val3, __local volatile float* smem1, __local volatile float* smem2, __local volatile float* smem3, int tid) { smem1[tid] = val1; @@ -150,7 +150,7 @@ inline void reduce3(float val1, float val2, float val3, barrier(CLK_LOCAL_MEM_FENCE); } -inline void reduce2(float val1, float val2, __local volatile float* smem1, __local volatile float* smem2, int tid) +void reduce2(float val1, float val2, __local volatile float* smem1, __local volatile float* smem2, int tid) { smem1[tid] = val1; smem2[tid] = val2; @@ -189,7 +189,7 @@ inline void reduce2(float val1, float val2, __local volatile float* smem1, __loc barrier(CLK_LOCAL_MEM_FENCE); } -inline void reduce1(float val1, __local volatile float* smem1, int tid) +void reduce1(float val1, __local volatile float* smem1, int tid) { smem1[tid] = val1; barrier(CLK_LOCAL_MEM_FENCE); @@ -225,7 +225,7 @@ inline void reduce1(float val1, __local volatile float* smem1, int tid) // Image read mode __constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR; -inline void SetPatch(image2d_t I, float x, float y, +void SetPatch(image2d_t I, float x, float y, float* Pch, float* Dx, float* Dy, float* A11, float* A12, float* A22) { @@ -246,7 +246,7 @@ inline void SetPatch(image2d_t I, float x, float y, *A22 += dIdy * dIdy; } -inline void GetPatch(image2d_t J, float x, float y, +void GetPatch(image2d_t J, float x, float y, float* Pch, float* Dx, float* Dy, float* b1, float* b2) { @@ -256,13 +256,13 @@ inline void GetPatch(image2d_t J, float x, float y, *b2 += diff**Dy; } -inline void GetError(image2d_t J, const float x, const float y, const float* Pch, float* errval) +void GetError(image2d_t J, const float x, const float y, const float* Pch, float* errval) { float diff = read_imagef(J, sampler, (float2)(x,y)).x-*Pch; *errval += fabs(diff); } -inline void SetPatch4(image2d_t I, const float x, const float y, +void SetPatch4(image2d_t I, const float x, const float y, float4* Pch, float4* Dx, float4* Dy, float* A11, float* A12, float* A22) { @@ -285,7 +285,7 @@ inline void SetPatch4(image2d_t I, const float x, const float y, *A22 += sqIdx.x + sqIdx.y + sqIdx.z; } -inline void GetPatch4(image2d_t J, const float x, const float y, +void GetPatch4(image2d_t J, const float x, const float y, const float4* Pch, const float4* Dx, const float4* Dy, float* b1, float* b2) { @@ -297,7 +297,7 @@ inline void GetPatch4(image2d_t J, const float x, const float y, *b2 += xdiff.x + xdiff.y + xdiff.z; } -inline void GetError4(image2d_t J, const float x, const float y, const float4* Pch, float* errval) +void GetError4(image2d_t J, const float x, const float y, const float4* Pch, float* errval) { float4 diff = read_imagef(J, sampler, (float2)(x,y))-*Pch; *errval += fabs(diff.x) + fabs(diff.y) + fabs(diff.z); diff --git a/modules/ocl/src/opencl/stereobm.cl b/modules/ocl/src/opencl/stereobm.cl index d3efb5eb4c..3ba165249f 100644 --- a/modules/ocl/src/opencl/stereobm.cl +++ b/modules/ocl/src/opencl/stereobm.cl @@ -11,7 +11,7 @@ // For Open Source Computer Vision Library // // Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved. -// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Copyright (C) 2010,2014, Advanced Micro Devices, Inc., all rights reserved. // Third party copyrights are property of their respective owners. // // @Authors @@ -56,7 +56,7 @@ #define radius 64 #endif -inline unsigned int CalcSSD(__local unsigned int *col_ssd) +unsigned int CalcSSD(__local unsigned int *col_ssd) { unsigned int cache = col_ssd[0]; @@ -67,7 +67,7 @@ inline unsigned int CalcSSD(__local unsigned int *col_ssd) return cache; } -inline uint2 MinSSD(__local unsigned int *col_ssd) +uint2 MinSSD(__local unsigned int *col_ssd) { unsigned int ssd[N_DISPARITIES]; const int win_size = (radius << 1); @@ -95,7 +95,7 @@ inline uint2 MinSSD(__local unsigned int *col_ssd) return (uint2)(mssd, bestIdx); } -inline void StepDown(int idx1, int idx2, __global unsigned char* imageL, +void StepDown(int idx1, int idx2, __global unsigned char* imageL, __global unsigned char* imageR, int d, __local unsigned int *col_ssd) { uint8 imgR1 = convert_uint8(vload8(0, imageR + (idx1 - d - 7))); @@ -114,7 +114,7 @@ inline void StepDown(int idx1, int idx2, __global unsigned char* imageL, col_ssd[7 * (BLOCK_W + win_size)] += res.s0; } -inline void InitColSSD(int x_tex, int y_tex, int im_pitch, __global unsigned char* imageL, +void InitColSSD(int x_tex, int y_tex, int im_pitch, __global unsigned char* imageL, __global unsigned char* imageR, int d, __local unsigned int *col_ssd) { @@ -241,7 +241,7 @@ __kernel void prefilter_xsobel(__global unsigned char *input, __global unsigned /////////////////////////////////// Textureness filtering //////////////////////////////////////// ////////////////////////////////////////////////////////////////////////////////////////////////// -inline float sobel(__global unsigned char *input, int x, int y, int rows, int cols) +float sobel(__global unsigned char *input, int x, int y, int rows, int cols) { float conv = 0; int y1 = y==0? 0 : y-1; @@ -256,7 +256,7 @@ inline float sobel(__global unsigned char *input, int x, int y, int rows, int co return fabs(conv); } -inline float CalcSums(__local float *cols, __local float *cols_cache, int winsz) +float CalcSums(__local float *cols, __local float *cols_cache, int winsz) { unsigned int cache = cols[0]; diff --git a/modules/ocl/src/opencl/stereobp.cl b/modules/ocl/src/opencl/stereobp.cl index 5a1bf088c9..0fe00521df 100644 --- a/modules/ocl/src/opencl/stereobp.cl +++ b/modules/ocl/src/opencl/stereobp.cl @@ -11,7 +11,7 @@ // For Open Source Computer Vision Library // Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. // Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved. -// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Copyright (C) 2010,2014, Advanced Micro Devices, Inc., all rights reserved. // Third party copyrights are property of their respective owners. // // @Authors @@ -63,7 +63,7 @@ /////////////////////////////////////////////////////////////// /////////////////common/////////////////////////////////////// ///////////////////////////////////////////////////////////// -inline T saturate_cast(float v){ +T saturate_cast(float v){ #ifdef T_SHORT return convert_short_sat_rte(v); #else @@ -71,7 +71,7 @@ inline T saturate_cast(float v){ #endif } -inline T4 saturate_cast4(float4 v){ +T4 saturate_cast4(float4 v){ #ifdef T_SHORT return convert_short4_sat_rte(v); #else @@ -92,12 +92,12 @@ typedef struct ////////////////////////// comp data ////////////////////////// /////////////////////////////////////////////////////////////// -inline float pix_diff_1(const uchar4 l, __global const uchar *rs) +float pix_diff_1(const uchar4 l, __global const uchar *rs) { return abs((int)(l.x) - *rs); } -inline float pix_diff_4(const uchar4 l, __global const uchar *rs) +float pix_diff_4(const uchar4 l, __global const uchar *rs) { uchar4 r; r = *((__global uchar4 *)rs); @@ -115,7 +115,7 @@ inline float pix_diff_4(const uchar4 l, __global const uchar *rs) return val; } -inline float pix_diff_3(const uchar4 l, __global const uchar *rs) +float pix_diff_3(const uchar4 l, __global const uchar *rs) { return pix_diff_4(l, rs); } @@ -233,7 +233,7 @@ __kernel void level_up_message(__global T *src, int src_rows, int src_step, /////////////////////////////////////////////////////////////// //////////////////// calc all iterations ///////////////////// /////////////////////////////////////////////////////////////// -inline void message(__global T *us_, __global T *ds_, __global T *ls_, __global T *rs_, +void message(__global T *us_, __global T *ds_, __global T *ls_, __global T *rs_, const __global T *dt, int u_step, int msg_disp_step, int data_disp_step, float4 cmax_disc_term, float4 cdisc_single_jump) diff --git a/modules/ocl/src/opencl/stereocsbp.cl b/modules/ocl/src/opencl/stereocsbp.cl index 23fc814817..865854ee74 100644 --- a/modules/ocl/src/opencl/stereocsbp.cl +++ b/modules/ocl/src/opencl/stereocsbp.cl @@ -12,7 +12,7 @@ // // Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. // Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved. -// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Copyright (C) 2010,2014, Advanced Micro Devices, Inc., all rights reserved. // Third party copyrights are property of their respective owners. // // @Authors @@ -242,7 +242,7 @@ __kernel void get_first_k_initial_local_1(__global float *data_cost_selected_, _ /////////////////////// init data cost //////////////////////// /////////////////////////////////////////////////////////////// -inline float compute_3(__global uchar* left, __global uchar* right, +float compute_3(__global uchar* left, __global uchar* right, float cdata_weight, float cmax_data_term) { float tb = 0.114f * abs((int)left[0] - right[0]); @@ -252,13 +252,13 @@ inline float compute_3(__global uchar* left, __global uchar* right, return fmin(cdata_weight * (tr + tg + tb), cdata_weight * cmax_data_term); } -inline float compute_1(__global uchar* left, __global uchar* right, +float compute_1(__global uchar* left, __global uchar* right, float cdata_weight, float cmax_data_term) { return fmin(cdata_weight * abs((int)*left - (int)*right), cdata_weight * cmax_data_term); } -inline short round_short(float v) +short round_short(float v) { return convert_short_sat_rte(v); } @@ -1000,7 +1000,7 @@ __kernel void compute_data_cost_reduce_1(__global const float *selected_disp_pyr //////////////////////// init message ///////////////////////// /////////////////////////////////////////////////////////////// -inline void get_first_k_element_increase_0(__global short* u_new, __global short *d_new, __global short *l_new, +void get_first_k_element_increase_0(__global short* u_new, __global short *d_new, __global short *l_new, __global short *r_new, __global const short *u_cur, __global const short *d_cur, __global const short *l_cur, __global const short *r_cur, __global short *data_cost_selected, __global short *disparity_selected_new, @@ -1165,7 +1165,7 @@ __kernel void init_message_1(__global float *u_new_, __global float *d_new_, __g //////////////////// calc all iterations ///////////////////// /////////////////////////////////////////////////////////////// -inline void message_per_pixel_0(__global const short *data, __global short *msg_dst, __global const short *msg1, +void message_per_pixel_0(__global const short *data, __global short *msg_dst, __global const short *msg1, __global const short *msg2, __global const short *msg3, __global const short *dst_disp, __global const short *src_disp, int nr_plane, __global short *temp, @@ -1202,7 +1202,7 @@ inline void message_per_pixel_0(__global const short *data, __global short *msg_ msg_dst[d * cdisp_step1] = convert_short_sat_rte(temp[d * cdisp_step1] - sum); } -inline void message_per_pixel_1(__global const float *data, __global float *msg_dst, __global const float *msg1, +void message_per_pixel_1(__global const float *data, __global float *msg_dst, __global const float *msg1, __global const float *msg2, __global const float *msg3, __global const float *dst_disp, __global const float *src_disp, int nr_plane, __global float *temp, diff --git a/modules/ocl/src/opencl/tvl1flow.cl b/modules/ocl/src/opencl/tvl1flow.cl index b488e89696..9d946a6233 100644 --- a/modules/ocl/src/opencl/tvl1flow.cl +++ b/modules/ocl/src/opencl/tvl1flow.cl @@ -11,7 +11,7 @@ // For Open Source Computer Vision Library // // Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. -// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Copyright (C) 2010,2014, Advanced Micro Devices, Inc., all rights reserved. // Third party copyrights are property of their respective owners. // // @Authors @@ -62,7 +62,7 @@ __kernel void centeredGradientKernel(__global const float* src, int src_col, int } -inline float bicubicCoeff(float x_) +float bicubicCoeff(float x_) { float x = fabs(x_); @@ -156,7 +156,7 @@ __kernel void warpBackwardKernel(__global const float* I0, int I0_step, int I0_c } -inline float readImage(__global float *image, int x, int y, int rows, int cols, int elemCntPerRow) +float readImage(__global float *image, int x, int y, int rows, int cols, int elemCntPerRow) { int i0 = clamp(x, 0, cols - 1); int j0 = clamp(y, 0, rows - 1); @@ -284,7 +284,7 @@ __kernel void estimateDualVariablesKernel(__global const float* u1, int u1_col, } -inline float divergence(__global const float* v1, __global const float* v2, int y, int x, int v1_step, int v2_step) +float divergence(__global const float* v1, __global const float* v2, int y, int x, int v1_step, int v2_step) { if (x > 0 && y > 0)