     int my_index = index_array_32[threadIdx.x];
 
     BitonicSort::Array<int> bs( array_32 );
-    my_index = bs.sort( my_index );
+    my_index = bs.sort32( my_index );
 
     index_array_32[threadIdx.x] = my_index;
+}
+
+__global__
+void sort_64( )
+{
+    BitonicSort::Array<int> bs( array_64 );
+    int2 my_indices = make_int2( index_array_64[threadIdx.x],
+                                 index_array_64[threadIdx.x + 32] );
+    bs.sort64( my_indices );
+
+    index_array_64[threadIdx.x +  0] = my_indices.x;
+    index_array_64[threadIdx.x + 32] = my_indices.y;
 }
 
 __global__

    }

    __device__ inline
    int sort( int my_index )
    {
        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;
    }

private:
    __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 );

    sort_32<<<1,32>>>( );
    print_32<<<1,1>>>( );
    cudaDeviceSynchronize( );

    init_64<<<1,1>>>( 0 );
    print_64<<<1,1>>>( );
    sort_64<<<1,32>>>( );
    print_64<<<1,1>>>( );
    cudaDeviceSynchronize( );
    init_64<<<1,1>>>( 1 );
    print_64<<<1,1>>>( );
    sort_64<<<1,32>>>( );
    print_64<<<1,1>>>( );
    cudaDeviceSynchronize( );
    init_64<<<1,1>>>( 2 );
    print_64<<<1,1>>>( );
    sort_64<<<1,32>>>( );
    print_64<<<1,1>>>( );
    cudaDeviceSynchronize( );

    return 0;
}

