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
1f51e6c0
Commit
1f51e6c0
authored
Oct 10, 2013
by
Ilya Lavrenov
Browse files
Options
Browse Files
Download
Email Patches
Plain Diff
fixed kernel compilation warnings on MacOSX
parent
ad1ba56f
Hide whitespace changes
Inline
Side-by-side
Showing
5 changed files
with
35 additions
and
12 deletions
+35
-12
cl_operations.cpp
modules/ocl/src/cl_operations.cpp
+25
-2
imgproc.cpp
modules/ocl/src/imgproc.cpp
+1
-1
arithm_absdiff_nonsaturate.cl
modules/ocl/src/opencl/arithm_absdiff_nonsaturate.cl
+2
-2
arithm_sum.cl
modules/ocl/src/opencl/arithm_sum.cl
+6
-6
pyrlk.cpp
modules/ocl/src/pyrlk.cpp
+1
-1
No files found.
modules/ocl/src/cl_operations.cpp
View file @
1f51e6c0
...
...
@@ -212,13 +212,35 @@ void openCLVerifyKernel(const Context *ctx, cl_kernel kernel, size_t *localThrea
static
double
total_execute_time
=
0
;
static
double
total_kernel_time
=
0
;
#endif
static
std
::
string
removeDuplicatedWhiteSpaces
(
const
char
*
buildOptions
)
{
if
(
buildOptions
==
NULL
)
return
""
;
size_t
length
=
strlen
(
buildOptions
),
didx
=
0
,
sidx
=
0
;
while
(
sidx
<
length
&&
buildOptions
[
sidx
]
==
0
)
++
sidx
;
std
::
string
opt
;
opt
.
resize
(
length
);
for
(
;
sidx
<
length
;
++
sidx
)
if
(
buildOptions
[
sidx
]
!=
' '
)
opt
[
didx
++
]
=
buildOptions
[
sidx
];
else
if
(
!
(
didx
>
0
&&
opt
[
didx
-
1
]
==
' '
)
)
opt
[
didx
++
]
=
buildOptions
[
sidx
];
return
opt
;
}
void
openCLExecuteKernel_
(
Context
*
ctx
,
const
cv
::
ocl
::
ProgramEntry
*
source
,
string
kernelName
,
size_t
globalThreads
[
3
],
size_t
localThreads
[
3
],
vector
<
pair
<
size_t
,
const
void
*>
>
&
args
,
int
channels
,
int
depth
,
const
char
*
build_options
)
{
//construct kernel name
//The rule is functionName_Cn_Dn, C represent Channels, D Represent DataType Depth, n represent an integer number
//for ex
maple split_C2_D2, represent the split kernel with channels =2 and dataType Depth = 2(Data type is char
)
//for ex
ample split_C2_D3, represent the split kernel with channels = 2 and dataType Depth = 3(Data type is short
)
stringstream
idxStr
;
if
(
channels
!=
-
1
)
idxStr
<<
"_C"
<<
channels
;
...
...
@@ -227,7 +249,8 @@ void openCLExecuteKernel_(Context *ctx, const cv::ocl::ProgramEntry* source, str
kernelName
+=
idxStr
.
str
();
cl_kernel
kernel
;
kernel
=
openCLGetKernelFromSource
(
ctx
,
source
,
kernelName
,
build_options
);
std
::
string
fixedOptions
=
removeDuplicatedWhiteSpaces
(
build_options
);
kernel
=
openCLGetKernelFromSource
(
ctx
,
source
,
kernelName
,
fixedOptions
.
c_str
());
if
(
localThreads
!=
NULL
)
{
...
...
modules/ocl/src/imgproc.cpp
View file @
1f51e6c0
...
...
@@ -1497,7 +1497,7 @@ namespace cv
openCLSafeCall
(
clReleaseKernel
(
kernel
));
static
char
opt
[
20
]
=
{
0
};
sprintf
(
opt
,
"
-D WAVE_SIZE=%d"
,
(
int
)
wave_size
);
sprintf
(
opt
,
"-D WAVE_SIZE=%d"
,
(
int
)
wave_size
);
openCLExecuteKernel
(
Context
::
getContext
(),
&
imgproc_clahe
,
kernelName
,
globalThreads
,
localThreads
,
args
,
-
1
,
-
1
,
opt
);
}
}
...
...
modules/ocl/src/opencl/arithm_absdiff_nonsaturate.cl
View file @
1f51e6c0
...
...
@@ -70,7 +70,7 @@ __kernel void arithm_absdiff_nonsaturate_binary(__global srcT *src1, int src1_st
dstT
t1
=
convertToDstT
(
src2[src2_index]
)
;
dstT
t2
=
t0
-
t1
;
dst[dst_index]
=
t2
>=
0
?
t2
:
-t2
;
dst[dst_index]
=
t2
>=
(
dstT
)(
0
)
?
t2
:
-t2
;
}
}
...
...
@@ -88,6 +88,6 @@ __kernel void arithm_absdiff_nonsaturate(__global srcT *src1, int src1_step, int
dstT
t0
=
convertToDstT
(
src1[src1_index]
)
;
dst[dst_index]
=
t0
>=
0
?
t0
:
-t0
;
dst[dst_index]
=
t0
>=
(
dstT
)(
0
)
?
t0
:
-t0
;
}
}
modules/ocl/src/opencl/arithm_sum.cl
View file @
1f51e6c0
...
...
@@ -51,14 +51,14 @@
#
endif
#
endif
#
if
defined
(
FUNC_SUM
)
#
if
FUNC_SUM
#
define
FUNC
(
a,
b
)
b
+=
a
;
#
endif
#
if
defined
(
FUNC_ABS_SUM
)
#
define
FUNC
(
a,
b
)
b
+=
a
>=
0
?
a
:
-a
;
#
endif
#
if
defined
(
FUNC_SQR_SUM
)
#
elif
FUNC_ABS_SUM
#
define
FUNC
(
a,
b
)
b
+=
a
>=
(
dstT
)(
0
)
?
a
:
-a
;
#
elif
FUNC_SQR_SUM
#
define
FUNC
(
a,
b
)
b
+=
a
*
a
;
#
else
#
error
No
sum
function
#
endif
/**************************************Array
buffer
SUM**************************************/
...
...
modules/ocl/src/pyrlk.cpp
View file @
1f51e6c0
...
...
@@ -134,7 +134,7 @@ static void lkSparse_run(oclMat &I, oclMat &J,
openCLSafeCall
(
clReleaseKernel
(
kernel
));
static
char
opt
[
32
]
=
{
0
};
sprintf
(
opt
,
"
-D WAVE_SIZE=%d"
,
wave_size
);
sprintf
(
opt
,
"-D WAVE_SIZE=%d"
,
wave_size
);
openCLExecuteKernel
(
clCxt
,
&
pyrlk
,
kernelName
,
globalThreads
,
localThreads
,
args
,
I
.
oclchannels
(),
I
.
depth
(),
opt
);
...
...
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