                 qx = (j.x+(float)(Nx*ix));
                 kn1 = b1 * qx;
                 argx = Scalar(0.5)*qx*kH.x;
-                if (argx != 0.0f)
-                    wx = powf(sinf(argx)/argx,order);
-                else
-                    wx = Scalar(1.0);
+                wx = powf(gpu_sinc(argx),order);
 
                for (iy = -nby; iy <= nby; iy++) {
                     qy = (j.y+(float)(Ny*iy));
                     kn2 = b2 * qy;
                     argy = Scalar(0.5)*qy*kH.y;
-                    if (argy != 0.0f)
-                        wy = powf(sinf(argy)/argy,order);
-                    else
-                        wy = Scalar(1.0);
+                    wy = powf(gpu_sinc(argy),order);
 
                     for (iz = -nbz; iz <= nbz; iz++) {
                         qz = (j.z+(float)(Nz*iz));

    Scalar3 b2 = Scalar(2.0*M_PI)*make_scalar3(a3.y*a1.z-a3.z*a1.y, a3.z*a1.x-a3.x*a1.z, a3.x*a1.y-a3.y*a1.x)/V_box;
    Scalar3 b3 = Scalar(2.0*M_PI)*make_scalar3(a1.y*a2.z-a1.z*a2.y, a1.z*a2.x-a1.x*a2.z, a1.x*a2.y-a1.y*a2.x)/V_box;

    dim3 grid( (int)ceil((double)Nx*Ny*Nz / (double)block_size), 1, 1);
    dim3 threads(block_size, 1, 1);
    reset_kvec_green_hat_kernel <<< grid, threads >>> (box, b1, b2, b3, Nx, Ny, Nz, order, kappa, kvec, green_hat, vg, nbx, nby, nbz, gf_b);

        kvec_array[idx] = k;

        float sqk = dot(k,k);
        // omit DC term
        if(idx == 0) {
            vg[0+6*idx] = 0.0f;
            vg[1+6*idx] = 0.0f;
            vg[2+6*idx] = 0.0f;
