mirror of
https://github.com/tesseract-ocr/tesseract.git
synced 2025-01-18 14:41:36 +08:00
opencl: Clean whitespace issues in OpenCL kernel code
* Remove whitespace at line endings * Replace tabs by spaces Signed-off-by: Stefan Weil <sw@weilnetz.de>
This commit is contained in:
Stefan Weil
parent
a8f444112e
commit
0020fbc0bd
@ -58,7 +58,7 @@ KERNEL(
|
||||
)
|
||||
|
||||
KERNEL(
|
||||
\n__kernel void pixSubtract(__global int *dword, __global int *sword,
|
||||
\n__kernel void pixSubtract(__global int *dword, __global int *sword,
|
||||
const int wpl, const int h, __global int *outword)
|
||||
{
|
||||
const unsigned int row = get_global_id(1);
|
||||
@ -113,15 +113,15 @@ KERNEL(
|
||||
unsigned int prevword, nextword, currword,tempword;
|
||||
unsigned int destword;
|
||||
const int col = pos % wpl;
|
||||
|
||||
|
||||
//Ignore the execss
|
||||
if (pos >= (wpl * h))
|
||||
return;
|
||||
|
||||
|
||||
currword = *(sword + pos);
|
||||
|
||||
|
||||
currword = *(sword + pos);
|
||||
destword = currword;
|
||||
|
||||
|
||||
//Handle boundary conditions
|
||||
if(col==0)
|
||||
prevword=0;
|
||||
@ -132,9 +132,9 @@ KERNEL(
|
||||
nextword=0;
|
||||
else
|
||||
nextword = *(sword + pos + 1);
|
||||
|
||||
|
||||
//Loop unrolled
|
||||
|
||||
|
||||
//1 bit to left and 1 bit to right
|
||||
//Get the max value on LHS of every pixel
|
||||
tempword = (prevword << (31)) | ((currword >> 1));
|
||||
@ -150,10 +150,10 @@ KERNEL(
|
||||
//Get max value on RHS of every pixel
|
||||
tempword = (currword << 2) | (nextword >> (30));
|
||||
destword |= tempword;
|
||||
|
||||
|
||||
|
||||
|
||||
*(dword + pos) = destword;
|
||||
|
||||
|
||||
}\n
|
||||
)
|
||||
|
||||
@ -208,7 +208,7 @@ KERNEL(
|
||||
unsigned int destword, tempword, lastword, currword;
|
||||
unsigned int lnextword, lprevword, rnextword, rprevword, firstword, secondword;
|
||||
int i, j, siter, eiter;
|
||||
|
||||
|
||||
//Ignore the execss
|
||||
if (pos >= (wpl*h) || (xn < 1 && xp < 1))
|
||||
return;
|
||||
@ -233,7 +233,7 @@ KERNEL(
|
||||
firstword = 0x0;
|
||||
else
|
||||
firstword = *(sword + pos - 1);
|
||||
|
||||
|
||||
//Get next word
|
||||
if (col == (wpl - 1))
|
||||
secondword = 0x0;
|
||||
@ -245,7 +245,7 @@ KERNEL(
|
||||
{
|
||||
//Get the max value on LHS of every pixel
|
||||
tempword = ((i == parbitsxp) && (parbitsxp != parbitsxn)) ? 0x0 : (firstword << (32-i)) | ((currword >> i));
|
||||
|
||||
|
||||
destword |= tempword;
|
||||
|
||||
//Get max value on RHS of every pixel
|
||||
@ -274,11 +274,11 @@ KERNEL(
|
||||
else
|
||||
firstword = *(sword + row*wpl + siter);
|
||||
|
||||
if (eiter >= wpl)
|
||||
if (eiter >= wpl)
|
||||
lastword = 0x0;
|
||||
else
|
||||
lastword = *(sword + row*wpl + eiter);
|
||||
|
||||
|
||||
for ( i = 1; i < nwords; i++)
|
||||
{
|
||||
//Gets LHS words
|
||||
@ -288,14 +288,14 @@ KERNEL(
|
||||
secondword = *(sword + row*wpl + siter + i);
|
||||
|
||||
lprevword = firstword << (32 - parbitsxn) | secondword >> parbitsxn;
|
||||
|
||||
|
||||
firstword = secondword;
|
||||
|
||||
if ((siter + i + 1) < 0)
|
||||
secondword = 0x0;
|
||||
else
|
||||
secondword = *(sword + row*wpl + siter + i + 1);
|
||||
|
||||
|
||||
lnextword = firstword << (32 - parbitsxn) | secondword >> parbitsxn;
|
||||
|
||||
//Gets RHS words
|
||||
@ -303,7 +303,7 @@ KERNEL(
|
||||
firstword = 0x0;
|
||||
else
|
||||
firstword = *(sword + row*wpl + eiter - i);
|
||||
|
||||
|
||||
rnextword = firstword << parbitsxp | lastword >> (32 - parbitsxp);
|
||||
|
||||
lastword = firstword;
|
||||
@ -333,7 +333,7 @@ KERNEL(
|
||||
lastword = firstword;
|
||||
firstword = secondword;
|
||||
}
|
||||
|
||||
|
||||
*(dword + pos) = destword;
|
||||
}\n
|
||||
)
|
||||
@ -350,14 +350,14 @@ KERNEL(
|
||||
unsigned int prevword, nextword, currword,tempword;
|
||||
unsigned int destword;
|
||||
int i;
|
||||
|
||||
|
||||
//Ignore the execss
|
||||
if (pos >= (wpl * h))
|
||||
return;
|
||||
|
||||
currword = *(sword + pos);
|
||||
currword = *(sword + pos);
|
||||
destword = currword;
|
||||
|
||||
|
||||
//Handle boundary conditions
|
||||
if(col==0)
|
||||
prevword=0;
|
||||
@ -368,7 +368,7 @@ KERNEL(
|
||||
nextword=0;
|
||||
else
|
||||
nextword = *(sword + pos + 1);
|
||||
|
||||
|
||||
for (i = 1; i <= halfwidth; i++)
|
||||
{
|
||||
//Get the max value on LHS of every pixel
|
||||
@ -385,7 +385,7 @@ KERNEL(
|
||||
|
||||
//Get max value on RHS of every pixel
|
||||
tempword = (currword << i) | (nextword >> (32 - i));
|
||||
|
||||
|
||||
destword |= tempword;
|
||||
}
|
||||
|
||||
@ -405,7 +405,7 @@ KERNEL(
|
||||
unsigned int tempword;
|
||||
unsigned int destword;
|
||||
int i, siter, eiter;
|
||||
|
||||
|
||||
//Ignore the execss
|
||||
if (row >= h || col >= wpl)
|
||||
return;
|
||||
@ -435,27 +435,27 @@ KERNEL(
|
||||
unsigned int prevword, nextword, currword,tempword;
|
||||
unsigned int destword;
|
||||
const int col = pos % wpl;
|
||||
|
||||
|
||||
//Ignore the execss
|
||||
if (pos >= (wpl * h))
|
||||
return;
|
||||
|
||||
currword = *(sword + pos);
|
||||
|
||||
currword = *(sword + pos);
|
||||
destword = currword;
|
||||
|
||||
|
||||
//Handle boundary conditions
|
||||
if(col==0)
|
||||
prevword=0xffffffff;
|
||||
else
|
||||
prevword = *(sword + pos - 1);
|
||||
|
||||
|
||||
if(col==(wpl - 1))
|
||||
nextword=0xffffffff;
|
||||
else
|
||||
nextword = *(sword + pos + 1);
|
||||
|
||||
|
||||
//Loop unrolled
|
||||
|
||||
|
||||
//1 bit to left and 1 bit to right
|
||||
//Get the min value on LHS of every pixel
|
||||
tempword = (prevword << (31)) | ((currword >> 1));
|
||||
@ -471,10 +471,10 @@ KERNEL(
|
||||
//Get min value on RHS of every pixel
|
||||
tempword = (currword << 2) | (nextword >> (30));
|
||||
destword &= tempword;
|
||||
|
||||
|
||||
|
||||
|
||||
*(dword + pos) = destword;
|
||||
|
||||
|
||||
}\n
|
||||
)
|
||||
|
||||
@ -499,7 +499,7 @@ KERNEL(
|
||||
if (row < 2 || row >= (h - 2))
|
||||
{
|
||||
destword = 0x0;
|
||||
}
|
||||
}
|
||||
else
|
||||
{
|
||||
//2 words above
|
||||
@ -526,7 +526,7 @@ KERNEL(
|
||||
tempword = *(sword + i*wpl + col);
|
||||
destword &= tempword;
|
||||
|
||||
if (col == 0)
|
||||
if (col == 0)
|
||||
{
|
||||
destword &= fwmask;
|
||||
}
|
||||
@ -542,7 +542,7 @@ KERNEL(
|
||||
)
|
||||
|
||||
KERNEL(
|
||||
\n__kernel void morphoErodeHor(__global int *sword,__global int *dword, const int xp, const int xn, const int wpl,
|
||||
\n__kernel void morphoErodeHor(__global int *sword,__global int *dword, const int xp, const int xn, const int wpl,
|
||||
const int h, const char isAsymmetric, const int rwmask, const int lwmask)
|
||||
{
|
||||
const int col = get_global_id(0);
|
||||
@ -577,7 +577,7 @@ KERNEL(
|
||||
firstword = 0xffffffff;
|
||||
else
|
||||
firstword = *(sword + pos - 1);
|
||||
|
||||
|
||||
//Get next word
|
||||
if (col == (wpl - 1))
|
||||
secondword = 0xffffffff;
|
||||
@ -593,7 +593,7 @@ KERNEL(
|
||||
|
||||
//Get max value on RHS of every pixel
|
||||
tempword = ((i == parbitsxp) && (parbitsxp != parbitsxn)) ? 0xffffffff : (currword << i) | (secondword >> (32 - i));
|
||||
|
||||
|
||||
//tempword = (currword << i) | (secondword >> (32 - i));
|
||||
destword &= tempword;
|
||||
}
|
||||
@ -622,18 +622,18 @@ KERNEL(
|
||||
*(dword + pos) = destword;
|
||||
return;
|
||||
}
|
||||
|
||||
|
||||
if (siter < 0)
|
||||
firstword = 0xffffffff;
|
||||
else
|
||||
firstword = *(sword + row*wpl + siter);
|
||||
|
||||
if (eiter >= wpl)
|
||||
if (eiter >= wpl)
|
||||
lastword = 0xffffffff;
|
||||
else
|
||||
lastword = *(sword + row*wpl + eiter);
|
||||
|
||||
|
||||
|
||||
|
||||
for ( i = 1; i < nwords; i++)
|
||||
{
|
||||
//Gets LHS words
|
||||
@ -643,14 +643,14 @@ KERNEL(
|
||||
secondword = *(sword + row*wpl + siter + i);
|
||||
|
||||
lprevword = firstword << (32 - parbitsxp) | secondword >> (parbitsxp);
|
||||
|
||||
|
||||
firstword = secondword;
|
||||
|
||||
if ((siter + i + 1) < 0)
|
||||
secondword = 0xffffffff;
|
||||
else
|
||||
secondword = *(sword + row*wpl + siter + i + 1);
|
||||
|
||||
|
||||
lnextword = firstword << (32 - parbitsxp) | secondword >> (parbitsxp);
|
||||
|
||||
//Gets RHS words
|
||||
@ -658,7 +658,7 @@ KERNEL(
|
||||
firstword = 0xffffffff;
|
||||
else
|
||||
firstword = *(sword + row*wpl + eiter - i);
|
||||
|
||||
|
||||
rnextword = firstword << parbitsxn | lastword >> (32 - parbitsxn);
|
||||
|
||||
lastword = firstword;
|
||||
@ -688,7 +688,7 @@ KERNEL(
|
||||
lastword = firstword;
|
||||
firstword = secondword;
|
||||
}
|
||||
|
||||
|
||||
if (isAsymmetric)
|
||||
{
|
||||
//Clear boundary pixels
|
||||
@ -708,8 +708,8 @@ KERNEL(
|
||||
|
||||
KERNEL(
|
||||
\n__kernel void morphoErodeHor_32word(__global int *sword,__global int *dword,
|
||||
const int halfwidth, const int wpl,
|
||||
const int h, const char clearBoundPixH,
|
||||
const int halfwidth, const int wpl,
|
||||
const int h, const char clearBoundPixH,
|
||||
const int rwmask, const int lwmask,
|
||||
const char isEven)
|
||||
{
|
||||
@ -723,25 +723,25 @@ KERNEL(
|
||||
if (pos >= (wpl * h))
|
||||
return;
|
||||
|
||||
currword = *(sword + pos);
|
||||
currword = *(sword + pos);
|
||||
destword = currword;
|
||||
|
||||
|
||||
//Handle boundary conditions
|
||||
if(col==0)
|
||||
prevword=0xffffffff;
|
||||
else
|
||||
prevword = *(sword + pos - 1);
|
||||
|
||||
|
||||
if(col==(wpl - 1))
|
||||
nextword=0xffffffff;
|
||||
else
|
||||
nextword = *(sword + pos + 1);
|
||||
|
||||
|
||||
for (i = 1; i <= halfwidth; i++)
|
||||
{
|
||||
//Get the min value on LHS of every pixel
|
||||
tempword = (prevword << (32-i)) | ((currword >> i));
|
||||
|
||||
|
||||
destword &= tempword;
|
||||
|
||||
//Get min value on RHS of every pixel
|
||||
@ -759,7 +759,7 @@ KERNEL(
|
||||
|
||||
if (clearBoundPixH)
|
||||
{
|
||||
if (col == 0)
|
||||
if (col == 0)
|
||||
{
|
||||
destword &= rwmask;
|
||||
}
|
||||
@ -775,7 +775,7 @@ KERNEL(
|
||||
|
||||
KERNEL(
|
||||
\n__kernel void morphoErodeVer(__global int *sword,__global int *dword,
|
||||
const int yp,
|
||||
const int yp,
|
||||
const int wpl, const int h,
|
||||
const char clearBoundPixV, const int yn)
|
||||
{
|
||||
@ -784,7 +784,7 @@ KERNEL(
|
||||
const unsigned int pos = row * wpl + col;
|
||||
unsigned int tempword, destword;
|
||||
int i, siter, eiter;
|
||||
|
||||
|
||||
//Ignore the execss
|
||||
if (row >= h || col >= wpl)
|
||||
return;
|
||||
@ -804,7 +804,7 @@ KERNEL(
|
||||
|
||||
//Clear boundary pixels
|
||||
if (clearBoundPixV && ((row < yp) || ((h - row) <= yn)))
|
||||
{
|
||||
{
|
||||
destword = 0x0;
|
||||
}
|
||||
|
||||
@ -892,23 +892,23 @@ KERNEL(
|
||||
\n __global const uchar* data,
|
||||
\n uint numPixels,
|
||||
\n __global uint *histBuffer) { // each wg will write HIST_SIZE*NUM_CHANNELS into this result; cpu will accumulate across wg's
|
||||
\n
|
||||
\n
|
||||
\n /* declare variables */
|
||||
\n
|
||||
\n
|
||||
\n // work indices
|
||||
\n size_t groupId = get_group_id(0);
|
||||
\n size_t localId = get_local_id(0); // 0 -> 256-1
|
||||
\n size_t globalId = get_global_id(0); // 0 -> 8*10*256-1=20480-1
|
||||
\n uint numThreads = get_global_size(0);
|
||||
\n
|
||||
\n
|
||||
\n /* accumulate in global memory */
|
||||
\n for ( uint pc = get_global_id(0); pc < numPixels; pc += get_global_size(0) ) {
|
||||
\n uchar value = data[ pc ];
|
||||
\n int idx = value * get_global_size(0) + get_global_id(0);
|
||||
\n histBuffer[ idx ]++;
|
||||
\n
|
||||
\n
|
||||
\n }
|
||||
\n
|
||||
\n
|
||||
\n } // kernel_HistogramRectAllChannels_Grey
|
||||
|
||||
)
|
||||
@ -1001,35 +1001,35 @@ void kernel_HistogramRectOneChannelReduction(
|
||||
|
||||
KERNEL(
|
||||
// unused
|
||||
// each work group (x256) handles a histogram bin
|
||||
// each work group (x256) handles a histogram bin
|
||||
\n __attribute__((reqd_work_group_size(256, 1, 1)))
|
||||
\n __kernel
|
||||
\n void kernel_HistogramRectAllChannelsReduction_Grey(
|
||||
\n int n, // pixel redundancy that needs to be accumulated
|
||||
\n __global uint *histBuffer,
|
||||
\n __global uint* histResult) { // each wg accumulates 1 bin
|
||||
\n
|
||||
\n
|
||||
\n /* declare variables */
|
||||
\n
|
||||
\n
|
||||
\n // work indices
|
||||
\n size_t groupId = get_group_id(0);
|
||||
\n size_t localId = get_local_id(0); // 0 -> 256-1
|
||||
\n size_t globalId = get_global_id(0); // 0 -> 8*10*256-1=20480-1
|
||||
\n uint numThreads = get_global_size(0);
|
||||
\n unsigned int hist = 0;
|
||||
\n
|
||||
\n
|
||||
\n /* accumulate in global memory */
|
||||
\n for ( uint p = 0; p < n; p+=GROUP_SIZE) {
|
||||
\n hist += histBuffer[ (get_group_id(0)*n + p)];
|
||||
\n }
|
||||
\n
|
||||
\n
|
||||
\n /* reduction in local memory */
|
||||
\n // populate local memory
|
||||
\n __local unsigned int localHist[GROUP_SIZE];
|
||||
|
||||
\n localHist[localId] = hist;
|
||||
\n barrier(CLK_LOCAL_MEM_FENCE);
|
||||
\n
|
||||
\n
|
||||
\n for (int stride = GROUP_SIZE/2; stride >= 1; stride /= 2) {
|
||||
\n if (localId < stride) {
|
||||
\n hist = localHist[ (localId+stride)];
|
||||
@ -1040,10 +1040,10 @@ KERNEL(
|
||||
\n }
|
||||
\n barrier(CLK_LOCAL_MEM_FENCE);
|
||||
\n }
|
||||
\n
|
||||
\n
|
||||
\n if (localId == 0)
|
||||
\n histResult[get_group_id(0)] = localHist[0];
|
||||
\n
|
||||
\n
|
||||
\n } // kernel_HistogramRectAllChannelsReduction_Grey
|
||||
|
||||
)
|
||||
@ -1155,10 +1155,10 @@ void kernel_ThresholdRectToPix_OneChan(
|
||||
|
||||
// for each pixel in burst
|
||||
for ( int p = 0; p < PIXELS_PER_BURST; p++) {
|
||||
|
||||
|
||||
//int littleEndianIdx = p ^ 3;
|
||||
//int bigEndianIdx = p;
|
||||
int idx =
|
||||
int idx =
|
||||
\n#ifdef __ENDIAN_LITTLE__\n
|
||||
p ^ 3;
|
||||
\n#else\n
|
||||
@ -1179,30 +1179,30 @@ void kernel_ThresholdRectToPix_OneChan(
|
||||
|
||||
|
||||
KERNEL(
|
||||
\n#define RED_SHIFT 24\n
|
||||
\n#define GREEN_SHIFT 16\n
|
||||
\n#define BLUE_SHIFT 8\n
|
||||
\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,
|
||||
__global unsigned char *dstData,
|
||||
int srcWPL,
|
||||
int dstWPL,
|
||||
int height,
|
||||
int width,
|
||||
float rwt,
|
||||
float gwt,
|
||||
float bwt ) {
|
||||
|
||||
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) +
|
||||
unsigned int word = srcData[pixelIdx];
|
||||
int output = (rwt * ((word >> RED_SHIFT) & 0xff) +
|
||||
gwt * ((word >> GREEN_SHIFT) & 0xff) +
|
||||
bwt * ((word >> BLUE_SHIFT) & 0xff) + 0.5);
|
||||
// SET_DATA_BYTE
|
||||
|
||||
Loading…
Reference in New Issue
Block a user