Chiebot-Mirror
/
opencv
mirror of https://github.com/opencv/opencv.git
8
0
Fork 0
Code Issues Projects Releases Wiki Activity

ocl: Don't use 'inline' attribute on functions

Browse Source 
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
pull/3244/head
Tom Stellard 11 years ago
parent 6c3cadbd73
commit 934394c5e2
22 changed files with 101 additions and 101 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Unified View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 22
      modules/ocl/src/opencl/bgfg_mog.cl
  2. 10
      modules/ocl/src/opencl/brute_force_match.cl
  3. 4
      modules/ocl/src/opencl/filtering_boxFilter.cl
  4. 4
      modules/ocl/src/opencl/filtering_filter2D.cl
  5. 10
      modules/ocl/src/opencl/imgproc_canny.cl
  6. 8
      modules/ocl/src/opencl/imgproc_clahe.cl
  7. 4
      modules/ocl/src/opencl/imgproc_warpAffine.cl
  8. 4
      modules/ocl/src/opencl/imgproc_warpPerspective.cl
  9. 4
      modules/ocl/src/opencl/interpolate_frames.cl
  10. 4
      modules/ocl/src/opencl/kernel_radix_sort_by_key.cl
  11. 6
      modules/ocl/src/opencl/kernel_stablesort_by_key.cl
  12. 4
      modules/ocl/src/opencl/kmeans_kernel.cl
  13. 6
      modules/ocl/src/opencl/match_template.cl
  14. 4
      modules/ocl/src/opencl/meanShift.cl
  15. 4
      modules/ocl/src/opencl/objdetect_hog.cl
  16. 14
      modules/ocl/src/opencl/optical_flow_farneback.cl
  17. 14
      modules/ocl/src/opencl/pyr_down.cl
  18. 26
      modules/ocl/src/opencl/pyrlk.cl
  19. 14
      modules/ocl/src/opencl/stereobm.cl
  20. 14
      modules/ocl/src/opencl/stereobp.cl
  21. 14
      modules/ocl/src/opencl/stereocsbp.cl
  22. 8
      modules/ocl/src/opencl/tvl1flow.cl

22
modules/ocl/src/opencl/bgfg_mog.cl
Unescape View File

@ -11,7 +11,7 @@
// For Open Source Computer Vision Library // For Open Source Computer Vision Library
// //
// Copyright (C) 2010-2013, Multicoreware, Inc., all rights reserved. // 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. // Third party copyrights are property of their respective owners.
// //
// @Authors // @Authors
@ -48,22 +48,22 @@
#define T_MEAN_VAR float #define T_MEAN_VAR float
#define CONVERT_TYPE convert_uchar_sat #define CONVERT_TYPE convert_uchar_sat
#define F_ZERO (0.0f) #define F_ZERO (0.0f)
inline float cvt(uchar val) float cvt(uchar val)
{ {
return val; return val;
} }
inline float sqr(float val) float sqr(float val)
{ {
return val * val; return val * val;
} }
inline float sum(float val) float sum(float val)
{ {
return 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); 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 CONVERT_TYPE convert_uchar4_sat
#define F_ZERO (0.0f, 0.0f, 0.0f, 0.0f) #define F_ZERO (0.0f, 0.0f, 0.0f, 0.0f)
inline float4 cvt(const uchar4 val) float4 cvt(const uchar4 val)
{ {
float4 result; float4 result;
result.x = val.x; result.x = val.x;
@ -86,17 +86,17 @@ inline float4 cvt(const uchar4 val)
return result; 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; 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); 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]; float4 val = ptr[(k * rows + y) * ptr_step + x];
ptr[(k * rows + y) * ptr_step + x] = ptr[((k + 1) * 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; float4 result;
result.x = fmax(var.x + learningRate * (diff.x * diff.x - var.x), minVar); result.x = fmax(var.x + learningRate * (diff.x * diff.x - var.x), minVar);
@ -128,7 +128,7 @@ typedef struct
uchar c_shadowVal; uchar c_shadowVal;
} con_srtuct_t; } 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]; float val = ptr[(k * rows + y) * ptr_step + x];
ptr[(k * rows + y) * ptr_step + x] = ptr[((k + 1) * rows + y) * ptr_step + x]; ptr[(k * rows + y) * ptr_step + x] = ptr[((k + 1) * rows + y) * ptr_step + x];

10
modules/ocl/src/opencl/brute_force_match.cl
Unescape View File

@ -11,7 +11,7 @@
// For Open Source Computer Vision Library // For Open Source Computer Vision Library
// //
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. // 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. // Third party copyrights are property of their respective owners.
// //
// @Authors // @Authors
@ -82,7 +82,7 @@ typedef float result_type;
#define DIST_RES(x) sqrt(x) #define DIST_RES(x) sqrt(x)
#elif (DIST_TYPE == 2) // Hamming #elif (DIST_TYPE == 2) // Hamming
//http://graphics.stanford.edu/~seander/bithacks.html#CountBitsSetParallel //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 - ((v >> 1) & 0x55555555); // reuse input as temporary
v = (v & 0x33333333) + ((v >> 2) & 0x33333333); // temp v = (v & 0x33333333) + ((v >> 2) & 0x33333333); // temp
@ -94,7 +94,7 @@ typedef int result_type;
#define DIST_RES(x) (x) #define DIST_RES(x) (x)
#endif #endif
inline result_type reduce_block( result_type reduce_block(
__local value_type *s_query, __local value_type *s_query,
__local value_type *s_train, __local value_type *s_train,
int lidx, int lidx,
@ -112,7 +112,7 @@ inline result_type reduce_block(
return DIST_RES(result); return DIST_RES(result);
} }
inline result_type reduce_block_match( result_type reduce_block_match(
__local value_type *s_query, __local value_type *s_query,
__local value_type *s_train, __local value_type *s_train,
int lidx, int lidx,
@ -130,7 +130,7 @@ inline result_type reduce_block_match(
return (result); return (result);
} }
inline result_type reduce_multi_block( result_type reduce_multi_block(
__local value_type *s_query, __local value_type *s_query,
__local value_type *s_train, __local value_type *s_train,
int block_index, int block_index,

4
modules/ocl/src/opencl/filtering_boxFilter.cl
Unescape View File

@ -10,7 +10,7 @@
// License Agreement // License Agreement
// For Open Source Computer Vision Library // 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. // Third party copyrights are property of their respective owners.
// //
// Redistribution and use in source and binary forms, with or without modification, // Redistribution and use in source and binary forms, with or without modification,
@ -225,7 +225,7 @@ struct RectCoords
#endif #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 #ifdef BORDER_CONSTANT
, SCALAR_TYPE borderValue , SCALAR_TYPE borderValue
#endif #endif

4
modules/ocl/src/opencl/filtering_filter2D.cl
Unescape View File

@ -10,7 +10,7 @@
// License Agreement // License Agreement
// For Open Source Computer Vision Library // 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. // Third party copyrights are property of their respective owners.
// //
// Redistribution and use in source and binary forms, with or without modification, // Redistribution and use in source and binary forms, with or without modification,
@ -222,7 +222,7 @@ struct RectCoords
#endif #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 #ifdef BORDER_CONSTANT
, SCALAR_TYPE borderValue , SCALAR_TYPE borderValue
#endif #endif

10
modules/ocl/src/opencl/imgproc_canny.cl
Unescape View File

@ -11,7 +11,7 @@
// For Open Source Computer Vision Library // For Open Source Computer Vision Library
// //
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. // 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. // Third party copyrights are property of their respective owners.
// //
// @Authors // @Authors
@ -44,12 +44,12 @@
//M*/ //M*/
#ifdef L2GRAD #ifdef L2GRAD
inline float calc(int x, int y) float calc(int x, int y)
{ {
return sqrt((float)(x * x + y * y)); return sqrt((float)(x * x + y * y));
} }
#else #else
inline float calc(int x, int y) float calc(int x, int y)
{ {
return (float)abs(x) + abs(y); return (float)abs(x) + abs(y);
} }
@ -381,8 +381,8 @@ struct PtrStepSz {
int step; int step;
int rows, cols; 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))); } 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; } 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 // do Hysteresis for pixel whose edge type is 1

8
modules/ocl/src/opencl/imgproc_clahe.cl
Unescape View File

@ -11,7 +11,7 @@
// For Open Source Computer Vision Library // For Open Source Computer Vision Library
// //
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. // 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. // Third party copyrights are property of their respective owners.
// //
// @Authors // @Authors
@ -47,7 +47,7 @@
#define WAVE_SIZE 1 #define WAVE_SIZE 1
#endif #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; smem[tid] = val;
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
@ -61,7 +61,7 @@ inline int calc_lut(__local int* smem, int val, int tid)
} }
#ifdef CPU #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; smem[tid] = val;
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
@ -101,7 +101,7 @@ inline void reduce(volatile __local int* smem, int val, int tid)
#else #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; smem[tid] = val;
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);

4
modules/ocl/src/opencl/imgproc_warpAffine.cl
Unescape View File

@ -11,7 +11,7 @@
// For Open Source Computer Vision Library // For Open Source Computer Vision Library
// //
// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, 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. // Third party copyrights are property of their respective owners.
// //
// @Authors // @Authors
@ -70,7 +70,7 @@ typedef float4 F4;
#define INTER_REMAP_COEF_BITS 15 #define INTER_REMAP_COEF_BITS 15
#define INTER_REMAP_COEF_SCALE (1 << INTER_REMAP_COEF_BITS) #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; const float A = -0.75f;

4
modules/ocl/src/opencl/imgproc_warpPerspective.cl
Unescape View File

@ -11,7 +11,7 @@
// For Open Source Computer Vision Library // For Open Source Computer Vision Library
// //
// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, 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. // Third party copyrights are property of their respective owners.
// //
// @Authors // @Authors
@ -71,7 +71,7 @@ typedef float4 F4;
#define INTER_REMAP_COEF_BITS 15 #define INTER_REMAP_COEF_BITS 15
#define INTER_REMAP_COEF_SCALE (1 << INTER_REMAP_COEF_BITS) #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; const float A = -0.75f;

4
modules/ocl/src/opencl/interpolate_frames.cl
Unescape View File

@ -11,7 +11,7 @@
// For Open Source Computer Vision Library // For Open Source Computer Vision Library
// //
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. // 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. // Third party copyrights are property of their respective owners.
// //
// @Authors // @Authors
@ -50,7 +50,7 @@
__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR; __constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR;
// atomic add for 32bit floating point // 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 { union {
unsigned int intVal; unsigned int intVal;
float floatVal; float floatVal;

4
modules/ocl/src/opencl/kernel_radix_sort_by_key.cl
Unescape View File

@ -11,7 +11,7 @@
// For Open Source Computer Vision Library // For Open Source Computer Vision Library
// //
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. // 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. // Third party copyrights are property of their respective owners.
// //
// @Authors // @Authors
@ -64,7 +64,7 @@
// from Thrust::b40c, link: // from Thrust::b40c, link:
// https://github.com/thrust/thrust/blob/master/thrust/system/cuda/detail/detail/b40c/radixsort_key_conversion.h // 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 #ifdef K_FLT
unsigned int mask = (converted_key & 0x80000000) ? 0xffffffff : 0x80000000; unsigned int mask = (converted_key & 0x80000000) ? 0xffffffff : 0x80000000;

6
modules/ocl/src/opencl/kernel_stablesort_by_key.cl
Unescape View File

@ -11,7 +11,7 @@
// For Open Source Computer Vision Library // For Open Source Computer Vision Library
// //
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. // 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. // Third party copyrights are property of their respective owners.
// //
// @Authors // @Authors
@ -65,7 +65,7 @@
// by a base pointer and left and right index for a particular candidate value. The comparison operator is // 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 // 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 // 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 // The values firstIndex and lastIndex get modified within the loop, narrowing down the potential sequence
uint firstIndex = left; 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 // 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 // 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 // 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 ); uint upperBound = lowerBoundBinary( data, left, right, searchVal );

4
modules/ocl/src/opencl/kmeans_kernel.cl
Unescape View File

@ -11,7 +11,7 @@
// For Open Source Computer Vision Library // For Open Source Computer Vision Library
// //
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. // 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. // Third party copyrights are property of their respective owners.
// //
// @Authors // @Authors
@ -44,7 +44,7 @@
// //
//M*/ //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; float res = 0;
float4 v0, v1, v2; float4 v0, v1, v2;

6
modules/ocl/src/opencl/match_template.cl
Unescape View File

@ -11,7 +11,7 @@
// For Open Source Computer Vision Library // For Open Source Computer Vision Library
// //
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. // 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. // Third party copyrights are property of their respective owners.
// //
// @Authors // @Authors
@ -66,7 +66,7 @@
#define SUMS_PTR(ox, oy) mad24(gidy + oy, img_sums_step, gidx + img_sums_offset + ox) #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 // normAcc* are accurate normalization routines which make GPU matchTemplate
// consistent with CPU one // consistent with CPU one
inline float normAcc(float num, float denum) float normAcc(float num, float denum)
{ {
if(fabs(num) < denum) if(fabs(num) < denum)
{ {
@ -79,7 +79,7 @@ inline float normAcc(float num, float denum)
return 0; return 0;
} }
inline float normAcc_SQDIFF(float num, float denum) float normAcc_SQDIFF(float num, float denum)
{ {
if(fabs(num) < denum) if(fabs(num) < denum)
{ {

4
modules/ocl/src/opencl/meanShift.cl
Unescape View File

@ -11,7 +11,7 @@
// For Open Source Computer Vision Library // For Open Source Computer Vision Library
// //
// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, 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.
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. // Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
// Third party copyrights are property of their respective owners. // Third party copyrights are property of their respective owners.
// //
@ -46,7 +46,7 @@
// //
//M*/ //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, __global uchar4* in, int in_step, int dst_off, int src_off,
int cols, int rows, int sp, int sr, int maxIter, float eps) int cols, int rows, int sp, int sr, int maxIter, float eps)
{ {

4
modules/ocl/src/opencl/objdetect_hog.cl
Unescape View File

@ -11,7 +11,7 @@
// For Open Source Computer Vision Library // For Open Source Computer Vision Library
// //
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. // 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. // Third party copyrights are property of their respective owners.
// //
// @Authors // @Authors
@ -208,7 +208,7 @@ __kernel void normalize_hists_36_kernel(__global float* block_hists,
//------------------------------------------------------------- //-------------------------------------------------------------
// Normalization of histograms via L2Hys_norm // 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); unsigned int tid = get_local_id(0);
float sum = smem[tid]; float sum = smem[tid];

14
modules/ocl/src/opencl/optical_flow_farneback.cl
Unescape View File

@ -11,7 +11,7 @@
// For Open Source Computer Vision Library // For Open Source Computer Vision Library
// //
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. // 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. // Third party copyrights are property of their respective owners.
// //
// @Authors // @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); 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); 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); 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); 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); 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); return idx_col_low(idx_col_high(x, last_col), last_col);
} }

14
modules/ocl/src/opencl/pyr_down.cl
Unescape View File

@ -11,7 +11,7 @@
// For Open Source Computer Vision Library // For Open Source Computer Vision Library
// //
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. // 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. // Third party copyrights are property of their respective owners.
// //
// @Authors // @Authors
@ -43,32 +43,32 @@
// //
//M*/ //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); 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); 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); 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); 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); 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); return idx_col_low(idx_col_high(x, last_col), last_col);
} }

26
modules/ocl/src/opencl/pyrlk.cl
Unescape View File

@ -11,7 +11,7 @@
// For Open Source Computer Vision Library // For Open Source Computer Vision Library
// //
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. // 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. // Third party copyrights are property of their respective owners.
// //
// @Authors // @Authors
@ -52,7 +52,7 @@
#endif #endif
#ifdef CPU #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; smem1[tid] = val1;
smem2[tid] = val2; 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; smem1[tid] = val1;
smem2[tid] = val2; 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; smem1[tid] = val1;
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
@ -103,7 +103,7 @@ inline void reduce1(float val1, volatile __local float* smem1, int tid)
} }
} }
#else #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) __local volatile float* smem1, __local volatile float* smem2, __local volatile float* smem3, int tid)
{ {
smem1[tid] = val1; smem1[tid] = val1;
@ -150,7 +150,7 @@ inline void reduce3(float val1, float val2, float val3,
barrier(CLK_LOCAL_MEM_FENCE); 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; smem1[tid] = val1;
smem2[tid] = val2; smem2[tid] = val2;
@ -189,7 +189,7 @@ inline void reduce2(float val1, float val2, __local volatile float* smem1, __loc
barrier(CLK_LOCAL_MEM_FENCE); 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; smem1[tid] = val1;
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
@ -225,7 +225,7 @@ inline void reduce1(float val1, __local volatile float* smem1, int tid)
// Image read mode // Image read mode
__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR; __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* Pch, float* Dx, float* Dy,
float* A11, float* A12, float* A22) float* A11, float* A12, float* A22)
{ {
@ -246,7 +246,7 @@ inline void SetPatch(image2d_t I, float x, float y,
*A22 += dIdy * dIdy; *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* Pch, float* Dx, float* Dy,
float* b1, float* b2) float* b1, float* b2)
{ {
@ -256,13 +256,13 @@ inline void GetPatch(image2d_t J, float x, float y,
*b2 += diff**Dy; *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; float diff = read_imagef(J, sampler, (float2)(x,y)).x-*Pch;
*errval += fabs(diff); *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, float4* Pch, float4* Dx, float4* Dy,
float* A11, float* A12, float* A22) 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; *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, const float4* Pch, const float4* Dx, const float4* Dy,
float* b1, float* b2) 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; *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; float4 diff = read_imagef(J, sampler, (float2)(x,y))-*Pch;
*errval += fabs(diff.x) + fabs(diff.y) + fabs(diff.z); *errval += fabs(diff.x) + fabs(diff.y) + fabs(diff.z);

14
modules/ocl/src/opencl/stereobm.cl
Unescape View File

@ -11,7 +11,7 @@
// For Open Source Computer Vision Library // For Open Source Computer Vision Library
// //
// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, 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. // Third party copyrights are property of their respective owners.
// //
// @Authors // @Authors
@ -56,7 +56,7 @@
#define radius 64 #define radius 64
#endif #endif
inline unsigned int CalcSSD(__local unsigned int *col_ssd) unsigned int CalcSSD(__local unsigned int *col_ssd)
{ {
unsigned int cache = col_ssd[0]; unsigned int cache = col_ssd[0];
@ -67,7 +67,7 @@ inline unsigned int CalcSSD(__local unsigned int *col_ssd)
return cache; return cache;
} }
inline uint2 MinSSD(__local unsigned int *col_ssd) uint2 MinSSD(__local unsigned int *col_ssd)
{ {
unsigned int ssd[N_DISPARITIES]; unsigned int ssd[N_DISPARITIES];
const int win_size = (radius << 1); const int win_size = (radius << 1);
@ -95,7 +95,7 @@ inline uint2 MinSSD(__local unsigned int *col_ssd)
return (uint2)(mssd, bestIdx); 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) __global unsigned char* imageR, int d, __local unsigned int *col_ssd)
{ {
uint8 imgR1 = convert_uint8(vload8(0, imageR + (idx1 - d - 7))); 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; 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, __global unsigned char* imageR, int d,
__local unsigned int *col_ssd) __local unsigned int *col_ssd)
{ {
@ -241,7 +241,7 @@ __kernel void prefilter_xsobel(__global unsigned char *input, __global unsigned
/////////////////////////////////// Textureness filtering //////////////////////////////////////// /////////////////////////////////// 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; float conv = 0;
int y1 = y==0? 0 : y-1; 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); 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]; unsigned int cache = cols[0];

14
modules/ocl/src/opencl/stereobp.cl
Unescape View File

@ -11,7 +11,7 @@
// For Open Source Computer Vision Library // For Open Source Computer Vision Library
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. // 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, 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. // Third party copyrights are property of their respective owners.
// //
// @Authors // @Authors
@ -63,7 +63,7 @@
/////////////////////////////////////////////////////////////// ///////////////////////////////////////////////////////////////
/////////////////common/////////////////////////////////////// /////////////////common///////////////////////////////////////
///////////////////////////////////////////////////////////// /////////////////////////////////////////////////////////////
inline T saturate_cast(float v){ T saturate_cast(float v){
#ifdef T_SHORT #ifdef T_SHORT
return convert_short_sat_rte(v); return convert_short_sat_rte(v);
#else #else
@ -71,7 +71,7 @@ inline T saturate_cast(float v){
#endif #endif
} }
inline T4 saturate_cast4(float4 v){ T4 saturate_cast4(float4 v){
#ifdef T_SHORT #ifdef T_SHORT
return convert_short4_sat_rte(v); return convert_short4_sat_rte(v);
#else #else
@ -92,12 +92,12 @@ typedef struct
////////////////////////// comp data ////////////////////////// ////////////////////////// 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); 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; uchar4 r;
r = *((__global uchar4 *)rs); r = *((__global uchar4 *)rs);
@ -115,7 +115,7 @@ inline float pix_diff_4(const uchar4 l, __global const uchar *rs)
return val; 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); 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 ///////////////////// //////////////////// 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, const __global T *dt,
int u_step, int msg_disp_step, int data_disp_step, int u_step, int msg_disp_step, int data_disp_step,
float4 cmax_disc_term, float4 cdisc_single_jump) float4 cmax_disc_term, float4 cdisc_single_jump)

14
modules/ocl/src/opencl/stereocsbp.cl
Unescape View File

@ -12,7 +12,7 @@
// //
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. // 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, 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. // Third party copyrights are property of their respective owners.
// //
// @Authors // @Authors
@ -242,7 +242,7 @@ __kernel void get_first_k_initial_local_1(__global float *data_cost_selected_, _
/////////////////////// init data cost //////////////////////// /////////////////////// 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 cdata_weight, float cmax_data_term)
{ {
float tb = 0.114f * abs((int)left[0] - right[0]); 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); 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) float cdata_weight, float cmax_data_term)
{ {
return fmin(cdata_weight * abs((int)*left - (int)*right), cdata_weight * 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); 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 ///////////////////////// //////////////////////// 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 short *r_new, __global const short *u_cur, __global const short *d_cur,
__global const short *l_cur, __global const short *r_cur, __global const short *l_cur, __global const short *r_cur,
__global short *data_cost_selected, __global short *disparity_selected_new, __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 ///////////////////// //////////////////// 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 *msg2, __global const short *msg3,
__global const short *dst_disp, __global const short *src_disp, __global const short *dst_disp, __global const short *src_disp,
int nr_plane, __global short *temp, 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); 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 *msg2, __global const float *msg3,
__global const float *dst_disp, __global const float *src_disp, __global const float *dst_disp, __global const float *src_disp,
int nr_plane, __global float *temp, int nr_plane, __global float *temp,

8
modules/ocl/src/opencl/tvl1flow.cl
Unescape View File

@ -11,7 +11,7 @@
// For Open Source Computer Vision Library // For Open Source Computer Vision Library
// //
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. // 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. // Third party copyrights are property of their respective owners.
// //
// @Authors // @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_); 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 i0 = clamp(x, 0, cols - 1);
int j0 = clamp(y, 0, rows - 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) if (x > 0 && y > 0)

Write Preview
Loading…
Cancel
Save