__kernel
void transpose(__constant long *param, __global float2 *src, __global float2 *dst)
{
    unsigned long xgid = (unsigned long)get_global_id(0);
    unsigned long ygid = (unsigned long)get_global_id(1);

    long n = (long)param[3];

    unsigned int iid = ygid * n + xgid;
    unsigned int oid = xgid * n + ygid;

    dst[oid] = src[iid];
}