# Quick sort depth

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);
}
``````

}

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

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);
``````

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

Yes