extern "C" {
    __global__ void
    norm(long* param, float* x)
    {
        unsigned long gid = blockIdx.x*blockDim.x+threadIdx.x; // (unsigned long)s->get_param(0);
        unsigned long nid = blockIdx.y*blockDim.y+threadIdx.y; //(unsigned long)s->get_param(1);
        
        long n = param[0];
        
        x[(nid*n+gid)*2] = x[(nid*n+gid)*2] / (float)n;
        x[(nid*n+gid)*2+1] = x[(nid*n+gid)*2+1] / (float)n;
    }
}