diff --git a/components/rgbd-sources/src/sources/snapshot/snapshot_source.cpp b/components/rgbd-sources/src/sources/snapshot/snapshot_source.cpp
index 136a2e7dfcc7b279228cdd5c6efc0bfb8d303baa..61acdb9d814560c72fcd15e7f05cdfaf8fda66cd 100644
--- a/components/rgbd-sources/src/sources/snapshot/snapshot_source.cpp
+++ b/components/rgbd-sources/src/sources/snapshot/snapshot_source.cpp
@@ -53,6 +53,8 @@ SnapshotSource::SnapshotSource(ftl::rgbd::Source *host, Snapshot &snapshot, cons
     host->setPose(pose);
 
 	mspf_ = 1000 / host_->value("fps", 20);
+
+	cudaStreamCreate(&stream_);
 }
 
 bool SnapshotSource::compute(int n, int b) {
@@ -61,11 +63,14 @@ bool SnapshotSource::compute(int n, int b) {
 
 	//snap_rgb_.copyTo(rgb_);
 	//snap_depth_.copyTo(depth_);
-	rgb_.upload(snap_rgb_);
-	depth_.upload(snap_depth_);
-
-	auto cb = host_->callback();
-	if (cb) cb(timestamp_, rgb_, depth_);
+	cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream_);
+	rgb_.upload(snap_rgb_, cvstream);
+	depth_.upload(snap_depth_, cvstream);
+	cudaStreamSynchronize(stream_);
+
+	//auto cb = host_->callback();
+	//if (cb) cb(timestamp_, rgb_, depth_);
+	host_->notify(timestamp_, rgb_, depth_);
 
 	frame_idx_ = (frame_idx_ + 1) % snapshot_.getFramesCount();
 
diff --git a/components/rgbd-sources/src/sources/snapshot/snapshot_source.hpp b/components/rgbd-sources/src/sources/snapshot/snapshot_source.hpp
index de1b0df48be79df732f51144226f5c7e6d2f0478..80a0bf392b39fb9d5215dd80034768d806ac7957 100644
--- a/components/rgbd-sources/src/sources/snapshot/snapshot_source.hpp
+++ b/components/rgbd-sources/src/sources/snapshot/snapshot_source.hpp
@@ -32,6 +32,7 @@ class SnapshotSource : public detail::Source {
 	cv::Mat snap_rgb_;
 	cv::Mat snap_depth_;
 	int mspf_;
+	cudaStream_t stream_;
 };
 
 }