diff --git a/internal/demod/gpudemod/gpudemod_windows.go b/internal/demod/gpudemod/gpudemod_windows.go
index ac6d16c..d2b173c 100644
--- a/internal/demod/gpudemod/gpudemod_windows.go
+++ b/internal/demod/gpudemod/gpudemod_windows.go
@@ -10,7 +10,11 @@ package gpudemod
 #include <cuda_runtime.h>
 
 typedef struct { float x; float y; } gpud_float2;
+typedef void* gpud_stream_handle;
 
+typedef int (__stdcall *gpud_stream_create_fn)(gpud_stream_handle* out);
+typedef int (__stdcall *gpud_stream_destroy_fn)(gpud_stream_handle stream);
+typedef int (__stdcall *gpud_stream_sync_fn)(gpud_stream_handle stream);
 typedef int (__stdcall *gpud_upload_fir_taps_fn)(const float* taps, int n);
 typedef int (__stdcall *gpud_launch_freq_shift_fn)(const gpud_float2* in, gpud_float2* out, int n, double phase_inc, double phase_start);
 typedef int (__stdcall *gpud_launch_fm_discrim_fn)(const gpud_float2* in, float* out, int n);
diff --git a/internal/demod/gpudemod/native/exports.cu b/internal/demod/gpudemod/native/exports.cu
index 2114b65..85f0ffc 100644
--- a/internal/demod/gpudemod/native/exports.cu
+++ b/internal/demod/gpudemod/native/exports.cu
@@ -9,6 +9,27 @@
 #define GPUD_CALL
 #endif
 
+typedef void* gpud_stream_handle;
+
+GPUD_API int GPUD_CALL gpud_stream_create(gpud_stream_handle* out) {
+    if (!out) return -1;
+    cudaStream_t stream;
+    cudaError_t err = cudaStreamCreate(&stream);
+    if (err != cudaSuccess) return (int)err;
+    *out = (gpud_stream_handle)stream;
+    return 0;
+}
+
+GPUD_API int GPUD_CALL gpud_stream_destroy(gpud_stream_handle stream) {
+    if (!stream) return 0;
+    return (int)cudaStreamDestroy((cudaStream_t)stream);
+}
+
+GPUD_API int GPUD_CALL gpud_stream_sync(gpud_stream_handle stream) {
+    if (!stream) return (int)cudaDeviceSynchronize();
+    return (int)cudaStreamSynchronize((cudaStream_t)stream);
+}
+
 __global__ void gpud_freq_shift_kernel(
     const float2* __restrict__ in,
     float2* __restrict__ out,