extern "C" {
    __global__ void 
    transpose(long* param, float* src, float* dst)
    {
        unsigned long xgid = blockIdx.x*blockDim.x+threadIdx.x; // (unsigned long)s->get_param(0);
        unsigned long ygid = blockIdx.y*blockDim.y*threadIdx.y; // (unsigned long)s->get_param(1);

        long n = param[0];

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

        dst[2*oid] = src[2*iid];
        dst[2*oid+1] = src[2*iid+1];
    }
}