pragma unroll error - Advisory: Loop was not unrolled, unexpected control flow construct

Accelerated Computing CUDA CUDA Programming and Performance
1

I’ve been trying to play around with the pragma unroll construct but I keep getting the following compiler warning:

Advisory: Loop was not unrolled, unexpected control flow construct

Any idea what can be causing it and how to get pragma unroll working?

One of the codes that exhibits the problem if it helps:

[codebox]#pragma unroll 4

	for (int kx = 0 ; kx < seWidth  ; kx++)

	{

		unsigned char v;

		v = tex2D(seTex, kx, ky);

		if (v)

		{

			// Erosion of image by structuring element

			fgResponse.x = fgResponse.x && PATCH(sx     + kx, threadIdx.y + ky);

			fgResponse.y = fgResponse.y && PATCH(sx + 1 + kx, threadIdx.y + ky);

			fgResponse.z = fgResponse.z && PATCH(sx + 2 + kx, threadIdx.y + ky);

			fgResponse.w = fgResponse.w && PATCH(sx + 3 + kx, threadIdx.y + ky);

		}

		v = tex2D(secTex, kx, ky);

		if (v)

		{

			// Erosion of image complement by background structuring element (not that

			// the background structuring element is not the structuring element complement

			// if we want non-square hit-or-miss transform

			bgResponse.x = bgResponse.x && !PATCH(sx     + kx, threadIdx.y + ky);

			bgResponse.y = bgResponse.y && !PATCH(sx + 1 + kx, threadIdx.y + ky);

			bgResponse.z = bgResponse.z && !PATCH(sx + 2 + kx, threadIdx.y + ky);

			bgResponse.w = bgResponse.w && !PATCH(sx + 3 + kx, threadIdx.y + ky);

		}

	}[/codebox]
2

I’m quite certain that it’s the texture fetch that you are doing.

Unfortunately these calls can’t be unrolled if my memory serves me.

So you might have to go manual :)

3

Thanks, did some tests after your suggestion, and it seems that it’s not the texture fetch functions, but rather the logical operations.

For some reason, if I change all the && in the code to * loop unrolling works fine. Unfortunately I also have a goto in the original code

(to break out of two loops at once) which kills loop unrolling as well

4

Aha, so with some changes you manage to unroll despite the texture fetch being there? I haven’t tried it myself but I’ve seen people on this forum complain about it not being possible.

Yeah the “goto” should make things tricky for you :)