Inexplicable while loop behavior

OpenCL
1

Hello,

I’m having some trouble figuring out why a while loop in my code is bailing. I’m wondering if someone can see something, hopefully, obvious I’ve missed.

Thanks!

Here’s my kernel code:

static uint append(__global char log, uint logPos, __constant char string) {
uint idx=0;
while(string[idx]!=0) {
log[logPos++] = string[idx++];
}
log[logPos]=0;
return logPos;
}

__kernel void calc_edit_distance(
const uint widthIn, // needle and each in haystack width
__constant ulong *needleIn, // needle uint64_t’s
__constant ulong *haystackIn, // haystack uint64_t’s
__global ulong *distancesOut, // results
__global ulong *operationsOut, // the operations to transform the haystack element into the needle
__global char *logOut
) {
uint haystackRowIdx = get_global_id(0);

uint logPos = haystackRowIdx * 2048;

uint needleIdx = 0;
uint haystackIdx = 0;
ulong needleLast = 0;
ulong haystackLast = 0;

ulong distanceTotal = 0;

logPos=append(logOut,logPos,"Starting

");
while(needleIdx<widthIn && haystackIdx<widthIn && (logPos=append(logOut,logPos,"Iteration Start
"))) {

    logPos=append(logOut,logPos,"Here

");

    ulong needleCur = needleIn[needleIdx];
    ulong haystackCur = haystackIn[(widthIn*haystackRowIdx)+haystackIdx];
    
    if(needleCur == haystackCur) {
        if(haystackCur == needleLast) {
            if(needleCur == haystackLast) {
                if(needleLast == haystackLast) { 
                    // 1111 - error: never been here before
                    logPos=append(logOut,logPos,"1111 - error: never been here before

");
} else {
// 1110 - error: never been here before
logPos=append(logOut,logPos,"1110 - error: never been here before
");
}
} else {
if(needleLast == haystackLast) {
// 1101 - error: never been here before
logPos=append(logOut,logPos,"1101 - error: never been here before
");
} else {
// 1100 - error: never been here before
logPos=append(logOut,logPos,"1100 - error: never been here before
");
}
}
} else {
if(needleCur == haystackLast) {
if(needleLast == haystackLast) {
// 1011 - error: never been here before
logPos=append(logOut,logPos,"1011 - error: never been here before
");
} else {
// 1010 - possible insertion
logPos=append(logOut,logPos,"1010 - possible insertion
");
distanceTotal++;
}
} else {
if(needleLast == haystackLast) {
// 1001 - continuing match
logPos=append(logOut,logPos,"1001 - continuing match
");
} else {
// 1000 - match restored
logPos=append(logOut,logPos,"1000 - match restored
");
}
}
}
} else {
if(haystackCur == needleLast) {
if(needleCur == haystackLast) {
if(needleLast == haystackLast) {
// 0111 - error: never been here before
logPos=append(logOut,logPos,"0111 - error: never been here before
");
} else {
// 0110 - transposition
logPos=append(logOut,logPos,"0110 - transposition
");
distanceTotal++;
}
} else {
if(needleLast == haystackLast) {
// 0101 - repeat haystack
logPos=append(logOut,logPos,"0101 - repeat haystack
");
distanceTotal++;
} else {
// 0100 - may be restoration of match
logPos=append(logOut,logPos,"0100 - may be restoration of match
");
needleIdx–;
continue;
}
}
} else {
if(needleCur == haystackLast) {
if(needleLast == haystackLast) {
// 0011 - replacement or deletion
logPos=append(logOut,logPos,"0011 - replacement or deletion
");
distanceTotal++;
} else {
// 0010 - restore match, last was actually ommission in haystack
logPos=append(logOut,logPos,"0010 - restore match, last was actually ommission in haystack
");
haystackIdx–;
continue;
}
} else {
if(needleLast == haystackLast) {
// 0001 - break match, replacement
logPos=append(logOut,logPos,"0001 - break match, replacement
");
distanceTotal++;
} else {
// 0000 - continue broken match, replacement
logPos=append(logOut,logPos,"0000 - continue broken match, replacement
");
distanceTotal++;
}
}
}
}

    needleIdx++;
    haystackIdx++;
    needleLast = needleCur;
    haystackLast = haystackCur;
    
    logPos=append(logOut,logPos,(needleIdx&lt;widthIn?"needleIdx&lt;widthIn = true; ":"needleIdx&lt;widthIn = false; "));
    logPos=append(logOut,logPos,(haystackIdx&lt;widthIn?"haystackIdx&lt;widthIn = true

":"haystackIdx<widthIn = false
"));
logPos=append(logOut,logPos,"Iteration End
");
}
logPos=append(logOut,logPos,"Ending
");
distancesOut[haystackRowIdx] = distanceTotal;
}

Here’s my first data set:
uint16_t width=4;
uint16_t haystackSize=1;
uint64_t needle[] = {1,2,3,4};
uint64_t haystack[] = {1,2,3,4};
uint64_t distancesOut[1];
uint64_t operationsOut[haystackSize*(2*width)];

With it’s log output:

Starting
Iteration Start
Here
1001 - continuing match
needleIdx<widthIn = true; haystackIdx<widthIn = true
Iteration End
Iteration Start
Here
1001 - continuing match
needleIdx<widthIn = true; haystackIdx<widthIn = true
Iteration End
Iteration Start
Here
1001 - continuing match
needleIdx<widthIn = true; haystackIdx<widthIn = true
Iteration End
Iteration Start
Here
1001 - continuing match
needleIdx<widthIn = false; haystackIdx<widthIn = false
Iteration End
Ending

And my second data set:

    uint16_t width=4;
    uint16_t haystackSize=1;
    uint64_t needle[] = {1,2,3,4};
    uint64_t haystack[] = {4,3,2,1};
    uint64_t distancesOut[1];
    uint64_t operationsOut[haystackSize*(2*width)];

And here’s the log result I got:

Starting
Iteration Start
Here
0001 - break match, replacement
needleIdx<widthIn = true; haystackIdx<widthIn = true
Iteration End
Iteration Start
Ending

I’m really baffled at how it got to the second “Iteration Start” log statement, but then failed to reenter the loop. Anyone got any ideas? Stumped. =(