Skip to content
体验新版
项目
组织
正在加载...
登录
切换导航
打开侧边栏
a7660331a
tesseract
提交
50f92c8f
T
tesseract
项目概览
a7660331a
/
tesseract
与 Fork 源项目一致
从无法访问的项目Fork
通知
1
Star
0
Fork
0
代码
文件
提交
分支
Tags
贡献者
分支图
Diff
Issue
0
列表
看板
标记
里程碑
合并请求
0
Wiki
0
Wiki
分析
仓库
DevOps
项目成员
Pages
T
tesseract
项目概览
项目概览
详情
发布
仓库
仓库
文件
提交
分支
标签
贡献者
分支图
比较
Issue
0
Issue
0
列表
看板
标记
里程碑
合并请求
0
合并请求
0
Pages
分析
分析
仓库分析
DevOps
Wiki
0
Wiki
成员
成员
收起侧边栏
关闭侧边栏
动态
分支图
创建新Issue
提交
Issue看板
前往新版Gitcode,体验更适合开发者的 AI 搜索 >>
提交
50f92c8f
编写于
4月 26, 2017
作者:
Z
zdenop
提交者:
GitHub
4月 26, 2017
浏览文件
操作
浏览文件
下载
差异文件
Merge pull request #843 from stweil/opencl
opencl: Remove unused code
上级
7bb00d9a
e8ca29fb
变更
3
隐藏空白更改
内联
并排
Showing
3 changed file
with
0 addition
and
204 deletion
+0
-204
opencl/oclkernels.h
opencl/oclkernels.h
+0
-32
opencl/openclwrapper.cpp
opencl/openclwrapper.cpp
+0
-166
opencl/openclwrapper.h
opencl/openclwrapper.h
+0
-6
未找到文件。
opencl/oclkernels.h
浏览文件 @
50f92c8f
...
...
@@ -1175,38 +1175,6 @@ void kernel_ThresholdRectToPix_OneChan(
pix
[
w
]
=
word
;
}
}
)
KERNEL
(
\
n
#
define
RED_SHIFT
24
\
n
\
n
#
define
GREEN_SHIFT
16
\
n
\
n
#
define
BLUE_SHIFT
8
\
n
\
n
#
define
SET_DATA_BYTE
(
pdata
,
n
,
val
)
(
*
(
l_uint8
*
)((
l_uintptr_t
)((
l_uint8
*
)(
pdata
)
+
(
n
))
^
3
)
=
(
val
))
\
n
\
n
\
n__attribute__
((
reqd_work_group_size
(
256
,
1
,
1
)))
\
n
\
n__kernel
\
n
\
nvoid
kernel_RGBToGray
(
__global
const
unsigned
int
*
srcData
,
__global
unsigned
char
*
dstData
,
int
srcWPL
,
int
dstWPL
,
int
height
,
int
width
,
float
rwt
,
float
gwt
,
float
bwt
)
{
// pixel index
int
pixelIdx
=
get_global_id
(
0
);
if
(
pixelIdx
>=
height
*
width
)
return
;
unsigned
int
word
=
srcData
[
pixelIdx
];
int
output
=
(
rwt
*
((
word
>>
RED_SHIFT
)
&
0xff
)
+
gwt
*
((
word
>>
GREEN_SHIFT
)
&
0xff
)
+
bwt
*
((
word
>>
BLUE_SHIFT
)
&
0xff
)
+
0
.
5
f
);
// SET_DATA_BYTE
dstData
[
pixelIdx
]
=
output
;
}
)
;
// close char*
...
...
opencl/openclwrapper.cpp
浏览文件 @
50f92c8f
...
...
@@ -3611,170 +3611,4 @@ bool OpenclDevice::selectedDeviceIsOpenCL() {
return
(
device
.
type
==
DS_DEVICE_OPENCL_DEVICE
);
}
bool
OpenclDevice
::
selectedDeviceIsNativeCPU
()
{
ds_device
device
=
getDeviceSelection
();
return
(
device
.
type
==
DS_DEVICE_NATIVE_CPU
);
}
/*!
* pixConvertRGBToGray() from leptonica, converted to opencl kernel
*
* Input: pix (32 bpp RGB)
* rwt, gwt, bwt (non-negative; these should add to 1.0,
* or use 0.0 for default)
* Return: 8 bpp pix, or null on error
*
* Notes:
* (1) Use a weighted average of the RGB values.
*/
#define SET_DATA_BYTE(pdata, n, val) \
(*(l_uint8 *)((l_uintptr_t)((l_uint8 *)(pdata) + (n)) ^ 3) = (val))
Pix
*
OpenclDevice
::
pixConvertRGBToGrayOCL
(
Pix
*
srcPix
,
// 32-bit source
float
rwt
,
float
gwt
,
float
bwt
)
{
PERF_COUNT_START
(
"pixConvertRGBToGrayOCL"
)
Pix
*
dstPix
;
// 8-bit destination
if
(
rwt
<
0.0
||
gwt
<
0.0
||
bwt
<
0.0
)
return
nullptr
;
if
(
rwt
==
0.0
&&
gwt
==
0.0
&&
bwt
==
0.0
)
{
// magic numbers from leptonica
rwt
=
0.3
;
gwt
=
0.5
;
bwt
=
0.2
;
}
// normalize
float
sum
=
rwt
+
gwt
+
bwt
;
rwt
/=
sum
;
gwt
/=
sum
;
bwt
/=
sum
;
// source pix
int
w
,
h
;
pixGetDimensions
(
srcPix
,
&
w
,
&
h
,
nullptr
);
// printf("Image is %i x %i\n", w, h);
unsigned
int
*
srcData
=
pixGetData
(
srcPix
);
int
srcWPL
=
pixGetWpl
(
srcPix
);
int
srcSize
=
srcWPL
*
h
*
sizeof
(
unsigned
int
);
// destination pix
if
((
dstPix
=
pixCreate
(
w
,
h
,
8
))
==
nullptr
)
return
nullptr
;
pixCopyResolution
(
dstPix
,
srcPix
);
unsigned
int
*
dstData
=
pixGetData
(
dstPix
);
int
dstWPL
=
pixGetWpl
(
dstPix
);
int
dstWords
=
dstWPL
*
h
;
int
dstSize
=
dstWords
*
sizeof
(
unsigned
int
);
// printf("dstSize = %i\n", dstSize);
PERF_COUNT_SUB
(
"pix setup"
)
// opencl objects
cl_int
clStatus
;
KernelEnv
kEnv
;
SetKernelEnv
(
&
kEnv
);
// source buffer
cl_mem
srcBuffer
=
clCreateBuffer
(
kEnv
.
mpkContext
,
CL_MEM_READ_ONLY
|
CL_MEM_USE_HOST_PTR
,
srcSize
,
srcData
,
&
clStatus
);
CHECK_OPENCL
(
clStatus
,
"clCreateBuffer srcBuffer"
);
// destination buffer
cl_mem
dstBuffer
=
clCreateBuffer
(
kEnv
.
mpkContext
,
CL_MEM_WRITE_ONLY
|
CL_MEM_USE_HOST_PTR
,
dstSize
,
dstData
,
&
clStatus
);
CHECK_OPENCL
(
clStatus
,
"clCreateBuffer dstBuffer"
);
// setup work group size parameters
int
block_size
=
256
;
int
numWorkGroups
=
((
h
*
w
+
block_size
-
1
)
/
block_size
);
int
numThreads
=
block_size
*
numWorkGroups
;
size_t
local_work_size
[]
=
{
static_cast
<
size_t
>
(
block_size
)};
size_t
global_work_size
[]
=
{
static_cast
<
size_t
>
(
numThreads
)};
// printf("Enqueueing %i threads for %i output pixels\n", numThreads, w*h);
/* compile kernel */
kEnv
.
mpkKernel
=
clCreateKernel
(
kEnv
.
mpkProgram
,
"kernel_RGBToGray"
,
&
clStatus
);
CHECK_OPENCL
(
clStatus
,
"clCreateKernel kernel_RGBToGray"
);
/* set kernel arguments */
clStatus
=
clSetKernelArg
(
kEnv
.
mpkKernel
,
0
,
sizeof
(
cl_mem
),
&
srcBuffer
);
CHECK_OPENCL
(
clStatus
,
"clSetKernelArg srcBuffer"
);
clStatus
=
clSetKernelArg
(
kEnv
.
mpkKernel
,
1
,
sizeof
(
cl_mem
),
&
dstBuffer
);
CHECK_OPENCL
(
clStatus
,
"clSetKernelArg dstBuffer"
);
clStatus
=
clSetKernelArg
(
kEnv
.
mpkKernel
,
2
,
sizeof
(
int
),
&
srcWPL
);
CHECK_OPENCL
(
clStatus
,
"clSetKernelArg srcWPL"
);
clStatus
=
clSetKernelArg
(
kEnv
.
mpkKernel
,
3
,
sizeof
(
int
),
&
dstWPL
);
CHECK_OPENCL
(
clStatus
,
"clSetKernelArg dstWPL"
);
clStatus
=
clSetKernelArg
(
kEnv
.
mpkKernel
,
4
,
sizeof
(
int
),
&
h
);
CHECK_OPENCL
(
clStatus
,
"clSetKernelArg height"
);
clStatus
=
clSetKernelArg
(
kEnv
.
mpkKernel
,
5
,
sizeof
(
int
),
&
w
);
CHECK_OPENCL
(
clStatus
,
"clSetKernelArg width"
);
clStatus
=
clSetKernelArg
(
kEnv
.
mpkKernel
,
6
,
sizeof
(
float
),
&
rwt
);
CHECK_OPENCL
(
clStatus
,
"clSetKernelArg rwt"
);
clStatus
=
clSetKernelArg
(
kEnv
.
mpkKernel
,
7
,
sizeof
(
float
),
&
gwt
);
CHECK_OPENCL
(
clStatus
,
"clSetKernelArg gwt"
);
clStatus
=
clSetKernelArg
(
kEnv
.
mpkKernel
,
8
,
sizeof
(
float
),
&
bwt
);
CHECK_OPENCL
(
clStatus
,
"clSetKernelArg bwt"
);
/* launch kernel & wait */
PERF_COUNT_SUB
(
"before"
)
clStatus
=
clEnqueueNDRangeKernel
(
kEnv
.
mpkCmdQueue
,
kEnv
.
mpkKernel
,
1
,
nullptr
,
global_work_size
,
local_work_size
,
0
,
nullptr
,
nullptr
);
CHECK_OPENCL
(
clStatus
,
"clEnqueueNDRangeKernel kernel_RGBToGray"
);
clFinish
(
kEnv
.
mpkCmdQueue
);
PERF_COUNT_SUB
(
"kernel"
)
/* map results back from gpu */
void
*
ptr
=
clEnqueueMapBuffer
(
kEnv
.
mpkCmdQueue
,
dstBuffer
,
CL_TRUE
,
CL_MAP_READ
,
0
,
dstSize
,
0
,
nullptr
,
nullptr
,
&
clStatus
);
CHECK_OPENCL
(
clStatus
,
"clEnqueueMapBuffer dstBuffer"
);
clEnqueueUnmapMemObject
(
rEnv
.
mpkCmdQueue
,
dstBuffer
,
ptr
,
0
,
nullptr
,
nullptr
);
#if 0
// validate: compute on cpu
Pix *cpuPix = pixCreate(w, h, 8);
pixCopyResolution(cpuPix, srcPix);
unsigned int *cpuData = pixGetData(cpuPix);
int cpuWPL = pixGetWpl(cpuPix);
unsigned int *cpuLine, *srcLine;
int i, j;
for (i = 0, srcLine = srcData, cpuLine = cpuData; i < h; i++) {
for (j = 0; j < w; j++) {
unsigned int word = *(srcLine + j);
int val = (l_int32)(rwt * ((word >> L_RED_SHIFT) & 0xff) +
gwt * ((word >> L_GREEN_SHIFT) & 0xff) +
bwt * ((word >> L_BLUE_SHIFT) & 0xff) + 0.5);
SET_DATA_BYTE(cpuLine, j, val);
}
srcLine += srcWPL;
cpuLine += cpuWPL;
}
// validate: compare
printf("converted 32-bit -> 8-bit image\n");
for (int row = 0; row < h; row++) {
for (int col = 0; col < w; col++) {
int idx = row*w + col;
unsigned int srcVal = srcData[idx];
unsigned char cpuVal = ((unsigned char *)cpuData)[idx];
unsigned char oclVal = ((unsigned char *)dstData)[idx];
if (srcVal > 0) {
printf("%4i,%4i: %u, %u, %u\n", row, col, srcVal, cpuVal, oclVal);
}
}
//printf("\n");
}
#endif
// release opencl objects
clReleaseMemObject
(
srcBuffer
);
clReleaseMemObject
(
dstBuffer
);
PERF_COUNT_END
// success
return
dstPix
;
}
#endif
opencl/openclwrapper.h
浏览文件 @
50f92c8f
...
...
@@ -313,16 +313,10 @@ public:
int
rect_height
,
int
rect_width
,
int
rect_top
,
int
rect_left
);
static
Pix
*
pixConvertRGBToGrayOCL
(
Pix
*
pix
,
float
weightRed
=
0.3
,
float
weightGreen
=
0.5
,
float
weightBlue
=
0.2
);
static
ds_device
getDeviceSelection
();
static
ds_device
selectedDevice
;
static
bool
deviceIsSelected
;
static
bool
selectedDeviceIsOpenCL
();
static
bool
selectedDeviceIsNativeCPU
();
};
#endif // USE_OPENCL
...
...
编辑
预览
Markdown
is supported
0%
请重试
或
添加新附件
.
添加附件
取消
You are about to add
0
people
to the discussion. Proceed with caution.
先完成此消息的编辑!
取消
想要评论请
注册
或
登录