typedef int magma_int_t;
typedef double2 magmaDoubleComplex;

#define MAX_PIVOTS 32
typedef struct {
    int npivots;
    int ipiv[MAX_PIVOTS];
} zlaswp_params_t;

__kernel void zlaswp_kernel(
    magma_int_t n,
    __global magmaDoubleComplex *dAT, unsigned long dAT_offset, magma_int_t ldda,
    zlaswp_params_t params )
{
    dAT += dAT_offset;

    int tid = get_local_id(0) + get_local_size(0)*get_group_id(0);
    if ( tid < n ) {
        dAT += tid;
        __global magmaDoubleComplex *A1  = dAT;
        
        for( int i1 = 0; i1 < params.npivots; ++i1 ) {
            int i2 = params.ipiv[i1];
            __global magmaDoubleComplex *A2 = dAT + i2*ldda;
            magmaDoubleComplex temp = *A1;
            *A1 = *A2;
            *A2 = temp;
            A1 += ldda;  // A1 = dA + i1*ldx
        }
    }
}