mirror of
https://github.com/tesseract-ocr/tesseract.git
synced 2025-01-18 06:30:14 +08:00
opencl: Remove more unused functions and related code
Signed-off-by: Stefan Weil <sw@weilnetz.de>
This commit is contained in:
Stefan Weil
parent
50f92c8f23
commit
582eaeba7d
@ -843,7 +843,6 @@ int OpenclDevice::InitOpenclRunEnv_DeviceSelection( int argc ) {
|
||||
//PERF_COUNT_START("InitOpenclRunEnv_DS")
|
||||
if (!isInited) {
|
||||
// after programs compiled, selects best device
|
||||
//printf("[DS] InitOpenclRunEnv_DS::Calling performDeviceSelection()\n");
|
||||
ds_device bestDevice_DS = getDeviceSelection( );
|
||||
//PERF_COUNT_SUB("called getDeviceSelection()")
|
||||
cl_device_id bestDevice = bestDevice_DS.oclDeviceID;
|
||||
@ -1065,22 +1064,6 @@ int OpenclDevice::GeneratBinFromKernelSource( cl_program program, const char * c
|
||||
return 1;
|
||||
}
|
||||
|
||||
void copyIntBuffer( KernelEnv rEnv, cl_mem xValues, const l_uint32 *_pValues, size_t nElements, cl_int *pStatus )
|
||||
{
|
||||
l_int32 *pValues = (l_int32 *)clEnqueueMapBuffer(
|
||||
rEnv.mpkCmdQueue, xValues, CL_TRUE, CL_MAP_WRITE, 0,
|
||||
nElements * sizeof(l_int32), 0, nullptr, nullptr, nullptr);
|
||||
clFinish(rEnv.mpkCmdQueue);
|
||||
if (_pValues != nullptr) {
|
||||
for (int i = 0; i < (int)nElements; i++) pValues[i] = (l_int32)_pValues[i];
|
||||
}
|
||||
|
||||
clEnqueueUnmapMemObject(rEnv.mpkCmdQueue, xValues, pValues, 0, nullptr,
|
||||
nullptr);
|
||||
//clFinish( rEnv.mpkCmdQueue );
|
||||
return;
|
||||
}
|
||||
|
||||
int OpenclDevice::CompileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
|
||||
{
|
||||
//PERF_COUNT_START("CompileKernelFile")
|
||||
@ -1317,575 +1300,6 @@ PERF_COUNT_END
|
||||
return pResult;
|
||||
}
|
||||
|
||||
PIX * OpenclDevice::pixReadTiffCl ( const char *filename, l_int32 n )
|
||||
{
|
||||
PERF_COUNT_START("pixReadTiffCL")
|
||||
FILE *fp;
|
||||
PIX *pix;
|
||||
|
||||
//printf("pixReadTiffCl file");
|
||||
PROCNAME("pixReadTiff");
|
||||
|
||||
if (!filename)
|
||||
return (PIX *)ERROR_PTR("filename not defined", procName, nullptr);
|
||||
|
||||
if ((fp = fopenReadStream(filename)) == nullptr)
|
||||
return (PIX *)ERROR_PTR("image file not found", procName, nullptr);
|
||||
if ((pix = pixReadStreamTiffCl(fp, n)) == nullptr) {
|
||||
fclose(fp);
|
||||
return (PIX *)ERROR_PTR("pix not read", procName, nullptr);
|
||||
}
|
||||
fclose(fp);
|
||||
PERF_COUNT_END
|
||||
return pix;
|
||||
}
|
||||
TIFF *
|
||||
OpenclDevice::fopenTiffCl(FILE *fp,
|
||||
const char *modestring)
|
||||
{
|
||||
l_int32 fd;
|
||||
|
||||
PROCNAME("fopenTiff");
|
||||
|
||||
if (!fp) return (TIFF *)ERROR_PTR("stream not opened", procName, nullptr);
|
||||
if (!modestring)
|
||||
return (TIFF *)ERROR_PTR("modestring not defined", procName, nullptr);
|
||||
|
||||
if ((fd = fileno(fp)) < 0)
|
||||
return (TIFF *)ERROR_PTR("invalid file descriptor", procName, nullptr);
|
||||
lseek(fd, 0, SEEK_SET);
|
||||
|
||||
return TIFFFdOpen(fd, "TIFFstream", modestring);
|
||||
}
|
||||
l_int32 OpenclDevice::getTiffStreamResolutionCl(TIFF *tif,
|
||||
l_int32 *pxres,
|
||||
l_int32 *pyres)
|
||||
{
|
||||
l_uint16 resunit;
|
||||
l_int32 foundxres, foundyres;
|
||||
l_float32 fxres, fyres;
|
||||
|
||||
PROCNAME("getTiffStreamResolution");
|
||||
|
||||
if (!tif)
|
||||
return ERROR_INT("tif not opened", procName, 1);
|
||||
if (!pxres || !pyres)
|
||||
return ERROR_INT("&xres and &yres not both defined", procName, 1);
|
||||
*pxres = *pyres = 0;
|
||||
|
||||
TIFFGetFieldDefaulted(tif, TIFFTAG_RESOLUTIONUNIT, &resunit);
|
||||
foundxres = TIFFGetField(tif, TIFFTAG_XRESOLUTION, &fxres);
|
||||
foundyres = TIFFGetField(tif, TIFFTAG_YRESOLUTION, &fyres);
|
||||
if (!foundxres && !foundyres) return 1;
|
||||
if (!foundxres && foundyres)
|
||||
fxres = fyres;
|
||||
else if (foundxres && !foundyres)
|
||||
fyres = fxres;
|
||||
|
||||
if (resunit == RESUNIT_CENTIMETER) { /* convert to ppi */
|
||||
*pxres = (l_int32)(2.54 * fxres + 0.5);
|
||||
*pyres = (l_int32)(2.54 * fyres + 0.5);
|
||||
}
|
||||
else {
|
||||
*pxres = (l_int32)fxres;
|
||||
*pyres = (l_int32)fyres;
|
||||
}
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
struct L_Memstream
|
||||
{
|
||||
l_uint8 *buffer; /* expands to hold data when written to; */
|
||||
/* fixed size when read from. */
|
||||
size_t bufsize; /* current size allocated when written to; */
|
||||
/* fixed size of input data when read from. */
|
||||
size_t offset; /* byte offset from beginning of buffer. */
|
||||
size_t hw; /* high-water mark; max bytes in buffer. */
|
||||
l_uint8 **poutdata; /* input param for writing; data goes here. */
|
||||
size_t *poutsize; /* input param for writing; data size goes here. */
|
||||
};
|
||||
typedef struct L_Memstream L_MEMSTREAM;
|
||||
|
||||
/* These are static functions for memory I/O */
|
||||
static L_MEMSTREAM *memstreamCreateForRead(l_uint8 *indata, size_t pinsize);
|
||||
static L_MEMSTREAM *memstreamCreateForWrite(l_uint8 **poutdata,
|
||||
size_t *poutsize);
|
||||
static tsize_t tiffReadCallback(thandle_t handle, tdata_t data, tsize_t length);
|
||||
static tsize_t tiffWriteCallback(thandle_t handle, tdata_t data,
|
||||
tsize_t length);
|
||||
static toff_t tiffSeekCallback(thandle_t handle, toff_t offset, l_int32 whence);
|
||||
static l_int32 tiffCloseCallback(thandle_t handle);
|
||||
static toff_t tiffSizeCallback(thandle_t handle);
|
||||
static l_int32 tiffMapCallback(thandle_t handle, tdata_t *data, toff_t *length);
|
||||
static void tiffUnmapCallback(thandle_t handle, tdata_t data, toff_t length);
|
||||
|
||||
|
||||
static L_MEMSTREAM *
|
||||
memstreamCreateForRead(l_uint8 *indata,
|
||||
size_t insize)
|
||||
{
|
||||
L_MEMSTREAM *mstream;
|
||||
|
||||
mstream = (L_MEMSTREAM *)CALLOC(1, sizeof(L_MEMSTREAM));
|
||||
mstream->buffer = indata; /* handle to input data array */
|
||||
mstream->bufsize = insize; /* amount of input data */
|
||||
mstream->hw = insize; /* high-water mark fixed at input data size */
|
||||
mstream->offset = 0; /* offset always starts at 0 */
|
||||
return mstream;
|
||||
}
|
||||
|
||||
|
||||
static L_MEMSTREAM *
|
||||
memstreamCreateForWrite(l_uint8 **poutdata,
|
||||
size_t *poutsize)
|
||||
{
|
||||
L_MEMSTREAM *mstream;
|
||||
|
||||
mstream = (L_MEMSTREAM *)CALLOC(1, sizeof(L_MEMSTREAM));
|
||||
mstream->buffer = (l_uint8 *)CALLOC(8 * 1024, 1);
|
||||
mstream->bufsize = 8 * 1024;
|
||||
mstream->poutdata = poutdata; /* used only at end of write */
|
||||
mstream->poutsize = poutsize; /* ditto */
|
||||
mstream->hw = mstream->offset = 0;
|
||||
return mstream;
|
||||
}
|
||||
|
||||
|
||||
static tsize_t
|
||||
tiffReadCallback(thandle_t handle,
|
||||
tdata_t data,
|
||||
tsize_t length)
|
||||
{
|
||||
L_MEMSTREAM *mstream;
|
||||
size_t amount;
|
||||
|
||||
mstream = (L_MEMSTREAM *)handle;
|
||||
amount = L_MIN((size_t)length, mstream->hw - mstream->offset);
|
||||
memcpy(data, mstream->buffer + mstream->offset, amount);
|
||||
mstream->offset += amount;
|
||||
return amount;
|
||||
}
|
||||
|
||||
|
||||
static tsize_t
|
||||
tiffWriteCallback(thandle_t handle,
|
||||
tdata_t data,
|
||||
tsize_t length)
|
||||
{
|
||||
L_MEMSTREAM *mstream;
|
||||
size_t newsize;
|
||||
|
||||
/* reallocNew() uses calloc to initialize the array.
|
||||
* If malloc is used instead, for some of the encoding methods,
|
||||
* not all the data in 'bufsize' bytes in the buffer will
|
||||
* have been initialized by the end of the compression. */
|
||||
mstream = (L_MEMSTREAM *)handle;
|
||||
if (mstream->offset + length > mstream->bufsize) {
|
||||
newsize = 2 * (mstream->offset + length);
|
||||
mstream->buffer = (l_uint8 *)reallocNew((void **)&mstream->buffer,
|
||||
mstream->offset, newsize);
|
||||
mstream->bufsize = newsize;
|
||||
}
|
||||
|
||||
memcpy(mstream->buffer + mstream->offset, data, length);
|
||||
mstream->offset += length;
|
||||
mstream->hw = L_MAX(mstream->offset, mstream->hw);
|
||||
return length;
|
||||
}
|
||||
|
||||
|
||||
static toff_t
|
||||
tiffSeekCallback(thandle_t handle,
|
||||
toff_t offset,
|
||||
l_int32 whence)
|
||||
{
|
||||
L_MEMSTREAM *mstream;
|
||||
|
||||
PROCNAME("tiffSeekCallback");
|
||||
mstream = (L_MEMSTREAM *)handle;
|
||||
switch (whence) {
|
||||
case SEEK_SET:
|
||||
/* fprintf(stderr, "seek_set: offset = %d\n", offset); */
|
||||
mstream->offset = offset;
|
||||
break;
|
||||
case SEEK_CUR:
|
||||
/* fprintf(stderr, "seek_cur: offset = %d\n", offset); */
|
||||
mstream->offset += offset;
|
||||
break;
|
||||
case SEEK_END:
|
||||
/* fprintf(stderr, "seek end: hw = %d, offset = %d\n",
|
||||
mstream->hw, offset); */
|
||||
mstream->offset = mstream->hw - offset; /* offset >= 0 */
|
||||
break;
|
||||
default:
|
||||
return (toff_t)ERROR_INT("bad whence value", procName,
|
||||
mstream->offset);
|
||||
}
|
||||
|
||||
return mstream->offset;
|
||||
}
|
||||
|
||||
|
||||
static l_int32
|
||||
tiffCloseCallback(thandle_t handle)
|
||||
{
|
||||
L_MEMSTREAM *mstream;
|
||||
|
||||
mstream = (L_MEMSTREAM *)handle;
|
||||
if (mstream->poutdata) { /* writing: save the output data */
|
||||
*mstream->poutdata = mstream->buffer;
|
||||
*mstream->poutsize = mstream->hw;
|
||||
}
|
||||
FREE(mstream); /* never free the buffer! */
|
||||
return 0;
|
||||
}
|
||||
|
||||
|
||||
static toff_t
|
||||
tiffSizeCallback(thandle_t handle)
|
||||
{
|
||||
L_MEMSTREAM *mstream;
|
||||
|
||||
mstream = (L_MEMSTREAM *)handle;
|
||||
return mstream->hw;
|
||||
}
|
||||
|
||||
|
||||
static l_int32
|
||||
tiffMapCallback(thandle_t handle,
|
||||
tdata_t *data,
|
||||
toff_t *length)
|
||||
{
|
||||
L_MEMSTREAM *mstream;
|
||||
|
||||
mstream = (L_MEMSTREAM *)handle;
|
||||
*data = mstream->buffer;
|
||||
*length = mstream->hw;
|
||||
return 0;
|
||||
}
|
||||
|
||||
|
||||
static void
|
||||
tiffUnmapCallback(thandle_t handle,
|
||||
tdata_t data,
|
||||
toff_t length)
|
||||
{
|
||||
return;
|
||||
}
|
||||
|
||||
|
||||
/*!
|
||||
* fopenTiffMemstream()
|
||||
*
|
||||
* Input: filename (for error output; can be "")
|
||||
* operation ("w" for write, "r" for read)
|
||||
* &data (<return> written data)
|
||||
* &datasize (<return> size of written data)
|
||||
* Return: tiff (data structure, opened for write to memory)
|
||||
*
|
||||
* Notes:
|
||||
* (1) This wraps up a number of callbacks for either:
|
||||
* * reading from tiff in memory buffer --> pix
|
||||
* * writing from pix --> tiff in memory buffer
|
||||
* (2) After use, the memstream is automatically destroyed when
|
||||
* TIFFClose() is called. TIFFCleanup() doesn't free the memstream.
|
||||
*/
|
||||
static TIFF *
|
||||
fopenTiffMemstream(const char *filename,
|
||||
const char *operation,
|
||||
l_uint8 **pdata,
|
||||
size_t *pdatasize)
|
||||
{
|
||||
L_MEMSTREAM *mstream;
|
||||
|
||||
PROCNAME("fopenTiffMemstream");
|
||||
|
||||
if (!filename)
|
||||
return (TIFF *)ERROR_PTR("filename not defined", procName, nullptr);
|
||||
if (!operation)
|
||||
return (TIFF *)ERROR_PTR("operation not defined", procName, nullptr);
|
||||
if (!pdata)
|
||||
return (TIFF *)ERROR_PTR("&data not defined", procName, nullptr);
|
||||
if (!pdatasize)
|
||||
return (TIFF *)ERROR_PTR("&datasize not defined", procName, nullptr);
|
||||
if (!strcmp(operation, "r") && !strcmp(operation, "w"))
|
||||
return (TIFF *)ERROR_PTR("operation not 'r' or 'w'}", procName,
|
||||
nullptr);
|
||||
|
||||
if (!strcmp(operation, "r"))
|
||||
mstream = memstreamCreateForRead(*pdata, *pdatasize);
|
||||
else
|
||||
mstream = memstreamCreateForWrite(pdata, pdatasize);
|
||||
|
||||
return TIFFClientOpen(filename, operation, mstream, tiffReadCallback,
|
||||
tiffWriteCallback, tiffSeekCallback,
|
||||
tiffCloseCallback, tiffSizeCallback,
|
||||
tiffMapCallback, tiffUnmapCallback);
|
||||
}
|
||||
|
||||
|
||||
|
||||
PIX *
|
||||
OpenclDevice::pixReadMemTiffCl(const l_uint8 *data,size_t size,l_int32 n)
|
||||
{
|
||||
l_int32 i, pagefound;
|
||||
PIX *pix;
|
||||
TIFF *tif;
|
||||
// L_MEMSTREAM *memStream;
|
||||
PROCNAME("pixReadMemTiffCl");
|
||||
|
||||
if (!data)
|
||||
return (PIX *)ERROR_PTR("data pointer is nullptr", procName, nullptr);
|
||||
|
||||
if ((tif = fopenTiffMemstream("", "r", (l_uint8 **)&data, &size)) ==
|
||||
nullptr)
|
||||
return (PIX *)ERROR_PTR("tif not opened", procName, nullptr);
|
||||
|
||||
pagefound = FALSE;
|
||||
pix = nullptr;
|
||||
for (i = 0; i < MAX_PAGES_IN_TIFF_FILE; i++) {
|
||||
if (i == n) {
|
||||
pagefound = TRUE;
|
||||
if ((pix = pixReadFromTiffStreamCl(tif)) == nullptr) {
|
||||
TIFFCleanup(tif);
|
||||
return (PIX *)ERROR_PTR("pix not read", procName, nullptr);
|
||||
}
|
||||
break;
|
||||
}
|
||||
if (TIFFReadDirectory(tif) == 0) break;
|
||||
}
|
||||
|
||||
if (pagefound == FALSE) {
|
||||
L_WARNING("tiff page %d not found\n", procName, i);
|
||||
TIFFCleanup(tif);
|
||||
return nullptr;
|
||||
}
|
||||
|
||||
TIFFCleanup(tif);
|
||||
return pix;
|
||||
}
|
||||
|
||||
PIX *
|
||||
OpenclDevice::pixReadStreamTiffCl(FILE *fp,
|
||||
l_int32 n)
|
||||
{
|
||||
l_int32 i, pagefound;
|
||||
PIX *pix;
|
||||
TIFF *tif;
|
||||
|
||||
PROCNAME("pixReadStreamTiff");
|
||||
|
||||
if (!fp) return (PIX *)ERROR_PTR("stream not defined", procName, nullptr);
|
||||
|
||||
if ((tif = fopenTiffCl(fp, "rb")) == nullptr)
|
||||
return (PIX *)ERROR_PTR("tif not opened", procName, nullptr);
|
||||
|
||||
pagefound = FALSE;
|
||||
pix = nullptr;
|
||||
for (i = 0; i < MAX_PAGES_IN_TIFF_FILE; i++) {
|
||||
if (i == n) {
|
||||
pagefound = TRUE;
|
||||
if ((pix = pixReadFromTiffStreamCl(tif)) == nullptr) {
|
||||
TIFFCleanup(tif);
|
||||
return (PIX *)ERROR_PTR("pix not read", procName, nullptr);
|
||||
}
|
||||
break;
|
||||
}
|
||||
if (TIFFReadDirectory(tif) == 0)
|
||||
break;
|
||||
}
|
||||
|
||||
if (pagefound == FALSE) {
|
||||
L_WARNING("tiff page %d not found", procName, n);
|
||||
TIFFCleanup(tif);
|
||||
return nullptr;
|
||||
}
|
||||
|
||||
TIFFCleanup(tif);
|
||||
return pix;
|
||||
}
|
||||
|
||||
static l_int32
|
||||
getTiffCompressedFormat(l_uint16 tiffcomp)
|
||||
{
|
||||
l_int32 comptype;
|
||||
|
||||
switch (tiffcomp)
|
||||
{
|
||||
case COMPRESSION_CCITTFAX4:
|
||||
comptype = IFF_TIFF_G4;
|
||||
break;
|
||||
case COMPRESSION_CCITTFAX3:
|
||||
comptype = IFF_TIFF_G3;
|
||||
break;
|
||||
case COMPRESSION_CCITTRLE:
|
||||
comptype = IFF_TIFF_RLE;
|
||||
break;
|
||||
case COMPRESSION_PACKBITS:
|
||||
comptype = IFF_TIFF_PACKBITS;
|
||||
break;
|
||||
case COMPRESSION_LZW:
|
||||
comptype = IFF_TIFF_LZW;
|
||||
break;
|
||||
case COMPRESSION_ADOBE_DEFLATE:
|
||||
comptype = IFF_TIFF_ZIP;
|
||||
break;
|
||||
default:
|
||||
comptype = IFF_TIFF;
|
||||
break;
|
||||
}
|
||||
return comptype;
|
||||
}
|
||||
|
||||
void compare(l_uint32 *cpu, l_uint32 *gpu,int size)
|
||||
{
|
||||
for(int i=0;i<size;i++)
|
||||
{
|
||||
if(cpu[i]!=gpu[i])
|
||||
{
|
||||
printf("\ndoesnot match\n");
|
||||
return;
|
||||
}
|
||||
}
|
||||
printf("\nit matches\n");
|
||||
}
|
||||
|
||||
//OpenCL implementation of pixReadFromTiffStream.
|
||||
//Similar to the CPU implentation of pixReadFromTiffStream
|
||||
PIX *
|
||||
OpenclDevice::pixReadFromTiffStreamCl(TIFF *tif)
|
||||
{
|
||||
l_uint8 *linebuf, *data;
|
||||
l_uint16 spp, bps, bpp, tiffbpl, photometry, tiffcomp, orientation;
|
||||
l_uint16 *redmap, *greenmap, *bluemap;
|
||||
l_int32 d, wpl, bpl, comptype, i, ncolors;
|
||||
l_int32 xres, yres;
|
||||
l_uint32 w, h;
|
||||
l_uint32 *line, *tiffdata;
|
||||
PIX *pix;
|
||||
PIXCMAP *cmap;
|
||||
|
||||
PROCNAME("pixReadFromTiffStream");
|
||||
|
||||
if (!tif) return (PIX *)ERROR_PTR("tif not defined", procName, nullptr);
|
||||
|
||||
TIFFGetFieldDefaulted(tif, TIFFTAG_BITSPERSAMPLE, &bps);
|
||||
TIFFGetFieldDefaulted(tif, TIFFTAG_SAMPLESPERPIXEL, &spp);
|
||||
bpp = bps * spp;
|
||||
if (bpp > 32)
|
||||
return (PIX *)ERROR_PTR("can't handle bpp > 32", procName, nullptr);
|
||||
if (spp == 1)
|
||||
d = bps;
|
||||
else if (spp == 3 || spp == 4)
|
||||
d = 32;
|
||||
else
|
||||
return (PIX *)ERROR_PTR("spp not in set {1,3,4}", procName, nullptr);
|
||||
|
||||
TIFFGetField(tif, TIFFTAG_IMAGEWIDTH, &w);
|
||||
TIFFGetField(tif, TIFFTAG_IMAGELENGTH, &h);
|
||||
tiffbpl = TIFFScanlineSize(tif);
|
||||
|
||||
if ((pix = pixCreate(w, h, d)) == nullptr)
|
||||
return (PIX *)ERROR_PTR("pix not made", procName, nullptr);
|
||||
data = (l_uint8 *)pixGetData(pix);
|
||||
wpl = pixGetWpl(pix);
|
||||
bpl = 4 * wpl;
|
||||
|
||||
if (spp == 1) {
|
||||
if ((linebuf = (l_uint8 *)CALLOC(tiffbpl + 1, sizeof(l_uint8))) ==
|
||||
nullptr)
|
||||
return (PIX *)ERROR_PTR("calloc fail for linebuf", procName, nullptr);
|
||||
|
||||
for (i = 0; i < h; i++) {
|
||||
if (TIFFReadScanline(tif, linebuf, i, 0) < 0) {
|
||||
FREE(linebuf);
|
||||
pixDestroy(&pix);
|
||||
return (PIX *)ERROR_PTR("line read fail", procName, nullptr);
|
||||
}
|
||||
memcpy((char *)data, (char *)linebuf, tiffbpl);
|
||||
data += bpl;
|
||||
}
|
||||
if (bps <= 8)
|
||||
pixEndianByteSwap(pix);
|
||||
else
|
||||
pixEndianTwoByteSwap(pix);
|
||||
FREE(linebuf);
|
||||
} else {
|
||||
if ((tiffdata = (l_uint32 *)CALLOC(w * h, sizeof(l_uint32))) == nullptr) {
|
||||
pixDestroy(&pix);
|
||||
return (PIX *)ERROR_PTR("calloc fail for tiffdata", procName, nullptr);
|
||||
}
|
||||
if (!TIFFReadRGBAImageOriented(tif, w, h, (uint32 *)tiffdata,
|
||||
ORIENTATION_TOPLEFT, 0)) {
|
||||
FREE(tiffdata);
|
||||
pixDestroy(&pix);
|
||||
return (PIX *)ERROR_PTR("failed to read tiffdata", procName, nullptr);
|
||||
}
|
||||
line = pixGetData(pix);
|
||||
|
||||
// Invoke the OpenCL kernel for pixReadFromTiff
|
||||
l_uint32 *output_gpu = pixReadFromTiffKernel(tiffdata, w, h, wpl, line);
|
||||
|
||||
pixSetData(pix, output_gpu);
|
||||
// pix already has data allocated, it now points to output_gpu?
|
||||
FREE(tiffdata);
|
||||
FREE(line);
|
||||
// FREE(output_gpu);
|
||||
}
|
||||
|
||||
if (getTiffStreamResolutionCl(tif, &xres, &yres) == 0) {
|
||||
pixSetXRes(pix, xres);
|
||||
pixSetYRes(pix, yres);
|
||||
}
|
||||
|
||||
|
||||
TIFFGetFieldDefaulted(tif, TIFFTAG_COMPRESSION, &tiffcomp);
|
||||
comptype = getTiffCompressedFormat(tiffcomp);
|
||||
pixSetInputFormat(pix, comptype);
|
||||
|
||||
if (TIFFGetField(tif, TIFFTAG_COLORMAP, &redmap, &greenmap, &bluemap)) {
|
||||
if ((cmap = pixcmapCreate(bps)) == nullptr) {
|
||||
pixDestroy(&pix);
|
||||
return (PIX *)ERROR_PTR("cmap not made", procName, nullptr);
|
||||
}
|
||||
ncolors = 1 << bps;
|
||||
for (i = 0; i < ncolors; i++)
|
||||
pixcmapAddColor(cmap, redmap[i] >> 8, greenmap[i] >> 8,
|
||||
bluemap[i] >> 8);
|
||||
pixSetColormap(pix, cmap);
|
||||
} else {
|
||||
if (!TIFFGetField(tif, TIFFTAG_PHOTOMETRIC, &photometry)) {
|
||||
if (tiffcomp == COMPRESSION_CCITTFAX3 ||
|
||||
tiffcomp == COMPRESSION_CCITTFAX4 ||
|
||||
tiffcomp == COMPRESSION_CCITTRLE ||
|
||||
tiffcomp == COMPRESSION_CCITTRLEW) {
|
||||
photometry = PHOTOMETRIC_MINISWHITE;
|
||||
} else
|
||||
photometry = PHOTOMETRIC_MINISBLACK;
|
||||
}
|
||||
if ((d == 1 && photometry == PHOTOMETRIC_MINISBLACK) ||
|
||||
(d == 8 && photometry == PHOTOMETRIC_MINISWHITE))
|
||||
pixInvert(pix, pix);
|
||||
}
|
||||
|
||||
if (TIFFGetField(tif, TIFFTAG_ORIENTATION, &orientation)) {
|
||||
if (orientation >= 1 && orientation <= 8) {
|
||||
struct tiff_transform *transform =
|
||||
&tiff_orientation_transforms[orientation - 1];
|
||||
if (transform->vflip) pixFlipTB(pix, pix);
|
||||
if (transform->hflip) pixFlipLR(pix, pix);
|
||||
if (transform->rotate) {
|
||||
PIX *oldpix = pix;
|
||||
pix = pixRotate90(oldpix, transform->rotate);
|
||||
pixDestroy(&oldpix);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
return pix;
|
||||
}
|
||||
|
||||
//Morphology Dilate operation for 5x5 structuring element. Invokes the relevant OpenCL kernels
|
||||
cl_int
|
||||
pixDilateCL_55(l_int32 wpl, l_int32 h)
|
||||
@ -2242,44 +1656,6 @@ cl_int pixErodeCL(l_int32 hsize, l_int32 vsize, l_uint32 wpl, l_uint32 h) {
|
||||
return status;
|
||||
}
|
||||
|
||||
// OpenCL implementation of Morphology Dilate
|
||||
//Note: Assumes the source and dest opencl buffer are initialized. No check done
|
||||
PIX *OpenclDevice::pixDilateBrickCL(PIX *pixd, PIX *pixs, l_int32 hsize,
|
||||
l_int32 vsize, bool reqDataCopy = false) {
|
||||
l_uint32 wpl, h;
|
||||
|
||||
wpl = pixGetWpl(pixs);
|
||||
h = pixGetHeight(pixs);
|
||||
|
||||
clStatus = pixDilateCL(hsize, vsize, wpl, h);
|
||||
|
||||
if (reqDataCopy) {
|
||||
pixd = mapOutputCLBuffer(rEnv, pixdCLBuffer, pixd, pixs, wpl * h,
|
||||
CL_MAP_READ, false);
|
||||
}
|
||||
|
||||
return pixd;
|
||||
}
|
||||
|
||||
// OpenCL implementation of Morphology Erode
|
||||
//Note: Assumes the source and dest opencl buffer are initialized. No check done
|
||||
PIX *OpenclDevice::pixErodeBrickCL(PIX *pixd, PIX *pixs, l_int32 hsize,
|
||||
l_int32 vsize, bool reqDataCopy = false) {
|
||||
l_uint32 wpl, h;
|
||||
|
||||
wpl = pixGetWpl(pixs);
|
||||
h = pixGetHeight(pixs);
|
||||
|
||||
clStatus = pixErodeCL(hsize, vsize, wpl, h);
|
||||
|
||||
if (reqDataCopy) {
|
||||
pixd =
|
||||
mapOutputCLBuffer(rEnv, pixdCLBuffer, pixd, pixs, wpl * h, CL_MAP_READ);
|
||||
}
|
||||
|
||||
return pixd;
|
||||
}
|
||||
|
||||
//Morphology Open operation. Invokes the relevant OpenCL kernels
|
||||
cl_int
|
||||
pixOpenCL(l_int32 hsize, l_int32 vsize, l_int32 wpl, l_int32 h)
|
||||
@ -2318,44 +1694,6 @@ pixCloseCL(l_int32 hsize, l_int32 vsize, l_int32 wpl, l_int32 h)
|
||||
return status;
|
||||
}
|
||||
|
||||
// OpenCL implementation of Morphology Close
|
||||
//Note: Assumes the source and dest opencl buffer are initialized. No check done
|
||||
PIX *OpenclDevice::pixCloseBrickCL(PIX *pixd, PIX *pixs, l_int32 hsize,
|
||||
l_int32 vsize, bool reqDataCopy = false) {
|
||||
l_uint32 wpl, h;
|
||||
|
||||
wpl = pixGetWpl(pixs);
|
||||
h = pixGetHeight(pixs);
|
||||
|
||||
clStatus = pixCloseCL(hsize, vsize, wpl, h);
|
||||
|
||||
if (reqDataCopy) {
|
||||
pixd =
|
||||
mapOutputCLBuffer(rEnv, pixdCLBuffer, pixd, pixs, wpl * h, CL_MAP_READ);
|
||||
}
|
||||
|
||||
return pixd;
|
||||
}
|
||||
|
||||
// OpenCL implementation of Morphology Open
|
||||
//Note: Assumes the source and dest opencl buffer are initialized. No check done
|
||||
PIX *OpenclDevice::pixOpenBrickCL(PIX *pixd, PIX *pixs, l_int32 hsize,
|
||||
l_int32 vsize, bool reqDataCopy = false) {
|
||||
l_uint32 wpl, h;
|
||||
|
||||
wpl = pixGetWpl(pixs);
|
||||
h = pixGetHeight(pixs);
|
||||
|
||||
clStatus = pixOpenCL(hsize, vsize, wpl, h);
|
||||
|
||||
if (reqDataCopy) {
|
||||
pixd =
|
||||
mapOutputCLBuffer(rEnv, pixdCLBuffer, pixd, pixs, wpl * h, CL_MAP_READ);
|
||||
}
|
||||
|
||||
return pixd;
|
||||
}
|
||||
|
||||
//pix OR operation: outbuffer = buffer1 | buffer2
|
||||
cl_int
|
||||
pixORCL_work(l_uint32 wpl, l_uint32 h, cl_mem buffer1, cl_mem buffer2, cl_mem outbuffer)
|
||||
@ -2394,45 +1732,6 @@ pixORCL_work(l_uint32 wpl, l_uint32 h, cl_mem buffer1, cl_mem buffer2, cl_mem ou
|
||||
return status;
|
||||
}
|
||||
|
||||
//pix AND operation: outbuffer = buffer1 & buffer2
|
||||
cl_int
|
||||
pixANDCL_work(l_uint32 wpl, l_uint32 h, cl_mem buffer1, cl_mem buffer2, cl_mem outbuffer)
|
||||
{
|
||||
cl_int status;
|
||||
size_t globalThreads[2];
|
||||
int gsize;
|
||||
size_t localThreads[] = {GROUPSIZE_X, GROUPSIZE_Y};
|
||||
|
||||
gsize = (wpl + GROUPSIZE_X - 1)/ GROUPSIZE_X * GROUPSIZE_X;
|
||||
globalThreads[0] = gsize;
|
||||
gsize = (h + GROUPSIZE_Y - 1)/ GROUPSIZE_Y * GROUPSIZE_Y;
|
||||
globalThreads[1] = gsize;
|
||||
|
||||
rEnv.mpkKernel = clCreateKernel( rEnv.mpkProgram, "pixAND", &status );
|
||||
CHECK_OPENCL(status, "clCreateKernel pixAND");
|
||||
|
||||
// Enqueue a kernel run call.
|
||||
status = clSetKernelArg(rEnv.mpkKernel,
|
||||
0,
|
||||
sizeof(cl_mem),
|
||||
&buffer1);
|
||||
status = clSetKernelArg(rEnv.mpkKernel,
|
||||
1,
|
||||
sizeof(cl_mem),
|
||||
&buffer2);
|
||||
status = clSetKernelArg(rEnv.mpkKernel,
|
||||
2,
|
||||
sizeof(cl_mem),
|
||||
&outbuffer);
|
||||
status = clSetKernelArg(rEnv.mpkKernel, 3, sizeof(wpl), &wpl);
|
||||
status = clSetKernelArg(rEnv.mpkKernel, 4, sizeof(h), &h);
|
||||
status = clEnqueueNDRangeKernel(rEnv.mpkCmdQueue, rEnv.mpkKernel, 2,
|
||||
nullptr, globalThreads, localThreads, 0,
|
||||
nullptr, nullptr);
|
||||
|
||||
return status;
|
||||
}
|
||||
|
||||
//output = buffer1 & ~(buffer2)
|
||||
cl_int pixSubtractCL_work(l_uint32 wpl, l_uint32 h, cl_mem buffer1,
|
||||
cl_mem buffer2, cl_mem outBuffer = nullptr) {
|
||||
@ -2470,81 +1769,6 @@ cl_int pixSubtractCL_work(l_uint32 wpl, l_uint32 h, cl_mem buffer1,
|
||||
return status;
|
||||
}
|
||||
|
||||
// OpenCL implementation of Subtract pix
|
||||
//Note: Assumes the source and dest opencl buffer are initialized. No check done
|
||||
PIX *OpenclDevice::pixSubtractCL(PIX *pixd, PIX *pixs1, PIX *pixs2,
|
||||
bool reqDataCopy = false) {
|
||||
l_uint32 wpl, h;
|
||||
|
||||
PROCNAME("pixSubtractCL");
|
||||
|
||||
if (!pixs1) return (PIX *)ERROR_PTR("pixs1 not defined", procName, pixd);
|
||||
if (!pixs2) return (PIX *)ERROR_PTR("pixs2 not defined", procName, pixd);
|
||||
if (pixGetDepth(pixs1) != pixGetDepth(pixs2))
|
||||
return (PIX *)ERROR_PTR("depths of pixs* unequal", procName, pixd);
|
||||
|
||||
#if EQUAL_SIZE_WARNING
|
||||
if (!pixSizesEqual(pixs1, pixs2))
|
||||
L_WARNING("pixs1 and pixs2 not equal sizes", procName);
|
||||
#endif /* EQUAL_SIZE_WARNING */
|
||||
|
||||
wpl = pixGetWpl(pixs1);
|
||||
h = pixGetHeight(pixs1);
|
||||
|
||||
clStatus = pixSubtractCL_work(wpl, h, pixdCLBuffer, pixsCLBuffer);
|
||||
|
||||
if (reqDataCopy)
|
||||
{
|
||||
//Read back output data from OCL buffer to cpu
|
||||
pixd = mapOutputCLBuffer(rEnv, pixdCLBuffer, pixd, pixs1, wpl*h, CL_MAP_READ);
|
||||
}
|
||||
|
||||
return pixd;
|
||||
}
|
||||
|
||||
// OpenCL implementation of Hollow pix
|
||||
//Note: Assumes the source and dest opencl buffer are initialized. No check done
|
||||
PIX *OpenclDevice::pixHollowCL(PIX *pixd, PIX *pixs, l_int32 close_hsize,
|
||||
l_int32 close_vsize, l_int32 open_hsize,
|
||||
l_int32 open_vsize, bool reqDataCopy = false) {
|
||||
l_uint32 wpl, h;
|
||||
cl_mem pixtemp;
|
||||
|
||||
wpl = pixGetWpl(pixs);
|
||||
h = pixGetHeight(pixs);
|
||||
|
||||
// First step : Close Morph operation: Dilate followed by Erode
|
||||
clStatus = pixCloseCL(close_hsize, close_vsize, wpl, h);
|
||||
|
||||
// Store the output of close operation in an intermediate buffer
|
||||
// this will be later used for pixsubtract
|
||||
clStatus =
|
||||
clEnqueueCopyBuffer(rEnv.mpkCmdQueue, pixdCLBuffer, pixdCLIntermediate, 0,
|
||||
0, sizeof(int) * wpl * h, 0, nullptr, nullptr);
|
||||
|
||||
// Second step: Open Operation - Erode followed by Dilate
|
||||
pixtemp = pixsCLBuffer;
|
||||
pixsCLBuffer = pixdCLBuffer;
|
||||
pixdCLBuffer = pixtemp;
|
||||
|
||||
clStatus = pixOpenCL(open_hsize, open_vsize, wpl, h);
|
||||
|
||||
// Third step: Subtract : (Close - Open)
|
||||
pixtemp = pixsCLBuffer;
|
||||
pixsCLBuffer = pixdCLBuffer;
|
||||
pixdCLBuffer = pixdCLIntermediate;
|
||||
pixdCLIntermediate = pixtemp;
|
||||
|
||||
clStatus = pixSubtractCL_work(wpl, h, pixdCLBuffer, pixsCLBuffer);
|
||||
|
||||
if (reqDataCopy) {
|
||||
// Read back output data from OCL buffer to cpu
|
||||
pixd =
|
||||
mapOutputCLBuffer(rEnv, pixdCLBuffer, pixd, pixs, wpl * h, CL_MAP_READ);
|
||||
}
|
||||
return pixd;
|
||||
}
|
||||
|
||||
// OpenCL implementation of Get Lines from pix function
|
||||
//Note: Assumes the source and dest opencl buffer are initialized. No check done
|
||||
void OpenclDevice::pixGetLinesCL(PIX *pixd, PIX *pixs, PIX **pix_vline,
|
||||
|
||||
@ -241,13 +241,7 @@ public:
|
||||
static int BinaryGenerated( const char * clFileName, FILE ** fhandle );
|
||||
//static int CompileKernelFile( const char *filename, GPUEnv *gpuInfo, const char *buildOption );
|
||||
static l_uint32* pixReadFromTiffKernel(l_uint32 *tiffdata,l_int32 w,l_int32 h,l_int32 wpl, l_uint32 *line);
|
||||
static Pix* pixReadTiffCl( const char *filename, l_int32 n );
|
||||
static PIX * pixReadStreamTiffCl ( FILE *fp, l_int32 n );
|
||||
static PIX * pixReadMemTiffCl(const l_uint8 *data, size_t size, l_int32 n);
|
||||
static PIX* pixReadFromTiffStreamCl(TIFF *tif);
|
||||
static int composeRGBPixelCl(int *tiffdata,int *line,int h,int w);
|
||||
static l_int32 getTiffStreamResolutionCl(TIFF *tif,l_int32 *pxres,l_int32 *pyres);
|
||||
static TIFF* fopenTiffCl(FILE *fp,const char *modestring);
|
||||
|
||||
/* OpenCL implementations of Morphological operations*/
|
||||
|
||||
@ -255,24 +249,6 @@ public:
|
||||
static int initMorphCLAllocations(l_int32 wpl, l_int32 h, PIX* pixs);
|
||||
static void releaseMorphCLBuffers();
|
||||
|
||||
// OpenCL implementation of Morphology Dilate
|
||||
static PIX* pixDilateBrickCL(PIX *pixd, PIX *pixs, l_int32 hsize, l_int32 vsize, bool reqDataCopy);
|
||||
|
||||
// OpenCL implementation of Morphology Erode
|
||||
static PIX* pixErodeBrickCL(PIX *pixd, PIX *pixs, l_int32 hsize, l_int32 vsize, bool reqDataCopy);
|
||||
|
||||
// OpenCL implementation of Morphology Close
|
||||
static PIX* pixCloseBrickCL(PIX *pixd, PIX *pixs, l_int32 hsize, l_int32 vsize, bool reqDataCopy);
|
||||
|
||||
// OpenCL implementation of Morphology Open
|
||||
static PIX* pixOpenBrickCL(PIX *pixd, PIX *pixs, l_int32 hsize, l_int32 vsize, bool reqDataCopy);
|
||||
|
||||
// OpenCL implementation of Morphology Open
|
||||
static PIX* pixSubtractCL(PIX *pixd, PIX *pixs1, PIX *pixs2, bool reqDataCopy);
|
||||
|
||||
// OpenCL implementation of Morphology (Hollow = Closed - Open)
|
||||
static PIX* pixHollowCL(PIX *pixd, PIX *pixs, l_int32 close_hsize, l_int32 close_vsize, l_int32 open_hsize, l_int32 open_vsize, bool reqDataCopy);
|
||||
|
||||
static void pixGetLinesCL(PIX *pixd, PIX *pixs, PIX **pix_vline,
|
||||
PIX **pix_hline, PIX **pixClosed,
|
||||
bool getpixClosed, l_int32 close_hsize,
|
||||
@ -290,8 +266,6 @@ public:
|
||||
//int RegisterKernelWrapper( const char *kernelName, cl_kernel_function function );
|
||||
//int RunKernelWrapper( cl_kernel_function function, const char * kernelName, void **usrdata );
|
||||
//int GetKernelEnvAndFunc( const char *kernelName, KernelEnv *env, cl_kernel_function *function );
|
||||
// static cl_device_id performDeviceSelection( );
|
||||
//static bool thresholdRectToPixMicroBench( TessScoreEvaluationInputData input, ds_device_type type);
|
||||
|
||||
static int LoadOpencl();
|
||||
#ifdef WIN32
|
||||
|
||||
Loading…
Reference in New Issue
Block a user