#pragma unroll not working?

Accelerated Computing CUDA CUDA Programming and Performance
1

Consider this very sophisticated and advanced sorting algorithm for 32 integer numbers :)

#define NUM 32

#define NUMTESTS 65536

__global__ static void bubbleSort(int * input, int *output) {

	__shared__ int shared[NUM+1];

	shared[NUM]=999999;

	const unsigned int tid = threadIdx.x;

	for (int test=0; test<NUMTESTS; ++test) {

		__syncthreads();

		shared[tid] = input[tid];

		__threadfence_block();

#pragma unroll 32

		for (int i=0; i<32; ++i) {

			if ((threadIdx.x&1) == (i&1)) {

				if (shared[tid+1]<shared[tid])

					swap(shared[tid+1],shared[tid]);

			}

			__threadfence_block();

		}

		output[tid] = shared[tid];

	}

}

I get the following running time:

Done! GPU execution time: 585.410156 ms

Now consider another code - I simply unroll the for loop manually, 32 times. I am not going to paste the code here for obvious reasons :)

Done! GPU execution time: 388.730560 ms

So my question is: is #pragma unroll really working?

Some answers you might want to hear:

    [*]Times are repetitive. It is not one of those strange nasty execution-time spikes.

    [*]I am using nvcc -O2 -arch sm_13 (although programming guide does not state if and when unrolling may not work)

    [*]Tested on GTX260, Win XP 32bit, Cuda 2.2

    [*]I call it in <<<1,32>>> configuration.

    [*]I launch the timer just before kernel call and I do use cudaThreadSynchronize() before stopping it.

    [*]Algorithm gives correct results in both cases

    [*]Bitonic sort in SDK version is slower for arrays of size 32 (600ms), but if you manually unroll all its for-s and squeeze if-else into if (it can be done), you can actually be faster. (In my case I have it in 306ms)

2

The easiest way to check is to call nvcc with the --ptx flag. Then examine the generated PTX file to see if unrolling has occurred.

3

You can also check the cubin file with decuda, then you really know what is going on.

I have seen cases where a loop didn’t get unrolled because the iterator variable ‘i’ was replaced by the compiler to a variable which was used inside the loop. No matter how hard I try to use the #pragma, the loop didn’t get unrolled. However, on the other hand, I did see loop unrolling on loops where I didn’t use the #pragma statement at all, so loop unrolling seems to be out of the programmers control.

4

It’s not entirely out of your control. The programming guide says that the default behavior of the compiler is to unroll short loops when possible. You can disable unrolling for a particular loop with #pragma unroll 1. There are some loops which the compiler cannot unroll for some reason, which is a compiler limitation, and also some loops which are not unrolled even when given an explicit trip count, which I believe is a bug.