Improve GPU phase precision

internal/demod/gpudemod/kernels.cu

@@ -12,6 +12,8 @@ extern "C" __global__ void gpud_freq_shift_kernel(
     if (idx >= n) return;
 
     double phase = phase_start + phase_inc * (double)idx;
+    const double TWO_PI = 6.283185307179586;
+    phase = phase - rint(phase / TWO_PI) * TWO_PI;
     float si, co;
     sincosf((float)phase, &si, &co);
 
@@ -161,6 +163,8 @@ extern "C" __global__ void gpud_ssb_product_kernel(
     int idx = blockIdx.x * blockDim.x + threadIdx.x;
     if (idx >= n) return;
     double phase = phase_start + phase_inc * (double)idx;
+    const double TWO_PI = 6.283185307179586;
+    phase = phase - rint(phase / TWO_PI) * TWO_PI;
     float si, co;
     sincosf((float)phase, &si, &co);
     float2 v = in[idx];

internal/demod/gpudemod/native/exports.cu

@@ -41,6 +41,12 @@ __global__ void gpud_freq_shift_kernel(
     if (idx >= n) return;
 
     double phase = phase_start + phase_inc * (double)idx;
+    // Reduce phase to [-pi, pi) BEFORE float cast to preserve precision.
+    // Without this, phase accumulates to millions of radians and the
+    // (float) cast loses ~0.03-0.1 rad, causing audible clicks at
+    // frame boundaries in streaming audio.
+    const double TWO_PI = 6.283185307179586;
+    phase = phase - rint(phase / TWO_PI) * TWO_PI;
     float si, co;
     sincosf((float)phase, &si, &co);
 
@@ -293,6 +299,8 @@ __global__ void gpud_ssb_product_kernel(
     int idx = blockIdx.x * blockDim.x + threadIdx.x;
     if (idx >= n) return;
     double phase = phase_start + phase_inc * (double)idx;
+    const double TWO_PI = 6.283185307179586;
+    phase = phase - rint(phase / TWO_PI) * TWO_PI;
     float si, co;
     sincosf((float)phase, &si, &co);
     float2 v = in[idx];