diff --git a/internal/demod/gpudemod/kernels.cu b/internal/demod/gpudemod/kernels.cu index f308b8d..8250792 100644 --- a/internal/demod/gpudemod/kernels.cu +++ b/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]; diff --git a/internal/demod/gpudemod/native/exports.cu b/internal/demod/gpudemod/native/exports.cu index 97118c8..6081b57 100644 --- a/internal/demod/gpudemod/native/exports.cu +++ b/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];