mirrors/gnss-sdr
1
0
Fork 0
mirror of https://github.com/gnss-sdr/gnss-sdr synced 2026-10-10 01:31:49 +00:00
Code Issues Releases Wiki Activity

Fix CUDA acquisition upload ordering and validation

Browse Source 
- 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.
This commit is contained in:
Carles Fernandez
2026-09-21 17:18:55 +02:00
parent a996687488
commit 65cca85178
8 changed files with 85 additions and 21 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/algorithms/acquisition/gnuradio_blocks/pcps_acquisition.cc
+1
View File
@@ -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()
src/algorithms/acquisition/gnuradio_blocks/pcps_acquisition.h
+22
View File
@@ -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<gnss_fft_complex_fwd> d_fft_if;
#if CUDA_GPU_ACCEL
std::unique_ptr<CudaPcpsEngine> d_cuda_engine; // null => CPU path
uint64_t d_cuda_grid_count{0};
#endif
};
src/algorithms/acquisition/libs/cuda_pcps_engine.cu
+24 -6
View File
@@ -403,11 +403,17 @@ bool CudaPcpsEngine::set_fft_codes(const std::complex<float>* 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<float>
// 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<size_t>(k) * p->fft_size, wipeoffs[k], row_bytes, cudaMemcpyHostToDevice);
if (e != cudaSuccess)
upload_error = cudaMemcpyAsync(p->d_wipe[grid] + static_cast<size_t>(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
src/algorithms/acquisition/libs/cuda_pcps_engine.h
+2
View File
@@ -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<float>* 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).
tests/benchmarks/benchmark_pcps_grid.cc
+5 -12
View File
@@ -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<std::vector<int64_t>> 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();
tests/unit-tests/signal-processing-blocks/acquisition/cuda_pcps_engine_test.cc
+10 -2
View File
@@ -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<const std::complex<float>*> 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<float>(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<float>(0.0F, 0.0F));
std::vector<fvec> gpu_grid(bins, fvec(E));
std::vector<float*> out_rows(bins);
tests/unit-tests/signal-processing-blocks/acquisition/gps_l1_ca_pcps_acquisition_cuda_test.cc
+20
View File
@@ -29,6 +29,7 @@
#include <pmt/pmt.h>
#include <chrono>
#include <cmath>
#include <cstdint>
#include <iostream>
#include <memory>
#include <string>
@@ -120,6 +121,7 @@ protected:
};
std::shared_ptr<InMemoryConfiguration> 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<InMemoryConfiguration> 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<InMemoryConfiguration> 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<const pcps_acquisition *>(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<PcpsAcquisitionAdapter>(config.get(), "Acquisition_1C", "GPS_L1_CA_PCPS_Acquisition", 1, 0, GPS_1C);
check_backend(*acquisition, true, 0);
}
utils/scripts/jetson-build.sh
+1 -1
View File
@@ -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