     __device__ inline
     int shiftit( int my_index, int shift, int direction )
     {
-        T   val         = _array[my_index];
-        int other_index = __shfl_xor( my_index, 1 << shift );
-        T   other_val   = _array[other_index];
+        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 );
 
-        const bool my_less = ( val < other_val );
-        bool       reverse = ( threadIdx.x && 2 << direction );
-        const bool id_less = ( threadIdx.x & ( 1 << shift ) == 0 );
-        if( not id_less ) reverse = ! reverse;
-        if( reverse )
-            return ( my_less ? val : other_val );
-        else
-            return ( my_less ? other_val : val );
+        return ( must_swap ? __shfl_xor( my_index, 1 << shift )
+                           : my_index );
     }
 };
 }

    __device__ inline
    int sort( int my_index )
    {
        my_index = shiftit( my_index, 1, 1 );
        return my_index;
    }


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

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

    index_array_32[threadIdx.x] = my_index;
}

__global__
void init_32( int i )
{
    for( int i=0; i<32; i++ ) index_array_32[i] = i;

    if( i == 0 )
        for( int i=0; i<32; i++ ) array_32[i] = 327 * i % 97;
    else if( i == 1 )
        for( int i=0; i<32; i++ ) array_32[i] = 32 - i;
    else
        for( int i=0; i<32; i++ ) array_32[i] = i;
}

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

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