Skip to content

O
opencv
Commit 4f6f6e8c authored by Alexander Smorkalov's avatar Alexander Smorkalov
Browse files
Options

static function qualifier replaced on inline to enable kernel compilation with… 

static function qualifier replaced on inline to enable kernel compilation with OpenCL 1.1 embedded profile.
parent f7b5e654
Show whitespace changes
Inline Side-by-side
Showing with 22 additions and 22 deletions
modules/ocl/src/opencl/bgfg_mog.cl
View file @ 4f6f6e8c
......@@ -63,7 +63,7 @@ inline float sum(float val)
return val;
}
static float clamp1(float var, float learningRate, float diff, float minVar)
inline float clamp1(float var, float learningRate, float diff, float minVar)
{
return fmax(var + learningRate * (diff * diff - var), minVar);
}
......@@ -96,7 +96,7 @@ inline float sum(const float4 val)
return (val.x + val.y + val.z);
}
static void swap4(__global float4* ptr, int x, int y, int k, int rows, int ptr_step)
inline 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 @@ static void swap4(__global float4* ptr, int x, int y, int k, int rows, int ptr_s
}
static float4 clamp1(const float4 var, float learningRate, const float4 diff, float minVar)
inline 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;
static void swap(__global float* ptr, int x, int y, int k, int rows, int ptr_step)
inline 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/kmeans_kernel.cl
View file @ 4f6f6e8c
......@@ -44,7 +44,7 @@
//
//M*/
static float distance_(__global const float * center, __global const float * src, int feature_length)
inline float distance_(__global const float * center, __global const float * src, int feature_length)
{
float res = 0;
float4 v0, v1, v2;
......
modules/ocl/src/opencl/meanShift.cl
View file @ 4f6f6e8c
......@@ -46,7 +46,7 @@
//
//M*/
static short2 do_mean_shift(int x0, int y0, __global uchar4* out,int out_step,
inline 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 @ 4f6f6e8c
......@@ -208,7 +208,7 @@ __kernel void normalize_hists_36_kernel(__global float* block_hists,
//-------------------------------------------------------------
// Normalization of histograms via L2Hys_norm
//
static float reduce_smem(volatile __local float* smem, int size)
inline float reduce_smem(volatile __local float* smem, int size)
{
unsigned int tid = get_local_id(0);
float sum = smem[tid];
......
modules/ocl/src/opencl/pyrlk.cl
View file @ 4f6f6e8c
......@@ -52,7 +52,7 @@
#endif
#ifdef CPU
static void reduce3(float val1, float val2, float val3, __local float* smem1, __local float* smem2, __local float* smem3, int tid)
inline 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 @@ static void reduce3(float val1, float val2, float val3, __local float* smem1,
}
}
static void reduce2(float val1, float val2, volatile __local float* smem1, volatile __local float* smem2, int tid)
inline 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 @@ static void reduce2(float val1, float val2, volatile __local float* smem1, volat
}
}
static void reduce1(float val1, volatile __local float* smem1, int tid)
inline void reduce1(float val1, volatile __local float* smem1, int tid)
{
smem1[tid] = val1;
barrier(CLK_LOCAL_MEM_FENCE);
......@@ -103,7 +103,7 @@ static void reduce1(float val1, volatile __local float* smem1, int tid)
}
}
#else
static void reduce3(float val1, float val2, float val3,
inline 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 @@ static void reduce3(float val1, float val2, float val3,
barrier(CLK_LOCAL_MEM_FENCE);
}
static void reduce2(float val1, float val2, __local volatile float* smem1, __local volatile float* smem2, int tid)
inline 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 @@ static void reduce2(float val1, float val2, __local volatile float* smem1, __loc
barrier(CLK_LOCAL_MEM_FENCE);
}
static void reduce1(float val1, __local volatile float* smem1, int tid)
inline void reduce1(float val1, __local volatile float* smem1, int tid)
{
smem1[tid] = val1;
barrier(CLK_LOCAL_MEM_FENCE);
......@@ -225,7 +225,7 @@ static 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;
static void SetPatch(image2d_t I, float x, float y,
inline void SetPatch(image2d_t I, float x, float y,
float* Pch, float* Dx, float* Dy,
float* A11, float* A12, float* A22)
{
......@@ -262,7 +262,7 @@ inline void GetError(image2d_t J, const float x, const float y, const float* Pch
*errval += fabs(diff);
}
static void SetPatch4(image2d_t I, const float x, const float y,
inline 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 @@ static void SetPatch4(image2d_t I, const float x, const float y,
*A22 += sqIdx.x + sqIdx.y + sqIdx.z;
}
static void GetPatch4(image2d_t J, const float x, const float y,
inline 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 @@ static void GetPatch4(image2d_t J, const float x, const float y,
*b2 += xdiff.x + xdiff.y + xdiff.z;
}
static void GetError4(image2d_t J, const float x, const float y, const float4* Pch, float* errval)
inline 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/stereobp.cl
View file @ 4f6f6e8c
......@@ -97,7 +97,7 @@ inline float pix_diff_1(const uchar4 l, __global const uchar *rs)
return abs((int)(l.x) - *rs);
}
static float pix_diff_4(const uchar4 l, __global const uchar *rs)
inline float pix_diff_4(const uchar4 l, __global const uchar *rs)
{
uchar4 r;
r = *((__global uchar4 *)rs);
......@@ -233,7 +233,7 @@ __kernel void level_up_message(__global T *src, int src_rows, int src_step,
///////////////////////////////////////////////////////////////
//////////////////// calc all iterations /////////////////////
///////////////////////////////////////////////////////////////
static void message(__global T *us_, __global T *ds_, __global T *ls_, __global T *rs_,
inline 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/tvl1flow.cl
View file @ 4f6f6e8c
......@@ -62,7 +62,7 @@ __kernel void centeredGradientKernel(__global const float* src, int src_col, int
}
static float bicubicCoeff(float x_)
inline 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
}
static float readImage(__global float *image, int x, int y, int rows, int cols, int elemCntPerRow)
inline 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,
}
static float divergence(__global const float* v1, __global const float* v2, int y, int x, int v1_step, int v2_step)
inline 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