Quick sort depth

Accelerated Computing CUDA CUDA Programming and Performance
1

Dear all;
I had got the cuda c code for quick sort, but there is a parameter named “depth”, what was meant by i

t, and what value i should set when i call the function??

global void cdp_simple_quicksort(unsigned int *data, int left, int right, int depth)
{
// If we’re too deep or there are few elements left, we use an insertion sort…
if (depth >= MAX_DEPTH || right - left <= INSERTION_SORT)
{
selection_sort(data, left, right);
return;
}

unsigned int *lptr = data + left;
unsigned int *rptr = data + right;
unsigned int  pivot = data[(left + right) / 2];

// Do the partitioning.
while (lptr <= rptr)
{
	// Find the next left- and right-hand values to swap
	unsigned int lval = *lptr;
	unsigned int rval = *rptr;

	// Move the left pointer as long as the pointed element is smaller than the pivot.
	while (lval < pivot)
	{
		lptr++;
		lval = *lptr;
	}

	// Move the right pointer as long as the pointed element is larger than the pivot.
	while (rval > pivot)
	{
		rptr--;
		rval = *rptr;
	}

	// If the swap points are valid, do the swap!
	if (lptr <= rptr)
	{
		*lptr++ = rval;
		*rptr-- = lval;
	}
}

// Now the recursive part
int nright = rptr - data;
int nleft = lptr - data;

// Launch a new block to sort the left part.
if (left < (rptr - data))
{
	cudaStream_t s;
	cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking);
	cdp_simple_quicksort << < 1, 1, 0, s >> >(data, left, nright, depth + 1);
	cudaStreamDestroy(s);
}

// Launch a new block to sort the right part.
if ((lptr - data) < right)
{
	cudaStream_t s1;
	cudaStreamCreateWithFlags(&s1, cudaStreamNonBlocking);
	cdp_simple_quicksort << < 1, 1, 0, s1 >> >(data, nleft, right, depth + 1);
	cudaStreamDestroy(s1);
}

}

////////////////////////////////////////////////////////////////////////////////

2

Probably it will help if you study the whole file and sample code that the function you have shown comes from:

http://docs.nvidia.com/cuda/cuda-samples/index.html#simple-quicksort–cuda-dynamic-parallelism-

This sorting algorithm calls the function you have shown recursively. CUDA dynamic parallelism can support recursive algorithms by having parent kernels call child kernels. However the process of a parent kernel calling a child kernel requires the storage of some state per call, and this limits the depth of recursion. The depth parameter simply reflects the depth of recursion. When a predefined depth is reached, then recursion is halted, and further sorting is switched to a non-recursive method.

If you study the run_qsort function in the same file that this code comes from, you’ll see how the depth parameter is set initially (it is set to 0):

cdp_simple_quicksort<<< 1, 1 >>>(data, left, right, 0);

Thereafter it is incremented (automatically) at each recursive call by the code in the kernel you have shown:

cdp_simple_quicksort << < 1, 1, 0, s >> >(data, left, nright, depth + 1);
3

Thank you
you mean that i set its value to 0 then automatically its incremented ?

4

Yes