Mirrors number at borders of interval, untill it lays in the interval

Question

I have the following simple function:

__device__
mirror(int index , int lB, int uB)
{
    while(index < lB || index >= uB)
    {
        if(index < lB) {
            index = lB + (lB-index);
        }

        if(index >= uB) {
            index = uB-1 -(index-uB);
        }
    }
    return index;
}

Assume lb <= uB.

It is used for every index to give it a mirror border behaviour and make sure the result is a valid index.

Are there ways to improve this, with or without changing the result for indices outside [lB-d,uB+d), where d >= uB-lB.

mirror(i,lB,uB) has to be in [lB,uB) of course.

What are concerns for performance?

• this function is called very often
• it is called from an kernel (cuda things relevant? thread divergence?)

Are there possible improvments for the special case lB == 0?

EDIT: added this small test example

__global__
void driver(unsigned char* img_dest, unsigned char* img_src, int width, int height)
{
    int x = blockIdx.x*blockDim.x+threadIdx.x;
    int y = blockIdx.y*blockDim.y+threadIdx.y;
    if(x >= 0 && x < width && y >= 0 && y < height)
    {
        int x_src = mirror(x*2-width/2,0,width);
        assert(x_src >= 0);
        assert(x_src < width);
        int y_src = mirror(y+x,0,height);
        assert(y_src >= 0);
        assert(y_src < height);
        img_dest[x+y*width] = img_src[x_src + y_src*width];
    }
}

// let the images be gray images, with no pitch.
void callDriver(unsigned char* img_dest, unsigned char* img_src, int width,int height)
{
    dim3 block(16,16);
    dim3 grid((width+block.x-1)/block.x,(height+block.y-1)/block.y);
    driver<<<grid,block>>>(img_dest,img_src,width,height);
    cudaDeviceSynchronize();
}


main()
{
    int width = 512;
    int height = 256;
    unsigned char* img_dest;
    cudaMalloc(&img_dest,width*height);
    unsigned char* img_src;
    cudaMalloc(&img_src,width*height);
    callDriver(img_dest,img_src,width,height);
}

My use case is similar to the driver example: I have functions deforming where I would have to load pixel that would lay outside of the defined region. To provide meaningful values I want to define them by mirroring at the borders.

My use case are distortion, and linear filters. In both cases I have to use values for pixels that are not defined in the image. While in the case of linear filters, there are other possible solutions (starting extra kernel for border), I decided to work with one kernel for this time.

Answer 1

__device__
mirror(int index , int lB, int uB)

You've got an extra space before the first comma, and I think you're missing the function's return type entirely. You should always compile your code with -Wall and fix all the warnings before posting it (or before running it). There's no point in shipping code with bugs.

Similarly funky whitespace on this line:

            index = uB-1 -(index-uB);

---

Anyway, if you're looking to speed up a piece of code with loops in it, your first thought should be, "How do I get rid of these loops?"

Consider that after you "mirror" index off of both boundaries, you'll simply have reduced its overall value by 2*(uB - lB). Proof:

index = uB-1 - (index-uB);
index = lB + (lB-index);

means

index = lB + lB - (uB - 1 - index + uB);

means

index = 2*lB - 2*uB + 1 + index;

means

index = index - (2*(uB - lB) - 1);

Hmm... I see you have an extra -1 in there. Was that intentional? Let's see if we can reproduce it in a test case.

assert(mirror(0, 0, 3) == 0);
assert(mirror(1, 0, 3) == 1);
assert(mirror(2, 0, 3) == 2);
assert(mirror(3, 0, 3) == 2);
assert(mirror(4, 0, 3) == 1);
assert(mirror(5, 0, 3) == 0);
assert(mirror(6, 0, 3) == 1);
assert(mirror(7, 0, 3) == 2);
assert(mirror(8, 0, 3) == 2);

Yep, the two boundaries behave differently! The moral of the story is: Always test your code.

Now that we know how the code behaves (which, honestly, probably isn't how you intended it to behave) — we can modify it safe in the knowledge that we won't introduce bugs. We just have to keep all our test cases passing. (Of course we should write some test cases with negative inputs, too.)

Any time you have code with repeated addition, you should think about whether it can be replaced with multiplication; and any time you have repeated subtraction, you should think about whether it can be replaced with division (which is to say, modulus).

int mirror(int index, int lB, int uB) {
    int n = uB - lB;
    int period = 2*n - 1;  // the pattern repeats with this period
    int mod_p = (index - lB) % period;
    if (mod_p < n) {
        return lB + mod_p;
    } else {
        return uB - 1 - (mod_p - n);
    }
}