     __device__ inline
     int sort( int my_index )
     {
-        my_index = shiftit( my_index, 1, 1 );
+        for( int outer=0; outer<5; outer++ ) {
+            for( int inner=outer; inner>=0; inner-- ) {
+                my_index = shiftit( my_index, inner, outer+1 );
+            }
+        }
         return my_index;
     }
 

__global__
void sort_32( )
{
    int my_index = threadIdx.x;

    BitonicSort::Array<int> bs( array_32 );
    my_index = bs.sort( my_index );
}

__global__
void init_32( )
{
    for( int i=0; i<32; i++ ) {
        array_32[i] = 327 * i % 17;
        index_array_32[i] = i;
    }
}

__global__
void print_32( )
{
    printf("Hello\n");
    for( int i=0; i<32; i++ ) {
        printf( "%d ", array_32[i] );
    }
    printf("\n");
}

int main( )
{
    init_32<<<1,1>>>( );
    print_32<<<1,1>>>( );
    sort_32<<<1,32>>>( );
    print_32<<<1,1>>>( );

    __device__ inline
    int shiftit( int my_index, int shift, int direction )
    {
        const T    my_val      = _array[my_index];
        const T    other_val   = __shfl_xor( my_val, 1 << shift );
        const bool my_less     = ( my_val < other_val );
        const bool reverse     = ( threadIdx.x & ( 1 << direction ) );
        const bool id_less     = ( ( threadIdx.x & ( 1 << shift ) ) == 0 );
        const bool must_swap   = not ( my_less ^ id_less ^ reverse );

        return ( must_swap ? __shfl_xor( my_index, 1 << shift )
                           : my_index );
    }
};
}
