typedef struct Data {
    int index;
    int ptr;
    int pad[2];
} Data, *DataPtr;

inline void
swap(__global Data *data, int left, int right )
{
    Data tmp    = data[left];
    data[left]  = data[right];
    data[right] = tmp;
}

__kernel void
quick_sort(__constant int *count,
           __global Data *data)
{
    int begin = 0;
    int end = count[0];

    int stack[1024];
    int sp = 0;
    int p;
    while (1) {
        while (begin < end) {
            int where = (begin + end) / 2;
            int pivot = data[where].index;
            data[where].index = data[begin].index;
            int i;
            p = begin;
            for (i=begin+1; i<=end; i++) {
                if (data[i].index < pivot) {
                    p++;
                    swap(data, p, i);
                }
            }
            data[begin].index = data[p].index;
            data[p].index = pivot;

            stack[sp++] = p + 1;
            stack[sp++] = end;
            end = p - 1;
        }
        if (sp == 0) return;
        end = stack[--sp];
        begin = stack[--sp];
        begin = p + 1;
    }
}