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

void
bubble_sort(__global Data *data, int begin, int end)
{
    for (int count=0;count<end;count++) {
        for (int c=end;c>count;c--) {
            if (data[c].index<data[c-1].index)swap(data,c-1,c);
        }
    }
}


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

    int stack[1024];
    int sp = 0;
    int p;

    while (1) {
        while (begin < end) {
            /*
             * if (end-begin <= 50) {
             *     bubble_sort(data,begin,end);
             *     break;
             * }
            */
            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;
    }

}