Skip to content
Projects
Groups
Snippets
Help
Loading...
Sign in / Register
Toggle navigation
O
opencv
Project
Project
Details
Activity
Cycle Analytics
Repository
Repository
Files
Commits
Branches
Tags
Contributors
Graph
Compare
Charts
Issues
0
Issues
0
List
Board
Labels
Milestones
Merge Requests
0
Merge Requests
0
CI / CD
CI / CD
Pipelines
Jobs
Schedules
Charts
Packages
Packages
Wiki
Wiki
Snippets
Snippets
Members
Members
Collapse sidebar
Close sidebar
Activity
Graph
Charts
Create a new issue
Jobs
Commits
Issue Boards
Open sidebar
submodule
opencv
Commits
ecb1f3c4
Commit
ecb1f3c4
authored
May 28, 2013
by
Vadim Pisarevsky
Committed by
OpenCV Buildbot
May 28, 2013
Browse files
Options
Browse Files
Download
Plain Diff
Merge pull request #923 from pengx17:2.4_macfix
parents
324cafdd
6fae02c0
Hide whitespace changes
Inline
Side-by-side
Showing
7 changed files
with
38 additions
and
20 deletions
+38
-20
filtering.cpp
modules/ocl/src/filtering.cpp
+1
-2
mcwutil.cpp
modules/ocl/src/mcwutil.cpp
+26
-8
arithm_add.cl
modules/ocl/src/opencl/arithm_add.cl
+5
-5
arithm_add_scalar_mask.cl
modules/ocl/src/opencl/arithm_add_scalar_mask.cl
+3
-3
filtering_morph.cl
modules/ocl/src/opencl/filtering_morph.cl
+1
-1
imgproc_threshold.cl
modules/ocl/src/opencl/imgproc_threshold.cl
+1
-1
precomp.hpp
modules/ocl/src/precomp.hpp
+1
-0
No files found.
modules/ocl/src/filtering.cpp
View file @
ecb1f3c4
...
@@ -356,8 +356,7 @@ static void GPUDilate(const oclMat &src, oclMat &dst, oclMat &mat_kernel,
...
@@ -356,8 +356,7 @@ static void GPUDilate(const oclMat &src, oclMat &dst, oclMat &mat_kernel,
char
compile_option
[
128
];
char
compile_option
[
128
];
sprintf
(
compile_option
,
"-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D DILATE %s %s"
,
sprintf
(
compile_option
,
"-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D DILATE %s %s"
,
anchor
.
x
,
anchor
.
y
,
(
int
)
localThreads
[
0
],
(
int
)
localThreads
[
1
],
anchor
.
x
,
anchor
.
y
,
(
int
)
localThreads
[
0
],
(
int
)
localThreads
[
1
],
rectKernel
?
"-D RECTKERNEL"
:
""
,
s
,
rectKernel
?
"-D RECTKERNEL"
:
""
);
s
);
vector
<
pair
<
size_t
,
const
void
*>
>
args
;
vector
<
pair
<
size_t
,
const
void
*>
>
args
;
args
.
push_back
(
make_pair
(
sizeof
(
cl_mem
),
(
void
*
)
&
src
.
data
));
args
.
push_back
(
make_pair
(
sizeof
(
cl_mem
),
(
void
*
)
&
src
.
data
));
args
.
push_back
(
make_pair
(
sizeof
(
cl_mem
),
(
void
*
)
&
dst
.
data
));
args
.
push_back
(
make_pair
(
sizeof
(
cl_mem
),
(
void
*
)
&
dst
.
data
));
...
...
modules/ocl/src/mcwutil.cpp
View file @
ecb1f3c4
...
@@ -43,9 +43,28 @@
...
@@ -43,9 +43,28 @@
//
//
//M*/
//M*/
#define CL_USE_DEPRECATED_OPENCL_1_1_APIS
#include "precomp.hpp"
#include "precomp.hpp"
#ifdef __GNUC__
#if ((__GNUC__ * 100) + __GNUC_MINOR__) >= 402
#define GCC_DIAG_STR(s) #s
#define GCC_DIAG_JOINSTR(x,y) GCC_DIAG_STR(x ## y)
# define GCC_DIAG_DO_PRAGMA(x) _Pragma (#x)
# define GCC_DIAG_PRAGMA(x) GCC_DIAG_DO_PRAGMA(GCC diagnostic x)
# if ((__GNUC__ * 100) + __GNUC_MINOR__) >= 406
# define GCC_DIAG_OFF(x) GCC_DIAG_PRAGMA(push) \
GCC_DIAG_PRAGMA(ignored GCC_DIAG_JOINSTR(-W,x))
# define GCC_DIAG_ON(x) GCC_DIAG_PRAGMA(pop)
# else
# define GCC_DIAG_OFF(x) GCC_DIAG_PRAGMA(ignored GCC_DIAG_JOINSTR(-W,x))
# define GCC_DIAG_ON(x) GCC_DIAG_PRAGMA(warning GCC_DIAG_JOINSTR(-W,x))
# endif
#else
# define GCC_DIAG_OFF(x)
# define GCC_DIAG_ON(x)
#endif
#endif
/* __GNUC__ */
using
namespace
std
;
using
namespace
std
;
namespace
cv
namespace
cv
...
@@ -121,6 +140,9 @@ namespace cv
...
@@ -121,6 +140,9 @@ namespace cv
build_options
,
finish_mode
);
build_options
,
finish_mode
);
}
}
#ifdef __GNUC__
GCC_DIAG_OFF
(
deprecated
-
declarations
)
#endif
cl_mem
bindTexture
(
const
oclMat
&
mat
)
cl_mem
bindTexture
(
const
oclMat
&
mat
)
{
{
cl_mem
texture
;
cl_mem
texture
;
...
@@ -180,10 +202,6 @@ namespace cv
...
@@ -180,10 +202,6 @@ namespace cv
else
else
#endif
#endif
{
{
#ifdef __GNUC__
#pragma GCC diagnostic push
#pragma GCC diagnostic ignored "-Wdeprecated-declarations"
#endif
texture
=
clCreateImage2D
(
texture
=
clCreateImage2D
(
(
cl_context
)
mat
.
clCxt
->
oclContext
(),
(
cl_context
)
mat
.
clCxt
->
oclContext
(),
CL_MEM_READ_WRITE
,
CL_MEM_READ_WRITE
,
...
@@ -193,9 +211,6 @@ namespace cv
...
@@ -193,9 +211,6 @@ namespace cv
0
,
0
,
NULL
,
NULL
,
&
err
);
&
err
);
#ifdef __GNUC__
#pragma GCC diagnostic pop
#endif
}
}
size_t
origin
[]
=
{
0
,
0
,
0
};
size_t
origin
[]
=
{
0
,
0
,
0
};
size_t
region
[]
=
{
mat
.
cols
,
mat
.
rows
,
1
};
size_t
region
[]
=
{
mat
.
cols
,
mat
.
rows
,
1
};
...
@@ -225,6 +240,9 @@ namespace cv
...
@@ -225,6 +240,9 @@ namespace cv
openCLSafeCall
(
err
);
openCLSafeCall
(
err
);
return
texture
;
return
texture
;
}
}
#ifdef __GNUC__
GCC_DIAG_ON
(
deprecated
-
declarations
)
#endif
void
releaseTexture
(
cl_mem
&
texture
)
void
releaseTexture
(
cl_mem
&
texture
)
{
{
openCLFree
(
texture
);
openCLFree
(
texture
);
...
...
modules/ocl/src/opencl/arithm_add.cl
View file @
ecb1f3c4
...
@@ -127,7 +127,7 @@ __kernel void arithm_add_D2 (__global ushort *src1, int src1_step, int src1_offs
...
@@ -127,7 +127,7 @@ __kernel void arithm_add_D2 (__global ushort *src1, int src1_step, int src1_offs
#
ifdef
dst_align
#
ifdef
dst_align
#
undef
dst_align
#
undef
dst_align
#
endif
#
endif
#
define
dst_align
((
dst_offset
>>
1
)
&
3
)
#
define
dst_align
((
dst_offset
/
2
)
&
3
)
int
src1_index
=
mad24
(
y,
src1_step,
(
x
<<
1
)
+
src1_offset
-
(
dst_align
<<
1
))
;
int
src1_index
=
mad24
(
y,
src1_step,
(
x
<<
1
)
+
src1_offset
-
(
dst_align
<<
1
))
;
int
src2_index
=
mad24
(
y,
src2_step,
(
x
<<
1
)
+
src2_offset
-
(
dst_align
<<
1
))
;
int
src2_index
=
mad24
(
y,
src2_step,
(
x
<<
1
)
+
src2_offset
-
(
dst_align
<<
1
))
;
...
@@ -165,7 +165,7 @@ __kernel void arithm_add_D3 (__global short *src1, int src1_step, int src1_offse
...
@@ -165,7 +165,7 @@ __kernel void arithm_add_D3 (__global short *src1, int src1_step, int src1_offse
#
ifdef
dst_align
#
ifdef
dst_align
#
undef
dst_align
#
undef
dst_align
#
endif
#
endif
#
define
dst_align
((
dst_offset
>>
1
)
&
3
)
#
define
dst_align
((
dst_offset
/
2
)
&
3
)
int
src1_index
=
mad24
(
y,
src1_step,
(
x
<<
1
)
+
src1_offset
-
(
dst_align
<<
1
))
;
int
src1_index
=
mad24
(
y,
src1_step,
(
x
<<
1
)
+
src1_offset
-
(
dst_align
<<
1
))
;
int
src2_index
=
mad24
(
y,
src2_step,
(
x
<<
1
)
+
src2_offset
-
(
dst_align
<<
1
))
;
int
src2_index
=
mad24
(
y,
src2_step,
(
x
<<
1
)
+
src2_offset
-
(
dst_align
<<
1
))
;
...
@@ -335,7 +335,7 @@ __kernel void arithm_add_with_mask_C1_D2 (__global ushort *src1, int src1_step,
...
@@ -335,7 +335,7 @@ __kernel void arithm_add_with_mask_C1_D2 (__global ushort *src1, int src1_step,
#
ifdef
dst_align
#
ifdef
dst_align
#
undef
dst_align
#
undef
dst_align
#
endif
#
endif
#
define
dst_align
((
dst_offset
>>
1
)
&
1
)
#
define
dst_align
((
dst_offset
/
2
)
&
1
)
int
src1_index
=
mad24
(
y,
src1_step,
(
x
<<
1
)
+
src1_offset
-
(
dst_align
<<
1
))
;
int
src1_index
=
mad24
(
y,
src1_step,
(
x
<<
1
)
+
src1_offset
-
(
dst_align
<<
1
))
;
int
src2_index
=
mad24
(
y,
src2_step,
(
x
<<
1
)
+
src2_offset
-
(
dst_align
<<
1
))
;
int
src2_index
=
mad24
(
y,
src2_step,
(
x
<<
1
)
+
src2_offset
-
(
dst_align
<<
1
))
;
int
mask_index
=
mad24
(
y,
mask_step,
x
+
mask_offset
-
dst_align
)
;
int
mask_index
=
mad24
(
y,
mask_step,
x
+
mask_offset
-
dst_align
)
;
...
@@ -375,7 +375,7 @@ __kernel void arithm_add_with_mask_C1_D3 (__global short *src1, int src1_step, i
...
@@ -375,7 +375,7 @@ __kernel void arithm_add_with_mask_C1_D3 (__global short *src1, int src1_step, i
#
ifdef
dst_align
#
ifdef
dst_align
#
undef
dst_align
#
undef
dst_align
#
endif
#
endif
#
define
dst_align
((
dst_offset
>>
1
)
&
1
)
#
define
dst_align
((
dst_offset
/
2
)
&
1
)
int
src1_index
=
mad24
(
y,
src1_step,
(
x
<<
1
)
+
src1_offset
-
(
dst_align
<<
1
))
;
int
src1_index
=
mad24
(
y,
src1_step,
(
x
<<
1
)
+
src1_offset
-
(
dst_align
<<
1
))
;
int
src2_index
=
mad24
(
y,
src2_step,
(
x
<<
1
)
+
src2_offset
-
(
dst_align
<<
1
))
;
int
src2_index
=
mad24
(
y,
src2_step,
(
x
<<
1
)
+
src2_offset
-
(
dst_align
<<
1
))
;
int
mask_index
=
mad24
(
y,
mask_step,
x
+
mask_offset
-
dst_align
)
;
int
mask_index
=
mad24
(
y,
mask_step,
x
+
mask_offset
-
dst_align
)
;
...
@@ -507,7 +507,7 @@ __kernel void arithm_add_with_mask_C2_D0 (__global uchar *src1, int src1_step, i
...
@@ -507,7 +507,7 @@ __kernel void arithm_add_with_mask_C2_D0 (__global uchar *src1, int src1_step, i
#
ifdef
dst_align
#
ifdef
dst_align
#
undef
dst_align
#
undef
dst_align
#
endif
#
endif
#
define
dst_align
((
dst_offset
>>
1
)
&
1
)
#
define
dst_align
((
dst_offset
/
2
)
&
1
)
int
src1_index
=
mad24
(
y,
src1_step,
(
x
<<
1
)
+
src1_offset
-
(
dst_align
<<
1
))
;
int
src1_index
=
mad24
(
y,
src1_step,
(
x
<<
1
)
+
src1_offset
-
(
dst_align
<<
1
))
;
int
src2_index
=
mad24
(
y,
src2_step,
(
x
<<
1
)
+
src2_offset
-
(
dst_align
<<
1
))
;
int
src2_index
=
mad24
(
y,
src2_step,
(
x
<<
1
)
+
src2_offset
-
(
dst_align
<<
1
))
;
int
mask_index
=
mad24
(
y,
mask_step,
x
+
mask_offset
-
dst_align
)
;
int
mask_index
=
mad24
(
y,
mask_step,
x
+
mask_offset
-
dst_align
)
;
...
...
modules/ocl/src/opencl/arithm_add_scalar_mask.cl
View file @
ecb1f3c4
...
@@ -126,7 +126,7 @@ __kernel void arithm_s_add_with_mask_C1_D2 (__global ushort *src1, int src1_st
...
@@ -126,7 +126,7 @@ __kernel void arithm_s_add_with_mask_C1_D2 (__global ushort *src1, int src1_st
#
ifdef
dst_align
#
ifdef
dst_align
#
undef
dst_align
#
undef
dst_align
#
endif
#
endif
#
define
dst_align
((
dst_offset
>>
1
)
&
1
)
#
define
dst_align
((
dst_offset
/
2
)
&
1
)
int
src1_index
=
mad24
(
y,
src1_step,
(
x
<<
1
)
+
src1_offset
-
(
dst_align
<<
1
))
;
int
src1_index
=
mad24
(
y,
src1_step,
(
x
<<
1
)
+
src1_offset
-
(
dst_align
<<
1
))
;
int
mask_index
=
mad24
(
y,
mask_step,
x
+
mask_offset
-
dst_align
)
;
int
mask_index
=
mad24
(
y,
mask_step,
x
+
mask_offset
-
dst_align
)
;
...
@@ -164,7 +164,7 @@ __kernel void arithm_s_add_with_mask_C1_D3 (__global short *src1, int src1_ste
...
@@ -164,7 +164,7 @@ __kernel void arithm_s_add_with_mask_C1_D3 (__global short *src1, int src1_ste
#
ifdef
dst_align
#
ifdef
dst_align
#
undef
dst_align
#
undef
dst_align
#
endif
#
endif
#
define
dst_align
((
dst_offset
>>
1
)
&
1
)
#
define
dst_align
((
dst_offset
/
2
)
&
1
)
int
src1_index
=
mad24
(
y,
src1_step,
(
x
<<
1
)
+
src1_offset
-
(
dst_align
<<
1
))
;
int
src1_index
=
mad24
(
y,
src1_step,
(
x
<<
1
)
+
src1_offset
-
(
dst_align
<<
1
))
;
int
mask_index
=
mad24
(
y,
mask_step,
x
+
mask_offset
-
dst_align
)
;
int
mask_index
=
mad24
(
y,
mask_step,
x
+
mask_offset
-
dst_align
)
;
...
@@ -288,7 +288,7 @@ __kernel void arithm_s_add_with_mask_C2_D0 (__global uchar *src1, int src1_ste
...
@@ -288,7 +288,7 @@ __kernel void arithm_s_add_with_mask_C2_D0 (__global uchar *src1, int src1_ste
#
ifdef
dst_align
#
ifdef
dst_align
#
undef
dst_align
#
undef
dst_align
#
endif
#
endif
#
define
dst_align
((
dst_offset
>>
1
)
&
1
)
#
define
dst_align
((
dst_offset
/
2
)
&
1
)
int
src1_index
=
mad24
(
y,
src1_step,
(
x
<<
1
)
+
src1_offset
-
(
dst_align
<<
1
))
;
int
src1_index
=
mad24
(
y,
src1_step,
(
x
<<
1
)
+
src1_offset
-
(
dst_align
<<
1
))
;
int
mask_index
=
mad24
(
y,
mask_step,
x
+
mask_offset
-
dst_align
)
;
int
mask_index
=
mad24
(
y,
mask_step,
x
+
mask_offset
-
dst_align
)
;
...
...
modules/ocl/src/opencl/filtering_morph.cl
View file @
ecb1f3c4
...
@@ -120,7 +120,7 @@ __kernel void morph_C1_D0(__global const uchar * restrict src,
...
@@ -120,7 +120,7 @@ __kernel void morph_C1_D0(__global const uchar * restrict src,
int
gidy
=
get_global_id
(
1
)
;
int
gidy
=
get_global_id
(
1
)
;
int
out_addr
=
mad24
(
gidy,dst_step_in_pixel,gidx+dst_offset_in_pixel
)
;
int
out_addr
=
mad24
(
gidy,dst_step_in_pixel,gidx+dst_offset_in_pixel
)
;
if
(
gidx+3<cols
&&
gidy<rows
&&
(
dst_offset_in_pixel&3
)
==0
)
if
(
gidx+3<cols
&&
gidy<rows
&&
(
(
dst_offset_in_pixel&3
)
==0
)
)
{
{
*
(
__global
uchar4*
)
&dst[out_addr]
=
res
;
*
(
__global
uchar4*
)
&dst[out_addr]
=
res
;
}
}
...
...
modules/ocl/src/opencl/imgproc_threshold.cl
View file @
ecb1f3c4
...
@@ -143,7 +143,7 @@ __kernel void threshold_C1_D5(__global const float * restrict src, __global floa
...
@@ -143,7 +143,7 @@ __kernel void threshold_C1_D5(__global const float * restrict src, __global floa
int4
dpos
=
(
int4
)(
dstart,
dstart+1,
dstart+2,
dstart+3
)
;
int4
dpos
=
(
int4
)(
dstart,
dstart+1,
dstart+2,
dstart+3
)
;
float4
dVal
=
*
(
__global
float4*
)(
dst+dst_offset+gy*dst_step+dstart
)
;
float4
dVal
=
*
(
__global
float4*
)(
dst+dst_offset+gy*dst_step+dstart
)
;
int4
con
=
dpos
>=
0
&&
dpos
<
dst_cols
;
int4
con
=
dpos
>=
0
&&
dpos
<
dst_cols
;
ddata
=
convert_float4
(
con
)
!=
0
?
ddata
:
dVal
;
ddata
=
convert_float4
(
con
)
!=
(
float4
)(
0
)
?
ddata
:
dVal
;
if
(
dstart
<
dst_cols
)
if
(
dstart
<
dst_cols
)
{
{
*
(
__global
float4*
)(
dst+dst_offset+gy*dst_step+dstart
)
=
ddata
;
*
(
__global
float4*
)(
dst+dst_offset+gy*dst_step+dstart
)
=
ddata
;
...
...
modules/ocl/src/precomp.hpp
View file @
ecb1f3c4
...
@@ -78,6 +78,7 @@
...
@@ -78,6 +78,7 @@
#if defined (HAVE_OPENCL)
#if defined (HAVE_OPENCL)
#define CL_USE_DEPRECATED_OPENCL_1_1_APIS
#include "opencv2/ocl/private/util.hpp"
#include "opencv2/ocl/private/util.hpp"
#include "safe_call.hpp"
#include "safe_call.hpp"
...
...
Write
Preview
Markdown
is supported
0%
Try again
or
attach a new file
Attach a file
Cancel
You are about to add
0
people
to the discussion. Proceed with caution.
Finish editing this message first!
Cancel
Please
register
or
sign in
to comment