Skip to content

O
opencv
Commit 19b3aab2 authored by Vadim Pisarevsky's avatar Vadim Pisarevsky
Browse files
Options

Merge pull request #3244 from tstellarAMD:2.4-ocl-inline-fix-v2

parents df28bb5c 934394c5
Hide whitespace changes
Inline Side-by-side
Showing with 101 additions and 101 deletions
modules/ocl/src/opencl/bgfg_mog.cl
View file @ 19b3aab2
......@@ -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];
......
modules/ocl/src/opencl/brute_force_match.cl
View file @ 19b3aab2
......@@ -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,
......
modules/ocl/src/opencl/filtering_boxFilter.cl
View file @ 19b3aab2
......@@ -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
......
modules/ocl/src/opencl/filtering_filter2D.cl
View file @ 19b3aab2
......@@ -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
......
modules/ocl/src/opencl/imgproc_canny.cl
View file @ 19b3aab2
......@@ -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
......
modules/ocl/src/opencl/imgproc_clahe.cl
View file @ 19b3aab2
......@@ -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);
......
modules/ocl/src/opencl/imgproc_warpAffine.cl
View file @ 19b3aab2
......@@ -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;
......
modules/ocl/src/opencl/imgproc_warpPerspective.cl
View file @ 19b3aab2
......@@ -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;
......
modules/ocl/src/opencl/interpolate_frames.cl
View file @ 19b3aab2
......@@ -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;
......
modules/ocl/src/opencl/kernel_radix_sort_by_key.cl
View file @ 19b3aab2
......@@ -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;
......
modules/ocl/src/opencl/kernel_stablesort_by_key.cl
View file @ 19b3aab2
......@@ -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 );
......
modules/ocl/src/opencl/kmeans_kernel.cl
View file @ 19b3aab2
......@@ -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;
......
modules/ocl/src/opencl/match_template.cl
View file @ 19b3aab2
......@@ -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)
{
......
modules/ocl/src/opencl/meanShift.cl
View file @ 19b3aab2
......@@ -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)
{
......
modules/ocl/src/opencl/objdetect_hog.cl
View file @ 19b3aab2
......@@ -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];
......
modules/ocl/src/opencl/optical_flow_farneback.cl
View file @ 19b3aab2
......@@ -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);
}
......
modules/ocl/src/opencl/pyr_down.cl
View file @ 19b3aab2
......@@ -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);
}
......
modules/ocl/src/opencl/pyrlk.cl
View file @ 19b3aab2
......@@ -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);
......
modules/ocl/src/opencl/stereobm.cl
View file @ 19b3aab2
......@@ -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];
......
modules/ocl/src/opencl/stereobp.cl
View file @ 19b3aab2
......@@ -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)
......
modules/ocl/src/opencl/stereocsbp.cl
View file @ 19b3aab2
......@@ -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,
......
modules/ocl/src/opencl/tvl1flow.cl
View file @ 19b3aab2
......@@ -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)
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment