__kernel void
norm(__constant long *param, __global float2 *in_x, __global float2 *out_x)
{
    unsigned long gid = (unsigned long)get_global_id(0);
    unsigned long nid = (unsigned long)get_global_id(1);

    long n = param[3];

    out_x[nid*n+gid].x = in_x[nid*n+gid].x / (float)n;
    out_x[nid*n+gid].y = in_x[nid*n+gid].y / (float)n;
}