Merge pull request #849 from stweil/opencl

opencl: Remove more unused functions and related code
This commit is contained in:
zdenop 2017-04-27 09:02:56 +02:00 committed by GitHub
commit 10e04ffe99
2 changed files with 0 additions and 802 deletions

View File

@ -843,7 +843,6 @@ int OpenclDevice::InitOpenclRunEnv_DeviceSelection( int argc ) {
//PERF_COUNT_START("InitOpenclRunEnv_DS") //PERF_COUNT_START("InitOpenclRunEnv_DS")
if (!isInited) { if (!isInited) {
// after programs compiled, selects best device // after programs compiled, selects best device
//printf("[DS] InitOpenclRunEnv_DS::Calling performDeviceSelection()\n");
ds_device bestDevice_DS = getDeviceSelection( ); ds_device bestDevice_DS = getDeviceSelection( );
//PERF_COUNT_SUB("called getDeviceSelection()") //PERF_COUNT_SUB("called getDeviceSelection()")
cl_device_id bestDevice = bestDevice_DS.oclDeviceID; cl_device_id bestDevice = bestDevice_DS.oclDeviceID;
@ -1065,22 +1064,6 @@ int OpenclDevice::GeneratBinFromKernelSource( cl_program program, const char * c
return 1; 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 ) int OpenclDevice::CompileKernelFile( GPUEnv *gpuInfo, const char *buildOption )
{ {
//PERF_COUNT_START("CompileKernelFile") //PERF_COUNT_START("CompileKernelFile")
@ -1317,575 +1300,6 @@ PERF_COUNT_END
return pResult; 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 //Morphology Dilate operation for 5x5 structuring element. Invokes the relevant OpenCL kernels
cl_int cl_int
pixDilateCL_55(l_int32 wpl, l_int32 h) 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; 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 //Morphology Open operation. Invokes the relevant OpenCL kernels
cl_int cl_int
pixOpenCL(l_int32 hsize, l_int32 vsize, l_int32 wpl, l_int32 h) 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; 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 //pix OR operation: outbuffer = buffer1 | buffer2
cl_int cl_int
pixORCL_work(l_uint32 wpl, l_uint32 h, cl_mem buffer1, cl_mem buffer2, cl_mem outbuffer) 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; 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) //output = buffer1 & ~(buffer2)
cl_int pixSubtractCL_work(l_uint32 wpl, l_uint32 h, cl_mem buffer1, cl_int pixSubtractCL_work(l_uint32 wpl, l_uint32 h, cl_mem buffer1,
cl_mem buffer2, cl_mem outBuffer = nullptr) { 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; 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 // OpenCL implementation of Get Lines from pix function
//Note: Assumes the source and dest opencl buffer are initialized. No check done //Note: Assumes the source and dest opencl buffer are initialized. No check done
void OpenclDevice::pixGetLinesCL(PIX *pixd, PIX *pixs, PIX **pix_vline, void OpenclDevice::pixGetLinesCL(PIX *pixd, PIX *pixs, PIX **pix_vline,

View File

@ -241,13 +241,7 @@ public:
static int BinaryGenerated( const char * clFileName, FILE ** fhandle ); static int BinaryGenerated( const char * clFileName, FILE ** fhandle );
//static int CompileKernelFile( const char *filename, GPUEnv *gpuInfo, const char *buildOption ); //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 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 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*/ /* OpenCL implementations of Morphological operations*/
@ -255,24 +249,6 @@ public:
static int initMorphCLAllocations(l_int32 wpl, l_int32 h, PIX* pixs); static int initMorphCLAllocations(l_int32 wpl, l_int32 h, PIX* pixs);
static void releaseMorphCLBuffers(); 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, static void pixGetLinesCL(PIX *pixd, PIX *pixs, PIX **pix_vline,
PIX **pix_hline, PIX **pixClosed, PIX **pix_hline, PIX **pixClosed,
bool getpixClosed, l_int32 close_hsize, bool getpixClosed, l_int32 close_hsize,
@ -290,8 +266,6 @@ public:
//int RegisterKernelWrapper( const char *kernelName, cl_kernel_function function ); //int RegisterKernelWrapper( const char *kernelName, cl_kernel_function function );
//int RunKernelWrapper( cl_kernel_function function, const char * kernelName, void **usrdata ); //int RunKernelWrapper( cl_kernel_function function, const char * kernelName, void **usrdata );
//int GetKernelEnvAndFunc( const char *kernelName, KernelEnv *env, cl_kernel_function *function ); //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(); static int LoadOpencl();
#ifdef WIN32 #ifdef WIN32