From 65cca851784170cdba5effc13852b480cc4f49a8 Mon Sep 17 00:00:00 2001 From: Carles Fernandez Date: Mon, 21 Sep 2026 17:18:55 +0200 Subject: [PATCH] Fix CUDA acquisition upload ordering and validation MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit - Order code and carrier uploads on the compute stream and wait before returning so callers can safely reuse host buffers. - Verify CUDA readiness and completed grid counts in adapter tests, rejecting unexpected CPU fallback. - Exercise host-buffer reuse in engine parity tests. - Correct the Jetson helper's test executable path. - Use ArgsProduct for compatibility with Google Benchmark 1.9.2–1.9.5. --- .../gnuradio_blocks/pcps_acquisition.cc | 1 + .../gnuradio_blocks/pcps_acquisition.h | 22 ++++++++++++++ .../acquisition/libs/cuda_pcps_engine.cu | 30 +++++++++++++++---- .../acquisition/libs/cuda_pcps_engine.h | 2 ++ tests/benchmarks/benchmark_pcps_grid.cc | 17 ++++------- .../acquisition/cuda_pcps_engine_test.cc | 12 ++++++-- .../gps_l1_ca_pcps_acquisition_cuda_test.cc | 20 +++++++++++++ utils/scripts/jetson-build.sh | 2 +- 8 files changed, 85 insertions(+), 21 deletions(-) diff --git a/src/algorithms/acquisition/gnuradio_blocks/pcps_acquisition.cc b/src/algorithms/acquisition/gnuradio_blocks/pcps_acquisition.cc index 8ed5df195..e6903d80b 100644 --- a/src/algorithms/acquisition/gnuradio_blocks/pcps_acquisition.cc +++ b/src/algorithms/acquisition/gnuradio_blocks/pcps_acquisition.cc @@ -745,6 +745,7 @@ void pcps_acquisition::doppler_grid(const gr_complex* in) { if (doppler_grid_cuda(in)) { + ++d_cuda_grid_count; return; } LOG(WARNING) << "CUDA acquisition failed in channel " << d_channel << " (" << d_cuda_engine->last_error() diff --git a/src/algorithms/acquisition/gnuradio_blocks/pcps_acquisition.h b/src/algorithms/acquisition/gnuradio_blocks/pcps_acquisition.h index e6009ce5a..344e946f8 100644 --- a/src/algorithms/acquisition/gnuradio_blocks/pcps_acquisition.h +++ b/src/algorithms/acquisition/gnuradio_blocks/pcps_acquisition.h @@ -165,6 +165,27 @@ public: */ void set_doppler_uncertainty(uint32_t doppler_uncertainty); + //! Whether the CUDA engine is available. Inspect only while acquisition is stopped. + inline bool cuda_ready() const + { +#if CUDA_GPU_ACCEL + return d_cuda_engine != nullptr; +#else + return false; +#endif + } + + //! Completed CUDA grids over the block's lifetime, excluding warm-up. + //! Inspect only while acquisition is stopped. + inline uint64_t cuda_grid_count() const + { +#if CUDA_GPU_ACCEL + return d_cuda_grid_count; +#else + return 0; +#endif + } + /*! * \brief Parallel Code Phase Search Acquisition signal processing. */ @@ -297,6 +318,7 @@ private: std::unique_ptr d_fft_if; #if CUDA_GPU_ACCEL std::unique_ptr d_cuda_engine; // null => CPU path + uint64_t d_cuda_grid_count{0}; #endif }; diff --git a/src/algorithms/acquisition/libs/cuda_pcps_engine.cu b/src/algorithms/acquisition/libs/cuda_pcps_engine.cu index 37daffabf..10beb114b 100644 --- a/src/algorithms/acquisition/libs/cuda_pcps_engine.cu +++ b/src/algorithms/acquisition/libs/cuda_pcps_engine.cu @@ -403,11 +403,17 @@ bool CudaPcpsEngine::set_fft_codes(const std::complex* fft_codes) return false; } cudaSetDevice(p->device); - // Synchronous copy: the caller may reuse its buffer immediately. - const cudaError_t e = cudaMemcpy(p->d_codes, fft_codes, p->fft_size * sizeof(float2), cudaMemcpyHostToDevice); + // Order uploads with kernels on the nonblocking stream, then wait so the + // caller may reuse its host buffer immediately. + cudaError_t e = cudaMemcpyAsync(p->d_codes, fft_codes, p->fft_size * sizeof(float2), cudaMemcpyHostToDevice, p->stream); if (e != cudaSuccess) { - return p->fail("cudaMemcpy(codes)", e); + return p->fail("cudaMemcpyAsync(codes)", e); + } + e = cudaStreamSynchronize(p->stream); + if (e != cudaSuccess) + { + return p->fail("cudaStreamSynchronize(codes)", e); } return true; } @@ -436,14 +442,26 @@ bool CudaPcpsEngine::set_doppler_wipeoffs(GridId grid, const std::complex // only done when the Doppler grid changes (PRN change on FDMA, assisted // Doppler center, or the two-step refinement), not per dwell. const size_t row_bytes = p->fft_size * sizeof(float2); + cudaError_t upload_error = cudaSuccess; for (uint32_t k = 0; k < bins; k++) { - const cudaError_t e = cudaMemcpy(p->d_wipe[grid] + static_cast(k) * p->fft_size, wipeoffs[k], row_bytes, cudaMemcpyHostToDevice); - if (e != cudaSuccess) + upload_error = cudaMemcpyAsync(p->d_wipe[grid] + static_cast(k) * p->fft_size, wipeoffs[k], row_bytes, cudaMemcpyHostToDevice, p->stream); + if (upload_error != cudaSuccess) { - return p->fail("cudaMemcpy(wipeoffs)", e); + break; } } + // Drain queued uploads even if a later row failed, before the caller can + // reuse its host buffers or plan creation can return an error. + const cudaError_t sync_error = cudaStreamSynchronize(p->stream); + if (upload_error != cudaSuccess) + { + return p->fail("cudaMemcpyAsync(wipeoffs)", upload_error); + } + if (sync_error != cudaSuccess) + { + return p->fail("cudaStreamSynchronize(wipeoffs)", sync_error); + } p->wipe_bins[grid] = bins; // Build the cuFFT plan and run the whole pipeline once on the current diff --git a/src/algorithms/acquisition/libs/cuda_pcps_engine.h b/src/algorithms/acquisition/libs/cuda_pcps_engine.h index 103baa8f4..1b0d5346b 100644 --- a/src/algorithms/acquisition/libs/cuda_pcps_engine.h +++ b/src/algorithms/acquisition/libs/cuda_pcps_engine.h @@ -86,11 +86,13 @@ public: /*! * \brief Upload conj(FFT(local code)), fft_size elements. + * The upload completes before returning; the caller may reuse fft_codes. */ bool set_fft_codes(const std::complex* fft_codes); /*! * \brief Upload the Doppler wipe-off carriers for one grid. + * Uploads complete before returning; the caller may reuse the host rows. * \param grid MAIN_GRID or STEP2_GRID. * \param wipeoffs Array of `bins` host pointers, each pointing at fft_size complex samples. * \param bins Number of Doppler bins (<= max_doppler_bins). diff --git a/tests/benchmarks/benchmark_pcps_grid.cc b/tests/benchmarks/benchmark_pcps_grid.cc index d95e2efa9..cdcb4104a 100644 --- a/tests/benchmarks/benchmark_pcps_grid.cc +++ b/tests/benchmarks/benchmark_pcps_grid.cc @@ -169,20 +169,13 @@ void bm_pcps_grid_cuda(benchmark::State& state) // fft_size x bins. 1 ms at 2/4/8/16/20 Msps, then 4 ms at 4 Msps (16000) and // 2 ms at 20 Msps (40000). Bins: +/-5 kHz at 500/250/125 Hz steps. -static void grid_args(benchmark::internal::Benchmark* b) -{ - for (int fft_size : {2000, 4000, 8000, 16000, 20000, 40000}) - { - for (int bins : {21, 41, 81}) - { - b->Args({fft_size, bins}); - } - } -} +const std::vector> grid_args{ + {2000, 4000, 8000, 16000, 20000, 40000}, + {21, 41, 81}}; -BENCHMARK(bm_pcps_grid_cpu)->Apply(grid_args)->Unit(benchmark::kMicrosecond)->UseRealTime(); +BENCHMARK(bm_pcps_grid_cpu)->ArgsProduct(grid_args)->Unit(benchmark::kMicrosecond)->UseRealTime(); #if CUDA_GPU_ACCEL -BENCHMARK(bm_pcps_grid_cuda)->Apply(grid_args)->Unit(benchmark::kMicrosecond)->UseRealTime(); +BENCHMARK(bm_pcps_grid_cuda)->ArgsProduct(grid_args)->Unit(benchmark::kMicrosecond)->UseRealTime(); #endif BENCHMARK_MAIN(); diff --git a/tests/unit-tests/signal-processing-blocks/acquisition/cuda_pcps_engine_test.cc b/tests/unit-tests/signal-processing-blocks/acquisition/cuda_pcps_engine_test.cc index c885580ba..19fc2bfef 100644 --- a/tests/unit-tests/signal-processing-blocks/acquisition/cuda_pcps_engine_test.cc +++ b/tests/unit-tests/signal-processing-blocks/acquisition/cuda_pcps_engine_test.cc @@ -218,13 +218,21 @@ void run_parity_case(const PcpsScenario& sc, uint32_t dwells) CudaPcpsEngine gpu(N, E, bins); ASSERT_TRUE(gpu.is_valid()) << gpu.last_error(); + auto wipeoffs = cpu.wipeoffs(); + auto fft_codes = cpu.fft_codes(); std::vector*> wipe_rows(bins); for (uint32_t k = 0; k < bins; k++) { - wipe_rows[k] = cpu.wipeoffs()[k].data(); + wipe_rows[k] = wipeoffs[k].data(); } ASSERT_TRUE(gpu.set_doppler_wipeoffs(CudaPcpsEngine::MAIN_GRID, wipe_rows.data(), bins)) << gpu.last_error(); - ASSERT_TRUE(gpu.set_fft_codes(cpu.fft_codes().data())) << gpu.last_error(); + // Uploads must finish before returning so host buffers can be reused. + for (auto& row : wipeoffs) + { + std::fill(row.begin(), row.end(), std::complex(0.0F, 0.0F)); + } + ASSERT_TRUE(gpu.set_fft_codes(fft_codes.data())) << gpu.last_error(); + std::fill(fft_codes.begin(), fft_codes.end(), std::complex(0.0F, 0.0F)); std::vector gpu_grid(bins, fvec(E)); std::vector out_rows(bins); diff --git a/tests/unit-tests/signal-processing-blocks/acquisition/gps_l1_ca_pcps_acquisition_cuda_test.cc b/tests/unit-tests/signal-processing-blocks/acquisition/gps_l1_ca_pcps_acquisition_cuda_test.cc index 127f05206..2d47033df 100644 --- a/tests/unit-tests/signal-processing-blocks/acquisition/gps_l1_ca_pcps_acquisition_cuda_test.cc +++ b/tests/unit-tests/signal-processing-blocks/acquisition/gps_l1_ca_pcps_acquisition_cuda_test.cc @@ -29,6 +29,7 @@ #include #include #include +#include #include #include #include @@ -120,6 +121,7 @@ protected: }; std::shared_ptr make_config(bool use_cuda, bool two_steps) const; + void check_backend(PcpsAcquisitionAdapter &acquisition, bool use_cuda, uint64_t expected_cuda_grids) const; Outcome run_once(bool use_cuda, bool two_steps = false); unsigned int doppler_max{5000}; @@ -134,6 +136,8 @@ std::shared_ptr GpsL1CaPcpsAcquisitionCudaTest::make_conf config->set_property("Acquisition_1C.implementation", "GPS_L1_CA_PCPS_Acquisition"); config->set_property("Acquisition_1C.item_type", "gr_complex"); config->set_property("Acquisition_1C.coherent_integration_time_ms", "1"); + config->set_property("Acquisition_1C.max_dwells", "1"); + config->set_property("Acquisition_1C.blocking", "true"); config->set_property("Acquisition_1C.dump", "false"); config->set_property("Acquisition_1C.threshold", "0.001"); config->set_property("Acquisition_1C.doppler_max", std::to_string(doppler_max)); @@ -150,8 +154,19 @@ std::shared_ptr GpsL1CaPcpsAcquisitionCudaTest::make_conf } +void GpsL1CaPcpsAcquisitionCudaTest::check_backend(PcpsAcquisitionAdapter &acquisition, bool use_cuda, uint64_t expected_cuda_grids) const +{ + const auto block = acquisition.get_right_block(); + const auto *pcps = dynamic_cast(block.get()); + ASSERT_NE(nullptr, pcps); + EXPECT_EQ(use_cuda, pcps->cuda_ready()) << "Unexpected acquisition backend (CUDA initialization or CPU fallback)."; + EXPECT_EQ(expected_cuda_grids, pcps->cuda_grid_count()) << "Acquisition did not execute the expected number of CUDA grids."; +} + + GpsL1CaPcpsAcquisitionCudaTest::Outcome GpsL1CaPcpsAcquisitionCudaTest::run_once(bool use_cuda, bool two_steps) { + SCOPED_TRACE(::testing::Message() << "use_cuda=" << use_cuda << ", two_steps=" << two_steps); Outcome out; auto config = make_config(use_cuda, two_steps); auto top_block = gr::make_top_block("Acquisition CUDA test"); @@ -178,12 +193,16 @@ GpsL1CaPcpsAcquisitionCudaTest::Outcome GpsL1CaPcpsAcquisitionCudaTest::run_once top_block->msg_connect(acquisition->get_right_block(), pmt::mp("events"), msg_rx, pmt::mp("events")); acquisition->set_local_code(); + check_backend(*acquisition, use_cuda, 0); acquisition->reset(); const auto start = std::chrono::steady_clock::now(); top_block->run(); const auto end = std::chrono::steady_clock::now(); + acquisition->stop_acquisition(); + check_backend(*acquisition, use_cuda, use_cuda ? (two_steps ? 2U : 1U) : 0U); + out.message = msg_rx->rx_message; out.doppler_hz = gnss_synchro.Acq_doppler_hz; out.delay_samples = gnss_synchro.Acq_delay_samples; @@ -196,6 +215,7 @@ TEST_F(GpsL1CaPcpsAcquisitionCudaTest /*unused*/, Instantiate /*unused*/) { auto config = make_config(true, false); auto acquisition = std::make_shared(config.get(), "Acquisition_1C", "GPS_L1_CA_PCPS_Acquisition", 1, 0, GPS_1C); + check_backend(*acquisition, true, 0); } diff --git a/utils/scripts/jetson-build.sh b/utils/scripts/jetson-build.sh index b83a7edb0..1769b02d6 100755 --- a/utils/scripts/jetson-build.sh +++ b/utils/scripts/jetson-build.sh @@ -77,7 +77,7 @@ cmake --build "$BUILD_DIR" -j"$JOBS" if [[ $DO_TESTS -eq 1 ]]; then echo "== Running CUDA unit tests" - "$BUILD_DIR/src/tests/run_tests" \ + "$BUILD_DIR/tests/run_tests" \ --gtest_filter='CudaPcpsEngineTest.*:GpsL1CaPcpsAcquisitionCudaTest.*:GpuMulticorrelatorTest.*' fi