diff --git a/CMakeLists.txt b/CMakeLists.txt index 1b6fc1068..fe04b6cbb 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -2758,14 +2758,61 @@ if(DEFINED ENV{CUDA_GPU_ACCEL}) endif() if(ENABLE_CUDA) - set(CMAKE_CUDA_STANDARD 14) + # Match the CUDA language standard to the host C++ standard so that headers + # shared between .cc and .cu files see the same language level. nvcc >= 11 + # supports C++17; older toolkits fall back to C++14. + if(CMAKE_CXX_STANDARD AND NOT (CMAKE_CXX_STANDARD LESS 17)) + set(CMAKE_CUDA_STANDARD 17) + else() + set(CMAKE_CUDA_STANDARD 14) + endif() + set(CMAKE_CUDA_STANDARD_REQUIRED ON) set(CMAKE_CUDA_EXTENSIONS ON) + # Select the target GPU architecture(s) if the user has not done so. + # CMake >= 3.24 can query the local device ("native"). On NVIDIA Jetson + # (L4T) systems CMake may be older, so we map the SoC to its SM version: + # Orin (AGX/NX/Nano) -> sm_87, Xavier -> sm_72, TX2 -> sm_62, Nano -> sm_53. + if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES) + if(EXISTS "/etc/nv_tegra_release" OR EXISTS "/proc/device-tree/compatible") + set(_tegra_compat "") + if(EXISTS "/proc/device-tree/compatible") + # The device-tree property is a list of NUL-separated strings; + # file(STRINGS) splits on the NULs, file(READ) would stop at the first. + file(STRINGS "/proc/device-tree/compatible" _tegra_compat) + endif() + if("${_tegra_compat}" MATCHES "tegra234") + set(CMAKE_CUDA_ARCHITECTURES 87) + elseif("${_tegra_compat}" MATCHES "tegra194") + set(CMAKE_CUDA_ARCHITECTURES 72) + elseif("${_tegra_compat}" MATCHES "tegra186") + set(CMAKE_CUDA_ARCHITECTURES 62) + elseif("${_tegra_compat}" MATCHES "tegra210") + set(CMAKE_CUDA_ARCHITECTURES 53) + elseif(NOT (CMAKE_VERSION VERSION_LESS 3.24)) + set(CMAKE_CUDA_ARCHITECTURES native) + endif() + elseif(NOT (CMAKE_VERSION VERSION_LESS 3.24)) + set(CMAKE_CUDA_ARCHITECTURES native) + endif() + endif() if(CMAKE_VERSION VERSION_GREATER 3.11) include(CheckLanguage) check_language(CUDA) if(CMAKE_CUDA_COMPILER) enable_language(CUDA) set(CUDA_FOUND TRUE) + # Imported targets CUDA::cudart, CUDA::cufft, ... (CMake >= 3.17) + if(NOT (CMAKE_VERSION VERSION_LESS 3.17)) + find_package(CUDAToolkit REQUIRED) + set_package_properties(CUDAToolkit PROPERTIES + URL "https://developer.nvidia.com/cuda-downloads" + DESCRIPTION "Library for parallel programming in Nvidia GPUs (found: v${CUDAToolkit_VERSION})" + PURPOSE "Used in some processing block implementations (GPU acquisition and tracking)." + TYPE REQUIRED + ) + else() + message(FATAL_ERROR "ENABLE_CUDA=ON requires CMake >= 3.17 (found ${CMAKE_VERSION}).") + endif() else() set(ENABLE_CUDA OFF) endif() @@ -2791,7 +2838,13 @@ if(ENABLE_CUDA) endif() endif() if(ENABLE_CUDA) + # Define CUDA_GPU_ACCEL for every translation unit: pcps_acquisition.h changes + # its class layout under this macro, so it has to be consistent project-wide. + add_compile_definitions(CUDA_GPU_ACCEL=1) message(STATUS "NVIDIA CUDA GPU Acceleration will be enabled.") + if(CMAKE_CUDA_ARCHITECTURES) + message(STATUS " Target CUDA architecture(s): ${CMAKE_CUDA_ARCHITECTURES}") + endif() message(STATUS " You can disable it with 'cmake -DENABLE_CUDA=OFF ..'") else() message(STATUS "NVIDIA CUDA GPU Acceleration will be not enabled.") @@ -3318,9 +3371,8 @@ if((CMAKE_CXX_COMPILER_ID STREQUAL "GNU") AND NOT WIN32) add_compile_options(-Wno-missing-field-initializers) endif() if(CMAKE_CROSSCOMPILING OR NOT ENABLE_PACKAGING) - if(NOT ENABLE_CUDA) - add_compile_options(-Wno-psabi) - endif() + # Restrict to C++ so the flag is never handed to nvcc when ENABLE_CUDA=ON + add_compile_options("$<$:-Wno-psabi>") endif() if(IS_ARM) if((CMAKE_CXX_COMPILER_VERSION VERSION_GREATER "7.1.0") AND (CMAKE_VERSION VERSION_GREATER "3.1")) @@ -3424,7 +3476,7 @@ add_feature_info(ENABLE_GPROF ENABLE_GPROF "Enables performance analysis with 'g add_feature_info(ENABLE_CLANG_TIDY ENABLE_CLANG_TIDY "Runs clang-tidy along with the compiler. Requires Clang.") add_feature_info(ENABLE_PROFILING ENABLE_PROFILING "Runs volk_gnsssdr_profile at the end of the building.") add_feature_info(ENABLE_OPENCL ENABLE_OPENCL "Enables GPS_L1_CA_PCPS_OpenCl_Acquisition (experimental). Requires OpenCL.") -add_feature_info(ENABLE_CUDA ENABLE_CUDA "Enables GPS_L1_CA_DLL_PLL_Tracking_GPU (experimental). Requires CUDA.") +add_feature_info(ENABLE_CUDA ENABLE_CUDA "Enables GPS_L1_CA_DLL_PLL_Tracking_GPU and the CUDA acquisition engine (Acquisition_XX.use_cuda=true). Requires CUDA.") add_feature_info(ENABLE_ARMA_NO_DEBUG ENABLE_ARMA_NO_DEBUG "Enables passing the ARMA_NO_DEBUG macro to Armadillo, hence disabling bound checking.") add_feature_info(ENABLE_PACKAGING ENABLE_PACKAGING "Enables software packaging.") add_feature_info(ENABLE_OWN_GLOG ENABLE_OWN_GLOG "Forces the downloading and building of Google glog.") diff --git a/README.md b/README.md index f2ef42a65..032dddc07 100644 --- a/README.md +++ b/README.md @@ -721,7 +721,18 @@ $ sudo cmake --install build ``` Of course, you will also need a GPU that -[supports CUDA](https://developer.nvidia.com/cuda-gpus "CUDA GPUs"). +[supports CUDA](https://developer.nvidia.com/cuda-gpus "CUDA GPUs"). The target +architecture is taken from `CMAKE_CUDA_ARCHITECTURES` if you set it (e.g. +`-DCMAKE_CUDA_ARCHITECTURES=87` for Jetson Orin), detected from the device tree +on NVIDIA Jetson modules, or `native` with CMake >= 3.24. + +With CUDA enabled, every PCPS acquisition block can evaluate its search grid on +the GPU by setting `Acquisition_XX.use_cuda=true` (or +`GNSS-SDR.use_cuda_acquisition=true` for all of them), and the experimental +`GPS_L1_CA_DLL_PLL_Tracking_GPU` tracking block becomes available. See +[docs/JETSON.md](./docs/JETSON.md) for a step-by-step guide on NVIDIA Jetson +(Orin, Xavier, TX2, Nano), including how to run the unit tests and the +CPU-vs-GPU acquisition benchmark. ## macOS @@ -1712,6 +1723,7 @@ Acquisition_1C.doppler_max=5000 ; Maximum expected Doppler shift [Hz] Acquisition_1C.doppler_step=250 ; Doppler step in the grid search [Hz] Acquisition_1C.dump=false ; Enables internal data file logging [true] or [false] Acquisition_1C.dump_filename=./acq_dump.dat ; Log path and filename +Acquisition_1C.use_cuda=false ; Evaluate the search grid on a CUDA GPU (requires -DENABLE_CUDA=ON) ``` and, for Galileo E1B channels: diff --git a/conf/Other/gnss-sdr_GPS_L1_gr_complex_gpu.conf b/conf/Other/gnss-sdr_GPS_L1_gr_complex_gpu.conf index f920c7cb6..f31266731 100644 --- a/conf/Other/gnss-sdr_GPS_L1_gr_complex_gpu.conf +++ b/conf/Other/gnss-sdr_GPS_L1_gr_complex_gpu.conf @@ -37,6 +37,7 @@ Channel.signal=1C ;######### ACQUISITION GLOBAL CONFIG ############ Acquisition_1C.implementation=GPS_L1_CA_PCPS_Acquisition +Acquisition_1C.use_cuda=true ; Evaluate the PCPS search grid on the GPU (any *_PCPS_Acquisition block supports it) Acquisition_1C.item_type=gr_complex Acquisition_1C.coherent_integration_time_ms=1 Acquisition_1C.threshold=0.005 diff --git a/docs/CHANGELOG.md b/docs/CHANGELOG.md index 81868c49c..ceafa74d4 100644 --- a/docs/CHANGELOG.md +++ b/docs/CHANGELOG.md @@ -179,6 +179,14 @@ All notable changes to GNSS-SDR will be documented in this file. days). A satellite that is already being tracked is never released because of this classification, and `PVT.elevation_mask` still decides which observations enter the navigation solution. Contributed by @joebre. +- New CUDA acquisition engine: with `-DENABLE_CUDA=ON`, any PCPS acquisition + block can evaluate its Doppler x code-phase search grid on the GPU with + batched cuFFTs by setting `Acquisition_XX.use_cuda=true` (or + `GNSS-SDR.use_cuda_acquisition=true`). Peak search and detection statistics + are unchanged, so results match the CPU implementation; the block falls back + to the CPU if the device cannot be initialized. Added `benchmark_pcps_grid` + (CPU baseline vs. GPU) and unit tests checking the GPU grid against the CPU + reference and running the full GPS L1 C/A adapter on a real capture. ### Improvements in Interoperability: @@ -424,6 +432,14 @@ All notable changes to GNSS-SDR will be documented in this file. - Refactored Python interpreter detection and improved CMake portability and robustness across dependency discovery, distro detection, and cross-compilation handling. +- The CUDA build (`-DENABLE_CUDA=ON`) works again with current toolkits and on + NVIDIA Jetson: removed the hardcoded `sm_30` (Kepler) architecture, which + CUDA >= 11 rejects; `CMAKE_CUDA_ARCHITECTURES` is now honored and detected + automatically on Jetson (Orin -> 87, Xavier -> 72, TX2 -> 62, Nano -> 53) or + set to `native` with CMake >= 3.24; the CUDA language standard follows the + host C++ standard (C++17); imported `CUDA::cudart`/`CUDA::cufft` targets are + linked explicitly; `-Wno-psabi` is no longer passed to `nvcc`. +- Added `docs/JETSON.md`, a build/verify/benchmark guide for NVIDIA Jetson. ### Improvements in Reliability: @@ -445,6 +461,13 @@ All notable changes to GNSS-SDR will be documented in this file. wall-clock GST alignment check for OSNMA tag processing, enabling replay of previously captured Galileo signals while keeping all other OSNMA verification steps active. +- `GPS_L1_CA_DLL_PLL_Tracking_GPU`: fixed a cross-block data race in the CUDA + multi-correlator kernel (the carrier wipe-off and the correlation were in the + same launch, synchronized only with `__syncthreads()`), fixed the + `cudaHostAlloc` flags (`cudaHostAllocMapped || cudaHostAllocWriteCombined` + evaluated to `cudaHostAllocPortable`), stopped calling `cudaDeviceReset()` + from a per-channel destructor (it tore down the context under the other + channels), and stopped `cudaFree()`-ing device aliases of host-mapped buffers. ### Improvements in Usability: diff --git a/docs/JETSON.md b/docs/JETSON.md new file mode 100644 index 000000000..d4de3099d --- /dev/null +++ b/docs/JETSON.md @@ -0,0 +1,300 @@ + +[comment]: # ( +SPDX-License-Identifier: GPL-3.0-or-later +) + +[comment]: # ( +SPDX-FileCopyrightText: 2026 Phillip Vu <36169227+phillipvu@users.noreply.github.com> +) + + +# Building GNSS-SDR with CUDA on NVIDIA Jetson + +This guide covers building GNSS-SDR with `-DENABLE_CUDA=ON` on NVIDIA Jetson +modules (Orin AGX / Orin NX / Orin Nano, and older Xavier / TX2 / Nano boards), +what the CUDA build gives you, and how to measure it against the CPU baseline. + +It was written and verified on a Jetson Orin Nano Super Developer Kit running +JetPack 7.2 (L4T r39, Ubuntu 24.04, CUDA 13.2, gcc 13). JetPack 6 (L4T r36, +Ubuntu 22.04, CUDA 12.x) differs only in package versions. + +## What the CUDA build enables + +| Block / feature | Selected with | Notes | +| ------------------------------------------ | --------------------------------------------------------------------------------- | -------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- | +| **GPU acquisition grid (all PCPS blocks)** | `Acquisition_XX.use_cuda=true` (or `GNSS-SDR.use_cuda_acquisition=true` globally) | The Doppler x code-phase search grid of every `*_PCPS_Acquisition` implementation is evaluated with batched cuFFTs. Peak search and statistics are unchanged, so results are identical to the CPU path. Falls back to the CPU automatically if the device cannot be initialized. | +| `GPS_L1_CA_DLL_PLL_Tracking_GPU` | `Tracking_1C.implementation=GPS_L1_CA_DLL_PLL_Tracking_GPU` | Legacy CUDA multi-correlator tracking block (GPS L1 C/A only, experimental). | + +Optional per-block override for multi-GPU hosts: `Acquisition_XX.cuda_device=N` +(`GNSS-SDR.cuda_device=N` globally). Jetson has a single device, so leave it +unset. + +## 1. Prerequisites on the Jetson + +JetPack already ships the CUDA toolkit, `nvcc`, and cuFFT under +`/usr/local/cuda`. Make sure `nvcc` is on your `PATH`: + +``` +$ echo 'export PATH=/usr/local/cuda/bin:$PATH' >> ~/.bashrc +$ echo 'export LD_LIBRARY_PATH=/usr/local/cuda/lib64:$LD_LIBRARY_PATH' >> ~/.bashrc +$ source ~/.bashrc +$ nvcc --version +``` + +Install the GNSS-SDR dependencies from the Ubuntu repositories (this is the +Debian/Ubuntu list from the main [README](../README.md), unchanged for arm64): + +``` +$ sudo apt-get install build-essential cmake git pkg-config libboost-dev \ + libboost-date-time-dev libboost-system-dev libboost-filesystem-dev \ + libboost-thread-dev libboost-chrono-dev libboost-serialization-dev \ + libabsl-dev libad9361-dev libarmadillo-dev libblas-dev \ + libgnutls-openssl-dev libgnutls28-dev libgtest-dev liblapack-dev \ + libmatio-dev libpcap-dev libprotobuf-dev libpugixml-dev libssl-dev \ + libuhd-dev gnuradio-dev gr-osmosdr protobuf-compiler python3-mako +``` + +Check that CMake is at least 3.17 (`cmake --version`); JetPack 6 ships 3.22 and +JetPack 7 ships 3.28. (CMake >= 3.24 is not required: on Jetson the build +detects the SoC from the device tree and selects the right `sm_XX` +automatically; see below.) + +## 2. Configure and build + +``` +$ git clone https://github.com/gnss-sdr/gnss-sdr +$ cd gnss-sdr +$ cmake -S . -B build -DENABLE_CUDA=ON -DENABLE_UNIT_TESTING=ON -DENABLE_BENCHMARKS=ON +$ cmake --build build -j$(nproc) +$ sudo cmake --install build +``` + +The same steps, plus dependency installation, tests and the benchmark, are +scripted in [utils/scripts/jetson-build.sh](../utils/scripts/jetson-build.sh): + +``` +$ utils/scripts/jetson-build.sh --deps --tests --bench +``` + +The configure log should show: + +``` +-- NVIDIA CUDA GPU Acceleration will be enabled. +-- Target CUDA architecture(s): 87 +-- The CUDA compiler identification is NVIDIA 13.2.86. Standard: C++17. +``` + +A full build with unit tests and benchmarks takes about 45 minutes on an Orin +Nano at `-j4` and peaks around 4 GB of RAM. + +### GPU architecture selection + +`CMAKE_CUDA_ARCHITECTURES` is chosen as follows (first match wins): + +1. Whatever you pass on the command line, e.g. `-DCMAKE_CUDA_ARCHITECTURES=87`. +2. On Jetson, from `/proc/device-tree/compatible`: `tegra234` (Orin) -> 87, + `tegra194` (Xavier) -> 72, `tegra186` (TX2) -> 62, `tegra210` (Nano/TX1) + -> 53. +3. `native` (CMake >= 3.24, any host). +4. Otherwise nvcc's default. + +To build a binary that runs on several Jetson generations, pass a list: +`-DCMAKE_CUDA_ARCHITECTURES="72;87"`. + +### Build tips for Jetson + +- Compile with all cores but watch memory on 8 GB modules: `-j4` is safer than + `-j$(nproc)` on Orin Nano / Orin NX 8 GB when building the unit tests. +- Put the board in max-performance mode before benchmarking: + `sudo nvpmodel -m 0 && sudo jetson_clocks`. +- `-DENABLE_UNIT_TESTING=OFF` roughly halves the build time if you only need the + receiver. + +## 3. Verify + +Run the CUDA-specific unit tests (they are built into `run_tests` only when +`ENABLE_CUDA=ON`): + +``` +$ cd build +$ ./tests/run_tests --gtest_filter='CudaPcpsEngineTest.*:GpsL1CaPcpsAcquisitionCudaTest.*:GpuMulticorrelatorTest.*' +``` + +Expected output (Orin Nano Super, JetPack 7.2): + +``` +CudaPcpsEngineTest.MatchesCpuReferenceSingleDwell + fft_size=4000 bins=40 dwells=1 max rel. error=2.7e-07 peak@(bin 27, idx 524) +CudaPcpsEngineTest.MatchesCpuReferenceNonCoherentAccumulation + fft_size=4000 bins=40 dwells=4 max rel. error=2.0e-07 +CudaPcpsEngineTest.MatchesCpuReferenceBitTransitionMode + fft_size=8000 bins=40 dwells=1 max rel. error=2.0e-07 +CudaPcpsEngineTest.MatchesCpuReferenceLongCoherentIntegration + fft_size=16000 bins=100 dwells=1 max rel. error=4.1e-07 +GpsL1CaPcpsAcquisitionCudaTest.SameEstimateAsCpu + CPU: Doppler=1700 Hz, delay=523 samples + CUDA: Doppler=1700 Hz, delay=523 samples +GpsL1CaPcpsAcquisitionCudaTest.SameEstimateAsCpuMakeTwoStep + CPU (two steps): Doppler=1740 Hz, delay=911 samples + CUDA (two steps): Doppler=1740 Hz, delay=911 samples +[ PASSED ] 11 tests. +``` + +The GPU grid matches the CPU grid to single-precision rounding (relative error a +few 1e-7), and the full adapter returns the same acquisition estimate on a real +capture. + +`CudaPcpsEngineTest.*` compares the GPU grid to a CPU reference sample by sample +(single dwell, non-coherent accumulation, bit-transition mode and 4 ms coherent +integration). `GpsL1CaPcpsAcquisitionCudaTest.*` runs the full +`GPS_L1_CA_PCPS_Acquisition` adapter on a real 4 Msps capture with +`use_cuda=true` and checks it returns the same Doppler/delay as the CPU path, +with and without the two-step (fine Doppler) search. + +## 4. Benchmark: GPU acquisition vs CPU baseline + +`benchmark_pcps_grid` (built with `-DENABLE_BENCHMARKS=ON`) times one complete +PCPS search grid (carrier wipe-off, forward FFT, code multiplication, inverse +FFT and squared magnitude for every Doppler bin) on the CPU (volk + FFTW via +`gr::fft`, the receiver's default) and on the GPU (`CudaPcpsEngine`, including +host <-> device transfers), for a sweep of FFT sizes and Doppler bin counts: + +``` +$ ./build/tests/benchmarks/benchmark_pcps_grid --benchmark_counters_tabular=true +``` + +The CPU baseline runs single-threaded, which is how `pcps_acquisition` uses it +(one acquisition worker thread per channel). Counters: + +- `dwells/s`: complete search grids per second. This is the figure that bounds + how many channels can acquire simultaneously in real time + (`Channels.in_acquisition`). +- `grid_cells/s`: Doppler x code-phase hypotheses evaluated per second. + +The `MeasureExecutionTime` case of `CudaPcpsEngineTest` prints a quick one-line +CPU/GPU comparison at the default 4 Msps / 1 ms / 40-bin geometry if you do not +want to build Google Benchmark. + +### Reference results + +Measured on a Jetson Orin Nano Super Developer Kit (8 GB, 6-core Cortex-A78AE, +1024-core Ampere GPU), JetPack 7.2 / CUDA 13.2, default `nvpmodel` 25 W mode +with CPU/GPU frequency scaling left on (i.e. no `jetson_clocks`), board +otherwise idle. Time is wall-clock per complete search grid; the GPU column +includes the host-to-device copy of the input, the device-to-host copy of the +magnitude grid, and the scatter into the acquisition block's per-bin buffers. + +| Samples per dwell (`fft_size`) | Doppler bins | CPU (µs) | GPU (µs) | Speed-up | GPU dwells/s | +| -----------------------------: | -----------: | -------: | -------: | -------: | -----------: | +| 2000 (2 Msps, 1 ms) | 21 | 759 | 128 | 5.9× | 7814 | +| 2000 | 41 | 1483 | 188 | 7.9× | 5314 | +| 2000 | 81 | 2934 | 321 | 9.1× | 3113 | +| 4000 (4 Msps, 1 ms) | 21 | 1581 | 200 | 7.9× | 5006 | +| 4000 | 41 | 3085 | 336 | 9.2× | 2976 | +| 4000 | 81 | 6152 | 611 | 10.1× | 1636 | +| 8000 (8 Msps, 1 ms) | 21 | 3355 | 367 | 9.1× | 2727 | +| 8000 | 41 | 6695 | 661 | 10.1× | 1513 | +| 8000 | 81 | 13600 | 1180 | 11.5× | 847 | +| 16000 (4 Msps, 4 ms) | 21 | 7699 | 756 | 10.2× | 1324 | +| 16000 | 41 | 15091 | 1334 | 11.3× | 750 | +| 16000 | 81 | 29888 | 2883 | 10.4× | 347 | +| 20000 (20 Msps, 1 ms) | 21 | 11962 | 912 | 13.1× | 1097 | +| 20000 | 41 | 23445 | 1697 | 13.8× | 589 | +| 20000 | 81 | 46426 | 3097 | 15.0× | 323 | +| 40000 (20 Msps, 2 ms) | 21 | 26399 | 2231 | 11.8× | 448 | +| 40000 | 41 | 51341 | 4194 | 12.2× | 238 | +| 40000 | 81 | 101573 | 8023 | 12.7× | 125 | + +Reading the table: with the default GPS L1 configuration (4 Msps, 1 ms, ±5 kHz +at 250 Hz = 41 bins) one Orin Nano CPU core sustains ~320 dwells/s, the GPU +~3000 dwells/s. At 20 Msps the CPU falls below real time for a single channel +with 81 bins (21 dwells/s of a 1 ms signal, i.e. 2% real time) while the GPU +stays at 323 dwells/s. Note the CPU figures are for one core; with several +channels in acquisition the CPU path scales with the cores you give it while the +channels' GPU engines share one device, so the per-channel speed-up shrinks as +`Channels.in_acquisition` grows. The numbers above are single-engine throughput; +concurrent multi-channel GPU throughput has not been characterized. + +Run the sweep yourself with: + +``` +$ ./build/tests/benchmarks/benchmark_pcps_grid --benchmark_counters_tabular=true --benchmark_min_time=0.5s +``` + +## 5. Run the receiver on the GPU + +Add to any configuration file: + +``` +GNSS-SDR.use_cuda_acquisition=true ; every PCPS acquisition block uses the GPU +``` + +or per block: + +``` +Acquisition_1C.use_cuda=true +Acquisition_1B.use_cuda=true +``` + +Look for this line in the log (`gnss-sdr --log_dir=/tmp` or stderr) to confirm +the GPU path is active: + +``` +PCPS acquisition grid will be computed on CUDA device Orin (FFT size 4000, up to 21 Doppler bins, ...) +``` + +If the device cannot be initialized the block logs a `WARNING` and continues on +the CPU, so a configuration that enables `use_cuda` stays usable on a machine +without a GPU (or built without `ENABLE_CUDA`). + +A ready-made example is +[conf/Other/gnss-sdr_GPS_L1_gr_complex_gpu.conf](../conf/Other/gnss-sdr_GPS_L1_gr_complex_gpu.conf). + +### End-to-end check on a real capture + +The public +[2013_04_04_GNSS_SIGNAL_at_CTTC_SPAIN](https://sourceforge.net/projects/gnss-sdr/files/data/2013_04_04_GNSS_SIGNAL_at_CTTC_SPAIN.tar.gz) +capture (GPS L1, 4 Msps `ishort`, 100 s) run through +`conf/File_input/GPS/gnss-sdr_GPS_L1_ishort.conf` with +`SignalSource.samples=480000000` (60 s), `Channels.in_acquisition=4`, and +`Acquisition_1C.use_cuda` set to `false` / `true`, on the Orin Nano Super (three +runs each, `next` branch): + +| Acquisition | Time to first fix (receiver time) | PVT solutions in 60 s | Processing time for 60 s of signal | +| ----------- | --------------------------------- | --------------------: | ---------------------------------: | +| CPU | 06:23:30.5 (3/3 runs) | 69 | 12.3 to 12.7 s | +| CUDA | 06:23:30.5 (3/3 runs) | 69 | 10.1 to 10.7 s | + +Both give the same position (41.2748 N, 1.9877 E, ~74 m, the CTTC roof). The +processing time is dominated by tracking on the CPU, so the GPU acquisition path +mostly frees CPU time here rather than shortening the run; the benefit grows +with the number of channels in acquisition, the sampling rate, and the Doppler +range. + +## 6. Notes on the implementation + +- `src/algorithms/acquisition/libs/cuda_pcps_engine.{h,cu}` is the engine. One + CUDA stream and one cuFFT batched plan per acquisition block (per channel), so + channels acquiring concurrently overlap on the device. All Doppler bins are + processed in a single batched forward/inverse FFT pair; three small kernels do + the wipe-off, code multiplication and magnitude/accumulation. Non-coherent + accumulation across dwells is kept on the device. +- `pcps_acquisition::doppler_grid()` dispatches to the engine when present and + otherwise to the unchanged CPU loop (`doppler_grid_cpu()`); peak search, CFAR + statistics, two-step refinement, dumping and the monitor output all run on the + host as before. +- The engine header contains no CUDA types, so it can be included from plain C++ + translation units. `CUDA_GPU_ACCEL=1` is defined project-wide when + `ENABLE_CUDA=ON`. +- Fixes to the existing CUDA tracking block that were needed to build and run on + JetPack 6 are listed in the [changelog](./CHANGELOG.md). + +## Troubleshooting + +| Symptom | Cause / fix | +| ------------------------------------------------------------------------------ | ------------------------------------------------------------------------------------------------------------------------------------ | +| `No CMAKE_CUDA_COMPILER could be found` | `nvcc` not on `PATH`; export `/usr/local/cuda/bin` or pass `-DCMAKE_CUDA_COMPILER=/usr/local/cuda/bin/nvcc`. | +| `nvcc fatal : Unsupported gpu architecture 'compute_30'` | Stale build directory from an older GNSS-SDR; delete `build/` and reconfigure. | +| `CUDA acquisition engine could not be initialized (... cudaErrorNoDevice ...)` | The process cannot see the GPU. Check `nvidia-smi`/`tegrastats`, and that the user is in the `video` group on Jetson. | +| `CUFFT_ALLOC_FAILED` for large FFTs | Not enough GPU memory for `bins x fft_size`. Reduce `doppler_max`/increase `doppler_step`, or lower `coherent_integration_time_ms`. | +| `unsupported GNU version! gcc versions later than N are not supported` | Host compiler newer than the toolkit supports; pass `-DCMAKE_CUDA_HOST_COMPILER=g++-N` (JetPack 6's gcc 11 is supported by CUDA 12). | diff --git a/src/algorithms/acquisition/gnuradio_blocks/pcps_acquisition.cc b/src/algorithms/acquisition/gnuradio_blocks/pcps_acquisition.cc index 3bb84413e..8ed5df195 100644 --- a/src/algorithms/acquisition/gnuradio_blocks/pcps_acquisition.cc +++ b/src/algorithms/acquisition/gnuradio_blocks/pcps_acquisition.cc @@ -40,6 +40,7 @@ #include #include #include +#include #if USE_GLOG_AND_GFLAGS #include @@ -196,6 +197,13 @@ pcps_acquisition::pcps_acquisition(const Acq_Conf& conf_) update_grid_doppler_wipeoffs(); +#if CUDA_GPU_ACCEL + if (d_acq_parameters.use_cuda) + { + init_cuda_engine(); + } +#endif + // While idle (not actively searching), general_work() only drains its input to avoid // stalling the upstream block, producing no output. Without a batching hint the TPB // scheduler wakes this block's thread for every small burst of new input, which is @@ -209,6 +217,68 @@ pcps_acquisition::pcps_acquisition(const Acq_Conf& conf_) } +#if CUDA_GPU_ACCEL +void pcps_acquisition::init_cuda_engine() +{ + const uint32_t max_bins = std::max(d_num_doppler_bins_step1_capacity, d_num_doppler_bins_step2); + auto engine = std::make_unique(d_fft_size, d_effective_fft_size, max_bins, d_acq_parameters.cuda_device); + if (!engine->is_valid()) + { + LOG(WARNING) << "CUDA acquisition engine could not be initialized (" << engine->last_error() + << "). Falling back to the CPU implementation."; + return; + } + d_cuda_engine = std::move(engine); + LOG(INFO) << "PCPS acquisition grid will be computed on CUDA device " << d_cuda_engine->device_name() + << " (FFT size " << d_fft_size << ", up to " << max_bins << " Doppler bins, " + << d_cuda_engine->device_bytes() / 1024 << " KiB device memory)"; + cuda_upload_wipeoffs(CudaPcpsEngine::MAIN_GRID); +} + + +void pcps_acquisition::cuda_upload_wipeoffs(CudaPcpsEngine::GridId grid) +{ + if (!d_cuda_engine) + { + return; + } + const bool step2 = (grid == CudaPcpsEngine::STEP2_GRID); + const uint32_t bins = step2 ? d_num_doppler_bins_step2 : d_num_doppler_bins_active; + if (bins == 0U || (step2 && d_grid_doppler_wipeoffs_step_two.empty())) + { + return; + } + std::vector*> rows(bins); + for (uint32_t k = 0; k < bins; k++) + { + rows[k] = step2 ? doppler_wipeoff_step_two_data(k) : doppler_wipeoff_data(k); + } + if (!d_cuda_engine->set_doppler_wipeoffs(grid, rows.data(), bins)) + { + LOG(WARNING) << "CUDA acquisition: failed to upload Doppler grid (" << d_cuda_engine->last_error() + << "). Falling back to the CPU implementation."; + d_cuda_engine.reset(); + } +} + + +bool pcps_acquisition::doppler_grid_cuda(const gr_complex* in) +{ + const auto bin_count = d_step_two ? d_num_doppler_bins_step2 : d_num_doppler_bins_active; + const auto grid_id = d_step_two ? CudaPcpsEngine::STEP2_GRID : CudaPcpsEngine::MAIN_GRID; + const uint32_t offset = (d_acq_parameters.bit_transition_flag ? d_effective_fft_size : 0); + const bool accumulate = (d_num_noncoherent_integrations_counter != 1); + + std::vector rows(bin_count); + for (uint32_t k = 0; k < bin_count; k++) + { + rows[k] = magnitude_grid_data(k); + } + return d_cuda_engine->compute_grid(in, grid_id, bin_count, offset, accumulate, rows.data()); +} +#endif + + pcps_acquisition::~pcps_acquisition() noexcept { try @@ -280,6 +350,14 @@ void pcps_acquisition::set_local_code(std::complex* code) d_fft_if->execute(); // We need the FFT of local code volk_32fc_conjugate_32fc(d_fft_codes.data(), d_fft_if->get_outbuf(), d_fft_size); +#if CUDA_GPU_ACCEL + if (d_cuda_engine && !d_cuda_engine->set_fft_codes(d_fft_codes.data())) + { + LOG(WARNING) << "CUDA acquisition: failed to upload local code (" << d_cuda_engine->last_error() + << "). Falling back to the CPU implementation."; + d_cuda_engine.reset(); + } +#endif } @@ -324,6 +402,9 @@ void pcps_acquisition::update_grid_doppler_wipeoffs() // keep using it as a noise-only reference without any change to its logic. update_local_carrier(own::span(doppler_wipeoff_data(0), d_fft_size), static_cast(d_doppler_bias + d_doppler_center)); update_local_carrier(own::span(doppler_wipeoff_data(1), d_fft_size), static_cast(d_doppler_bias + d_doppler_center + static_cast(d_doppler_max))); +#if CUDA_GPU_ACCEL + cuda_upload_wipeoffs(CudaPcpsEngine::MAIN_GRID); +#endif return; } for (uint32_t doppler_index = 0; doppler_index < d_num_doppler_bins_active; doppler_index++) @@ -331,6 +412,9 @@ void pcps_acquisition::update_grid_doppler_wipeoffs() const int32_t doppler = -static_cast(d_doppler_max) + d_doppler_center + d_doppler_step * doppler_index; update_local_carrier(own::span(doppler_wipeoff_data(doppler_index), d_fft_size), static_cast(d_doppler_bias + doppler)); } +#if CUDA_GPU_ACCEL + cuda_upload_wipeoffs(CudaPcpsEngine::MAIN_GRID); +#endif } @@ -343,6 +427,9 @@ void pcps_acquisition::update_grid_doppler_wipeoffs_step2() // stage, so the FDMA frequency offset must be added back to the wipeoff update_local_carrier(own::span(doppler_wipeoff_step_two_data(doppler_index), d_fft_size), static_cast(d_doppler_bias) + d_doppler_center_step_two + doppler); } +#if CUDA_GPU_ACCEL + cuda_upload_wipeoffs(CudaPcpsEngine::STEP2_GRID); +#endif } @@ -652,6 +739,24 @@ pcps_acquisition::AcquisitionResult pcps_acquisition::first_vs_second_peak_stati void pcps_acquisition::doppler_grid(const gr_complex* in) +{ +#if CUDA_GPU_ACCEL + if (d_cuda_engine) + { + if (doppler_grid_cuda(in)) + { + return; + } + LOG(WARNING) << "CUDA acquisition failed in channel " << d_channel << " (" << d_cuda_engine->last_error() + << "). Falling back to the CPU implementation."; + d_cuda_engine.reset(); + } +#endif + doppler_grid_cpu(in); +} + + +void pcps_acquisition::doppler_grid_cpu(const gr_complex* in) { const auto bin_count = d_step_two ? d_num_doppler_bins_step2 : d_num_doppler_bins_active; const auto* grid_doppler_wipeoffs = d_step_two ? d_grid_doppler_wipeoffs_step_two.data() : d_grid_doppler_wipeoffs.data(); diff --git a/src/algorithms/acquisition/gnuradio_blocks/pcps_acquisition.h b/src/algorithms/acquisition/gnuradio_blocks/pcps_acquisition.h index 0a27aa1b7..e6009ce5a 100644 --- a/src/algorithms/acquisition/gnuradio_blocks/pcps_acquisition.h +++ b/src/algorithms/acquisition/gnuradio_blocks/pcps_acquisition.h @@ -46,6 +46,9 @@ #include "acq_conf.h" #include "channel_fsm.h" #include "gnss_sdr_fft.h" +#if CUDA_GPU_ACCEL +#include "cuda_pcps_engine.h" +#endif #include #include #include // for gr_complex @@ -186,6 +189,12 @@ private: void update_grid_doppler_wipeoffs(); void update_grid_doppler_wipeoffs_step2(); void doppler_grid(const gr_complex* in); + void doppler_grid_cpu(const gr_complex* in); +#if CUDA_GPU_ACCEL + void init_cuda_engine(); + void cuda_upload_wipeoffs(CudaPcpsEngine::GridId grid); + bool doppler_grid_cuda(const gr_complex* in); +#endif AcquisitionResult compute_statistics(); void update_synchro(const AcquisitionResult& result); void handle_threshold_reached(AcquisitionResult& result); @@ -286,6 +295,9 @@ private: volk_gnsssdr::vector> d_grid_doppler_wipeoffs; volk_gnsssdr::vector> d_fft_codes; std::unique_ptr d_fft_if; +#if CUDA_GPU_ACCEL + std::unique_ptr d_cuda_engine; // null => CPU path +#endif }; diff --git a/src/algorithms/acquisition/libs/CMakeLists.txt b/src/algorithms/acquisition/libs/CMakeLists.txt index a2eb01230..88228a91f 100644 --- a/src/algorithms/acquisition/libs/CMakeLists.txt +++ b/src/algorithms/acquisition/libs/CMakeLists.txt @@ -15,6 +15,11 @@ if(ENABLE_FPGA) set(ACQUISITION_LIB_HEADERS ${ACQUISITION_LIB_HEADERS} fpga_acquisition.h) endif() +if(ENABLE_CUDA) + set(ACQUISITION_LIB_SOURCES ${ACQUISITION_LIB_SOURCES} cuda_pcps_engine.cu) + set(ACQUISITION_LIB_HEADERS ${ACQUISITION_LIB_HEADERS} cuda_pcps_engine.h) +endif() + list(SORT ACQUISITION_LIB_HEADERS) list(SORT ACQUISITION_LIB_SOURCES) @@ -71,3 +76,19 @@ if(ENABLE_FPGA) ) endif() +if(ENABLE_CUDA) + target_include_directories(acquisition_libs + PUBLIC ${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES} + ) + target_link_libraries(acquisition_libs + PUBLIC + CUDA::cudart + CUDA::cufft + ) + set_target_properties(acquisition_libs PROPERTIES + CUDA_SEPARABLE_COMPILATION ON + POSITION_INDEPENDENT_CODE ON + CUDA_RESOLVE_DEVICE_SYMBOLS ON + ) +endif() + diff --git a/src/algorithms/acquisition/libs/acq_conf.cc b/src/algorithms/acquisition/libs/acq_conf.cc index 30adaa17c..1e781a59a 100644 --- a/src/algorithms/acquisition/libs/acq_conf.cc +++ b/src/algorithms/acquisition/libs/acq_conf.cc @@ -93,6 +93,21 @@ void Acq_Conf::SetFromConfiguration(const ConfigurationInterface *configuration, enable_monitor_output = configuration->property("AcquisitionMonitor.enable_monitor", false); + // GPU offload of the search grid. A global GNSS-SDR.use_cuda_acquisition + // switch can be overridden per acquisition block with .use_cuda + use_cuda = configuration->property("GNSS-SDR.use_cuda_acquisition", use_cuda); + use_cuda = configuration->property(role + ".use_cuda", use_cuda); + cuda_device = configuration->property("GNSS-SDR.cuda_device", cuda_device); + cuda_device = configuration->property(role + ".cuda_device", cuda_device); +#if !CUDA_GPU_ACCEL + if (use_cuda) + { + LOG(WARNING) << "Parameter " << role << ".use_cuda is set but this build has no CUDA support " + << "(configure with -DENABLE_CUDA=ON). Falling back to the CPU implementation."; + use_cuda = false; + } +#endif + SetDerivedParams(); } diff --git a/src/algorithms/acquisition/libs/acq_conf.h b/src/algorithms/acquisition/libs/acq_conf.h index b54a9a771..8e6504cd4 100644 --- a/src/algorithms/acquisition/libs/acq_conf.h +++ b/src/algorithms/acquisition/libs/acq_conf.h @@ -100,6 +100,10 @@ public: // per-implementation in the .conf (e.g. Acquisition_1B.full_grid_search = true). bool full_grid_search{false}; + // Evaluate the PCPS grid on a CUDA GPU (requires ENABLE_CUDA at build time) + bool use_cuda{false}; + int32_t cuda_device{-1}; // CUDA device ordinal, -1 = default device + // Specific to some implementations bool acquire_pilot{false}; bool acquire_iq{false}; diff --git a/src/algorithms/acquisition/libs/cuda_pcps_engine.cu b/src/algorithms/acquisition/libs/cuda_pcps_engine.cu new file mode 100644 index 000000000..37daffabf --- /dev/null +++ b/src/algorithms/acquisition/libs/cuda_pcps_engine.cu @@ -0,0 +1,551 @@ +/*! + * \file cuda_pcps_engine.cu + * \brief GPU (CUDA + cuFFT) implementation of the Parallel Code Phase Search grid + * \author Phillip Vu, 2026. phillipvu(at)users.noreply.github.com + * + * ----------------------------------------------------------------------------- + * + * GNSS-SDR is a Global Navigation Satellite System software-defined receiver. + * This file is part of GNSS-SDR. + * + * Copyright (C) 2010-2026 (see AUTHORS file for a list of contributors) + * SPDX-License-Identifier: GPL-3.0-or-later + * + * ----------------------------------------------------------------------------- + */ + +#include "cuda_pcps_engine.h" +#include +#include +#include +#include +#include +#include + + +// --------------------------------------------------------------------------- +// Kernels +// --------------------------------------------------------------------------- + +namespace +{ +constexpr unsigned int THREADS_PER_BLOCK = 256U; + +__device__ __forceinline__ float2 cmul(float2 a, float2 b) +{ + return make_float2(a.x * b.x - a.y * b.y, a.x * b.y + a.y * b.x); +} + +// out[k*N + n] = in[n] * wipe[k*N + n] for k in [0, bins), n in [0, N) +__global__ void k_wipeoff(const float2* __restrict__ in, + const float2* __restrict__ wipe, + float2* __restrict__ out, + unsigned int fft_size, + unsigned int total) +{ + for (unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x; idx < total; idx += gridDim.x * blockDim.x) + { + out[idx] = cmul(in[idx % fft_size], wipe[idx]); + } +} + +// batch[k*N + n] *= codes[n] +__global__ void k_mult_codes(float2* __restrict__ batch, + const float2* __restrict__ codes, + unsigned int fft_size, + unsigned int total) +{ + for (unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x; idx < total; idx += gridDim.x * blockDim.x) + { + batch[idx] = cmul(batch[idx], codes[idx % fft_size]); + } +} + +// mag[k*E + i] (+)= |batch[k*N + offset + i]|^2 +__global__ void k_magnitude(const float2* __restrict__ batch, + float* __restrict__ mag, + unsigned int fft_size, + unsigned int effective_fft_size, + unsigned int offset, + unsigned int total, + int accumulate) +{ + for (unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x; idx < total; idx += gridDim.x * blockDim.x) + { + const unsigned int k = idx / effective_fft_size; + const unsigned int i = idx - k * effective_fft_size; + const float2 v = batch[k * fft_size + offset + i]; + const float m = v.x * v.x + v.y * v.y; + mag[idx] = accumulate ? (mag[idx] + m) : m; + } +} + +inline unsigned int grid_for(unsigned int total, int multiprocessors) +{ + const unsigned int needed = (total + THREADS_PER_BLOCK - 1U) / THREADS_PER_BLOCK; + // Enough blocks to fill the device several times over, but bounded so + // that very long FFTs (e.g. 20 ms at 20 Msps) do not launch huge grids. + const unsigned int cap = static_cast(std::max(1, multiprocessors)) * 8U; + return std::max(1U, std::min(needed, cap)); +} + +const char* cufft_error_string(cufftResult r) +{ + // Only codes present in every cuFFT release since 10.x: newer toolkits + // (cuFFT >= 12, CUDA 13) removed some of the older enumerators. + switch (r) + { + case CUFFT_SUCCESS: + return "CUFFT_SUCCESS"; + case CUFFT_INVALID_PLAN: + return "CUFFT_INVALID_PLAN"; + case CUFFT_ALLOC_FAILED: + return "CUFFT_ALLOC_FAILED"; + case CUFFT_INVALID_TYPE: + return "CUFFT_INVALID_TYPE"; + case CUFFT_INVALID_VALUE: + return "CUFFT_INVALID_VALUE"; + case CUFFT_INTERNAL_ERROR: + return "CUFFT_INTERNAL_ERROR"; + case CUFFT_EXEC_FAILED: + return "CUFFT_EXEC_FAILED"; + case CUFFT_SETUP_FAILED: + return "CUFFT_SETUP_FAILED"; + case CUFFT_INVALID_SIZE: + return "CUFFT_INVALID_SIZE"; + case CUFFT_UNALIGNED_DATA: + return "CUFFT_UNALIGNED_DATA"; + case CUFFT_INVALID_DEVICE: + return "CUFFT_INVALID_DEVICE"; + case CUFFT_NO_WORKSPACE: + return "CUFFT_NO_WORKSPACE"; + case CUFFT_NOT_IMPLEMENTED: + return "CUFFT_NOT_IMPLEMENTED"; + case CUFFT_NOT_SUPPORTED: + return "CUFFT_NOT_SUPPORTED"; + default: + return "CUFFT_ERROR (see cufft.h for the numeric code)"; + } +} +} // namespace + + +// --------------------------------------------------------------------------- +// Private state +// --------------------------------------------------------------------------- + +struct CudaPcpsEngine::Impl +{ + uint32_t fft_size{0}; + uint32_t effective_fft_size{0}; + uint32_t max_bins{0}; + int device{0}; + int multiprocessors{1}; + bool valid{false}; + std::string error; + std::string dev_name; + size_t dev_bytes{0}; + + cudaStream_t stream{nullptr}; + + float2* d_in{nullptr}; // fft_size + float2* d_codes{nullptr}; // fft_size + std::array d_wipe{{nullptr, nullptr}}; // max_bins * fft_size each + std::array wipe_bins{{0U, 0U}}; // bins uploaded per grid + float2* d_batch{nullptr}; // max_bins * fft_size (in-place FFT) + float* d_mag{nullptr}; // max_bins * effective_fft_size + + float2* h_in{nullptr}; // pinned staging, fft_size + float* h_mag{nullptr}; // pinned staging, max_bins * effective_fft_size + + // One cuFFT plan per grid; (re)created when the batch size changes. + std::array plan{{0, 0}}; + std::array plan_batch{{0U, 0U}}; + + bool fail(const char* what, cudaError_t e) + { + std::ostringstream ss; + ss << what << ": " << cudaGetErrorName(e) << " (" << cudaGetErrorString(e) << ")"; + error = ss.str(); + return false; + } + + bool fail(const char* what, cufftResult r) + { + std::ostringstream ss; + ss << what << ": " << cufft_error_string(r) << " [" << static_cast(r) << "]"; + error = ss.str(); + return false; + } + + template + bool dev_alloc(T** ptr, size_t count, const char* what) + { + const size_t bytes = count * sizeof(T); + const cudaError_t e = cudaMalloc(reinterpret_cast(ptr), bytes); + if (e != cudaSuccess) + { + return fail(what, e); + } + dev_bytes += bytes; + return true; + } + + bool ensure_plan(int grid, uint32_t bins) + { + if (plan_batch[grid] == bins && plan[grid] != 0) + { + return true; + } + if (plan[grid] != 0) + { + cufftDestroy(plan[grid]); + plan[grid] = 0; + plan_batch[grid] = 0U; + } + cufftResult r = cufftPlan1d(&plan[grid], static_cast(fft_size), CUFFT_C2C, static_cast(bins)); + if (r != CUFFT_SUCCESS) + { + plan[grid] = 0; + return fail("cufftPlan1d", r); + } + r = cufftSetStream(plan[grid], stream); + if (r != CUFFT_SUCCESS) + { + return fail("cufftSetStream", r); + } + plan_batch[grid] = bins; + return true; + } + + void release() + { + cudaSetDevice(device); + if (stream != nullptr) + { + cudaStreamSynchronize(stream); + } + for (int g = 0; g < CudaPcpsEngine::NUM_GRIDS; g++) + { + if (plan[g] != 0) + { + cufftDestroy(plan[g]); + plan[g] = 0; + } + if (d_wipe[g] != nullptr) + { + cudaFree(d_wipe[g]); + d_wipe[g] = nullptr; + } + } + if (d_in != nullptr) cudaFree(d_in); + if (d_codes != nullptr) cudaFree(d_codes); + if (d_batch != nullptr) cudaFree(d_batch); + if (d_mag != nullptr) cudaFree(d_mag); + if (h_in != nullptr) cudaFreeHost(h_in); + if (h_mag != nullptr) cudaFreeHost(h_mag); + d_in = d_codes = d_batch = nullptr; + d_mag = nullptr; + h_in = nullptr; + h_mag = nullptr; + if (stream != nullptr) + { + cudaStreamDestroy(stream); + stream = nullptr; + } + valid = false; + } +}; + + +// --------------------------------------------------------------------------- +// Public interface +// --------------------------------------------------------------------------- + +CudaPcpsEngine::CudaPcpsEngine(uint32_t fft_size, uint32_t effective_fft_size, uint32_t max_doppler_bins, int device) + : p(new Impl()) +{ + p->fft_size = fft_size; + p->effective_fft_size = effective_fft_size; + p->max_bins = std::max(1U, max_doppler_bins); + + if (fft_size == 0U || effective_fft_size == 0U || effective_fft_size > fft_size) + { + p->error = "invalid FFT geometry"; + return; + } + + int count = 0; + cudaError_t e = cudaGetDeviceCount(&count); + if (e != cudaSuccess || count == 0) + { + p->fail("cudaGetDeviceCount", e == cudaSuccess ? cudaErrorNoDevice : e); + return; + } + if (device < 0) + { + e = cudaGetDevice(&p->device); + if (e != cudaSuccess) + { + p->fail("cudaGetDevice", e); + return; + } + } + else + { + if (device >= count) + { + p->error = "requested CUDA device does not exist"; + return; + } + p->device = device; + } + e = cudaSetDevice(p->device); + if (e != cudaSuccess) + { + p->fail("cudaSetDevice", e); + return; + } + + cudaDeviceProp prop{}; + if (cudaGetDeviceProperties(&prop, p->device) == cudaSuccess) + { + p->dev_name = prop.name; + p->multiprocessors = prop.multiProcessorCount; + } + + e = cudaStreamCreateWithFlags(&p->stream, cudaStreamNonBlocking); + if (e != cudaSuccess) + { + p->fail("cudaStreamCreate", e); + return; + } + + const size_t n = static_cast(p->max_bins) * fft_size; + const size_t m = static_cast(p->max_bins) * effective_fft_size; + if (!p->dev_alloc(&p->d_in, fft_size, "cudaMalloc(d_in)") || + !p->dev_alloc(&p->d_codes, fft_size, "cudaMalloc(d_codes)") || + !p->dev_alloc(&p->d_batch, n, "cudaMalloc(d_batch)") || + !p->dev_alloc(&p->d_mag, m, "cudaMalloc(d_mag)")) + { + p->release(); + return; + } + // The wipe-off grids are allocated on first upload so that a block which + // never runs the two-step search does not pay for the second grid. + + e = cudaHostAlloc(reinterpret_cast(&p->h_in), fft_size * sizeof(float2), cudaHostAllocDefault); + if (e != cudaSuccess) + { + p->fail("cudaHostAlloc(h_in)", e); + p->release(); + return; + } + e = cudaHostAlloc(reinterpret_cast(&p->h_mag), m * sizeof(float), cudaHostAllocDefault); + if (e != cudaSuccess) + { + p->fail("cudaHostAlloc(h_mag)", e); + p->release(); + return; + } + + e = cudaMemsetAsync(p->d_codes, 0, fft_size * sizeof(float2), p->stream); + if (e == cudaSuccess) e = cudaMemsetAsync(p->d_mag, 0, m * sizeof(float), p->stream); + if (e == cudaSuccess) e = cudaStreamSynchronize(p->stream); + if (e != cudaSuccess) + { + p->fail("cudaMemset", e); + p->release(); + return; + } + + p->valid = true; +} + + +CudaPcpsEngine::~CudaPcpsEngine() +{ + if (p) + { + p->release(); + } +} + + +bool CudaPcpsEngine::is_valid() const +{ + return p->valid; +} + + +const std::string& CudaPcpsEngine::last_error() const +{ + return p->error; +} + + +const std::string& CudaPcpsEngine::device_name() const +{ + return p->dev_name; +} + + +size_t CudaPcpsEngine::device_bytes() const +{ + return p->dev_bytes; +} + + +bool CudaPcpsEngine::set_fft_codes(const std::complex* fft_codes) +{ + if (!p->valid) + { + 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); + if (e != cudaSuccess) + { + return p->fail("cudaMemcpy(codes)", e); + } + return true; +} + + +bool CudaPcpsEngine::set_doppler_wipeoffs(GridId grid, const std::complex* const* wipeoffs, uint32_t bins) +{ + if (!p->valid || grid < 0 || grid >= NUM_GRIDS) + { + return false; + } + if (bins == 0U || bins > p->max_bins) + { + p->error = "set_doppler_wipeoffs: bins out of range"; + return false; + } + cudaSetDevice(p->device); + if (p->d_wipe[grid] == nullptr) + { + if (!p->dev_alloc(&p->d_wipe[grid], static_cast(p->max_bins) * p->fft_size, "cudaMalloc(d_wipe)")) + { + return false; + } + } + // Rows live in separate host allocations; copy them one by one. This is + // 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); + 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) + { + return p->fail("cudaMemcpy(wipeoffs)", e); + } + } + p->wipe_bins[grid] = bins; + + // Build the cuFFT plan and run the whole pipeline once on the current + // (zero) input. Plan creation and first kernel launches can take tens of + // milliseconds on Jetson; doing them here instead of on the first real + // dwell keeps acquisition latency deterministic when the block runs in + // non-blocking mode and samples flow past while the worker is busy. + if (!p->ensure_plan(grid, bins)) + { + return false; + } + return run_pipeline(grid, bins, 0U, false); +} + + +bool CudaPcpsEngine::run_pipeline(int grid, uint32_t bins, uint32_t offset, bool accumulate) +{ + const uint32_t N = p->fft_size; + const uint32_t E = p->effective_fft_size; + const unsigned int total_c = bins * N; + const unsigned int total_m = bins * E; + cudaStream_t s = p->stream; + + // Carrier wipe-off for every Doppler bin at once + k_wipeoff<<multiprocessors), THREADS_PER_BLOCK, 0, s>>>(p->d_in, p->d_wipe[grid], p->d_batch, N, total_c); + cudaError_t e = cudaGetLastError(); + if (e != cudaSuccess) return p->fail("k_wipeoff", e); + + // Batched forward FFT (in place) + cufftResult r = cufftExecC2C(p->plan[grid], reinterpret_cast(p->d_batch), reinterpret_cast(p->d_batch), CUFFT_FORWARD); + if (r != CUFFT_SUCCESS) return p->fail("cufftExecC2C(forward)", r); + + // Multiply by conj(FFT(code)) + k_mult_codes<<multiprocessors), THREADS_PER_BLOCK, 0, s>>>(p->d_batch, p->d_codes, N, total_c); + e = cudaGetLastError(); + if (e != cudaSuccess) return p->fail("k_mult_codes", e); + + // Batched inverse FFT (in place, unnormalized like FFTW/gr::fft) + r = cufftExecC2C(p->plan[grid], reinterpret_cast(p->d_batch), reinterpret_cast(p->d_batch), CUFFT_INVERSE); + if (r != CUFFT_SUCCESS) return p->fail("cufftExecC2C(inverse)", r); + + // Squared magnitude (+ non-coherent accumulation) into the device grid + k_magnitude<<multiprocessors), THREADS_PER_BLOCK, 0, s>>>(p->d_batch, p->d_mag, N, E, offset, total_m, accumulate ? 1 : 0); + e = cudaGetLastError(); + if (e != cudaSuccess) return p->fail("k_magnitude", e); + + e = cudaStreamSynchronize(s); + if (e != cudaSuccess) return p->fail("cudaStreamSynchronize", e); + return true; +} + + +bool CudaPcpsEngine::compute_grid(const std::complex* in, GridId grid, uint32_t bins, uint32_t offset, + bool accumulate, float* const* magnitude_out) +{ + if (!p->valid || grid < 0 || grid >= NUM_GRIDS) + { + return false; + } + if (bins == 0U || bins > p->wipe_bins[grid]) + { + p->error = "compute_grid: bins exceed the uploaded Doppler grid"; + return false; + } + if (offset + p->effective_fft_size > p->fft_size) + { + p->error = "compute_grid: offset out of range"; + return false; + } + cudaSetDevice(p->device); + if (!p->ensure_plan(grid, bins)) + { + return false; + } + + const uint32_t N = p->fft_size; + const uint32_t E = p->effective_fft_size; + const size_t in_bytes = static_cast(N) * sizeof(float2); + const size_t mag_bytes = static_cast(bins) * E * sizeof(float); + cudaStream_t s = p->stream; + + // Input: pageable -> pinned staging -> device + std::memcpy(p->h_in, in, in_bytes); + cudaError_t e = cudaMemcpyAsync(p->d_in, p->h_in, in_bytes, cudaMemcpyHostToDevice, s); + if (e != cudaSuccess) return p->fail("cudaMemcpyAsync(in)", e); + + // Wipe-off, FFT, code multiply, IFFT, magnitude (synchronizes the stream) + if (!run_pipeline(grid, bins, offset, accumulate)) + { + return false; + } + + // Device -> pinned staging -> caller's per-bin rows + e = cudaMemcpyAsync(p->h_mag, p->d_mag, mag_bytes, cudaMemcpyDeviceToHost, s); + if (e != cudaSuccess) return p->fail("cudaMemcpyAsync(mag)", e); + e = cudaStreamSynchronize(s); + if (e != cudaSuccess) return p->fail("cudaStreamSynchronize", e); + + const size_t row_bytes = static_cast(E) * sizeof(float); + for (uint32_t k = 0; k < bins; k++) + { + std::memcpy(magnitude_out[k], p->h_mag + static_cast(k) * E, row_bytes); + } + return true; +} diff --git a/src/algorithms/acquisition/libs/cuda_pcps_engine.h b/src/algorithms/acquisition/libs/cuda_pcps_engine.h new file mode 100644 index 000000000..103baa8f4 --- /dev/null +++ b/src/algorithms/acquisition/libs/cuda_pcps_engine.h @@ -0,0 +1,130 @@ +/*! + * \file cuda_pcps_engine.h + * \brief GPU (CUDA + cuFFT) implementation of the Parallel Code Phase Search grid + * \author Phillip Vu, 2026. phillipvu(at)users.noreply.github.com + * + * Computes, for a batch of Doppler bins, the squared magnitude of the circular + * cross-correlation between the carrier-wiped input and the local code, using + * batched FFTs on an NVIDIA GPU. It is the GPU counterpart of the inner loop of + * pcps_acquisition::doppler_grid(). The public interface deliberately contains + * no CUDA types so that it can be consumed from plain C++ translation units; + * the implementation lives in cuda_pcps_engine.cu. + * + * ----------------------------------------------------------------------------- + * + * GNSS-SDR is a Global Navigation Satellite System software-defined receiver. + * This file is part of GNSS-SDR. + * + * Copyright (C) 2010-2026 (see AUTHORS file for a list of contributors) + * SPDX-License-Identifier: GPL-3.0-or-later + * + * ----------------------------------------------------------------------------- + */ + +#ifndef GNSS_SDR_CUDA_PCPS_ENGINE_H +#define GNSS_SDR_CUDA_PCPS_ENGINE_H + +#include +#include +#include +#include + +/** \addtogroup Acquisition + * \{ */ +/** \addtogroup acquisition_libs + * \{ */ + + +/*! + * \brief Batched PCPS grid evaluation on a CUDA device. + * + * Typical usage (mirrors pcps_acquisition): + * 1. Construct with the FFT geometry of the acquisition block. + * 2. set_doppler_wipeoffs() whenever the Doppler grid changes. + * 3. set_fft_codes() whenever the local code (PRN) changes. + * 4. compute_grid() once per dwell; the magnitude grid is written to the + * host rows supplied by the caller, so the existing CPU peak-search and + * statistics code can be reused unchanged. + * + * Each instance owns a CUDA stream, so several channels acquiring at the same + * time run their grids concurrently on the device. + */ +class CudaPcpsEngine +{ +public: + //! Identifier of the Doppler grid uploaded with set_doppler_wipeoffs() + enum GridId : int + { + MAIN_GRID = 0, //!< Coarse grid (doppler_min .. doppler_max, doppler_step) + STEP2_GRID = 1, //!< Fine grid used when make_two_steps=true + NUM_GRIDS = 2 + }; + + /*! + * \param fft_size Length of the FFTs (input vector length). + * \param effective_fft_size Number of code-phase cells written per Doppler bin + * (fft_size, or fft_size/2 when bit_transition_flag is set). + * \param max_doppler_bins Upper bound on the number of bins in any grid. + * \param device CUDA device ordinal, or -1 for the current/default device. + */ + CudaPcpsEngine(uint32_t fft_size, uint32_t effective_fft_size, uint32_t max_doppler_bins, int device = -1); + ~CudaPcpsEngine(); + + CudaPcpsEngine(const CudaPcpsEngine&) = delete; + CudaPcpsEngine& operator=(const CudaPcpsEngine&) = delete; + CudaPcpsEngine(CudaPcpsEngine&&) = delete; + CudaPcpsEngine& operator=(CudaPcpsEngine&&) = delete; + + //! True when construction succeeded and the engine can be used. + bool is_valid() const; + + //! Human-readable description of the last error (empty if none). + const std::string& last_error() const; + + //! Name of the CUDA device in use (empty if not valid). + const std::string& device_name() const; + + /*! + * \brief Upload conj(FFT(local code)), fft_size elements. + */ + bool set_fft_codes(const std::complex* fft_codes); + + /*! + * \brief Upload the Doppler wipe-off carriers for one grid. + * \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). + */ + bool set_doppler_wipeoffs(GridId grid, const std::complex* const* wipeoffs, uint32_t bins); + + /*! + * \brief Evaluate the PCPS grid for `bins` Doppler bins. + * + * For every bin k: mag_k[i] (+)= | IFFT( FFT(in .* w_k) .* C )[offset + i] |^2, + * with C = conj(FFT(code)), i in [0, effective_fft_size). + * + * \param in Input samples (fft_size complex values, host memory). + * \param grid Which wipe-off grid to use. + * \param bins Number of Doppler bins to compute (<= bins uploaded for `grid`). + * \param offset First sample of the output window (fft_size/2 when bit_transition_flag). + * \param accumulate false: overwrite the output rows; true: add to their current contents + * (non-coherent integration across dwells). Accumulation is kept on + * the device, so the host rows are always the full running sum. + * \param magnitude_out Array of `bins` host pointers, each with room for effective_fft_size floats. + */ + bool compute_grid(const std::complex* in, GridId grid, uint32_t bins, uint32_t offset, + bool accumulate, float* const* magnitude_out); + + //! Number of bytes allocated on the device by this instance. + size_t device_bytes() const; + +private: + struct Impl; + std::unique_ptr p; + bool run_pipeline(int grid, uint32_t bins, uint32_t offset, bool accumulate); +}; + + +/** \} */ +/** \} */ +#endif // GNSS_SDR_CUDA_PCPS_ENGINE_H diff --git a/src/algorithms/tracking/adapters/CMakeLists.txt b/src/algorithms/tracking/adapters/CMakeLists.txt index 97ee94048..6f6dc6840 100644 --- a/src/algorithms/tracking/adapters/CMakeLists.txt +++ b/src/algorithms/tracking/adapters/CMakeLists.txt @@ -94,6 +94,7 @@ if(ENABLE_CUDA) target_include_directories(tracking_adapters PUBLIC ${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES} ) + target_link_libraries(tracking_adapters PUBLIC CUDA::cudart) else() target_link_libraries(tracking_adapters PUBLIC ${CUDA_LIBRARIES} diff --git a/src/algorithms/tracking/gnuradio_blocks/CMakeLists.txt b/src/algorithms/tracking/gnuradio_blocks/CMakeLists.txt index 4f1b724cb..69bc157b7 100644 --- a/src/algorithms/tracking/gnuradio_blocks/CMakeLists.txt +++ b/src/algorithms/tracking/gnuradio_blocks/CMakeLists.txt @@ -99,6 +99,7 @@ if(ENABLE_CUDA) target_include_directories(tracking_gr_blocks PUBLIC ${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES} ) + target_link_libraries(tracking_gr_blocks PUBLIC CUDA::cudart) else() target_link_libraries(tracking_gr_blocks PUBLIC cuda_correlator_lib diff --git a/src/algorithms/tracking/gnuradio_blocks/gps_l1_ca_dll_pll_tracking_gpu_cc.cc b/src/algorithms/tracking/gnuradio_blocks/gps_l1_ca_dll_pll_tracking_gpu_cc.cc index 83bc7e93b..79556b884 100644 --- a/src/algorithms/tracking/gnuradio_blocks/gps_l1_ca_dll_pll_tracking_gpu_cc.cc +++ b/src/algorithms/tracking/gnuradio_blocks/gps_l1_ca_dll_pll_tracking_gpu_cc.cc @@ -97,12 +97,12 @@ Gps_L1_Ca_Dll_Pll_Tracking_GPU_cc::Gps_L1_Ca_Dll_Pll_Tracking_GPU_cc( // pinned memory mode - use special function to get OS-pinned memory d_n_correlator_taps = 3; // Early, Prompt, and Late // Get space for a vector with the C/A code replica sampled 1x/chip - cudaHostAlloc(reinterpret_cast(&d_ca_code), (static_cast(GPS_L1_CA_CODE_LENGTH_CHIPS) * sizeof(gr_complex)), cudaHostAllocMapped || cudaHostAllocWriteCombined); + cudaHostAlloc(reinterpret_cast(&d_ca_code), (static_cast(GPS_L1_CA_CODE_LENGTH_CHIPS) * sizeof(gr_complex)), cudaHostAllocMapped); // Get space for the resampled early / prompt / late local replicas - cudaHostAlloc(reinterpret_cast(&d_local_code_shift_chips), d_n_correlator_taps * sizeof(float), cudaHostAllocMapped || cudaHostAllocWriteCombined); - cudaHostAlloc(reinterpret_cast(&in_gpu), 2 * d_vector_length * sizeof(gr_complex), cudaHostAllocMapped || cudaHostAllocWriteCombined); + cudaHostAlloc(reinterpret_cast(&d_local_code_shift_chips), d_n_correlator_taps * sizeof(float), cudaHostAllocMapped); + cudaHostAlloc(reinterpret_cast(&in_gpu), 2 * d_vector_length * sizeof(gr_complex), cudaHostAllocMapped); // correlator outputs (scalar) - cudaHostAlloc(reinterpret_cast(&d_correlator_outs), sizeof(gr_complex) * d_n_correlator_taps, cudaHostAllocMapped || cudaHostAllocWriteCombined); + cudaHostAlloc(reinterpret_cast(&d_correlator_outs), sizeof(gr_complex) * d_n_correlator_taps, cudaHostAllocMapped); // Set TAPs delay values [chips] d_local_code_shift_chips[0] = -d_early_late_spc_chips; diff --git a/src/algorithms/tracking/libs/CMakeLists.txt b/src/algorithms/tracking/libs/CMakeLists.txt index 682208433..c4402d1cd 100644 --- a/src/algorithms/tracking/libs/CMakeLists.txt +++ b/src/algorithms/tracking/libs/CMakeLists.txt @@ -43,7 +43,8 @@ set(TRACKING_LIB_HEADERS ) if(ENABLE_CUDA) - list(APPEND CUDA_NVCC_FLAGS "-gencode arch=compute_30,code=sm_30; -O3; -use_fast_math -default-stream per-thread") + # Architecture selection is driven by CMAKE_CUDA_ARCHITECTURES (see the + # top-level CMakeLists.txt); do not hardcode -gencode flags here. if(CMAKE_VERSION VERSION_GREATER 3.11) set(TRACKING_LIB_SOURCES ${TRACKING_LIB_SOURCES} cuda_multicorrelator.cu) set(TRACKING_LIB_HEADERS ${TRACKING_LIB_HEADERS} cuda_multicorrelator.h) @@ -105,6 +106,10 @@ if(ENABLE_CUDA) target_include_directories(tracking_libs PUBLIC ${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES} ) + target_link_libraries(tracking_libs PUBLIC CUDA::cudart) + target_compile_options(tracking_libs PRIVATE + $<$:--default-stream=per-thread> + ) else() target_link_libraries(tracking_libs PUBLIC ${CUDA_LIBRARIES} diff --git a/src/algorithms/tracking/libs/cuda_multicorrelator.cu b/src/algorithms/tracking/libs/cuda_multicorrelator.cu index 44599c6bd..ff96d659c 100644 --- a/src/algorithms/tracking/libs/cuda_multicorrelator.cu +++ b/src/algorithms/tracking/libs/cuda_multicorrelator.cu @@ -29,24 +29,17 @@ #define ACCUM_N 128 -__global__ void Doppler_wippe_scalarProdGPUCPXxN_shifts_chips( - GPU_Complex *d_corr_out, - GPU_Complex *d_sig_in, +// Stage 1: carrier wipe-off. Every block covers a grid-strided slice of the +// input, so this MUST be a separate launch from the correlator below: a +// __syncthreads() only orders threads within one block, and the correlator +// reads the whole wiped vector from every block. +__global__ void Doppler_wipeoff_kernel( + const GPU_Complex *d_sig_in, GPU_Complex *d_sig_wiped, - GPU_Complex *d_local_code_in, - float *d_shifts_chips, - int code_length_chips, - float code_phase_step_chips, - float rem_code_phase_chips, - int vectorN, int elementN, float rem_carrier_phase_in_rad, float phase_step_rad) { - //Accumulators cache - __shared__ GPU_Complex accumResult[ACCUM_N]; - - // CUDA version of floating point NCO and vector dot product integrated float sin; float cos; for (int i = blockIdx.x * blockDim.x + threadIdx.x; @@ -56,8 +49,24 @@ __global__ void Doppler_wippe_scalarProdGPUCPXxN_shifts_chips( __sincosf(rem_carrier_phase_in_rad + i * phase_step_rad, &sin, &cos); d_sig_wiped[i] = d_sig_in[i] * GPU_Complex(cos, -sin); } +} + + +// Stage 2: multi-tap correlator with integrated local code resampler. +__global__ void Doppler_wippe_scalarProdGPUCPXxN_shifts_chips( + GPU_Complex *d_corr_out, + const GPU_Complex *d_sig_wiped, + const GPU_Complex *d_local_code_in, + const float *d_shifts_chips, + int code_length_chips, + float code_phase_step_chips, + float rem_code_phase_chips, + int vectorN, + int elementN) +{ + //Accumulators cache + __shared__ GPU_Complex accumResult[ACCUM_N]; - __syncthreads(); //////////////////////////////////////////////////////////////////////////// // Cycle through every pair of vectors, // taking into account that vector counts can be different @@ -162,8 +171,8 @@ bool cuda_multicorrelator::init_cuda_integrated_resampler( printf("L2 Cache size= %u \n", prop.l2CacheSize); printf("maxThreadsPerBlock= %u \n", prop.maxThreadsPerBlock); printf("maxGridSize= %i \n", prop.maxGridSize[0]); - printf("sharedMemPerBlock= %lu \n", prop.sharedMemPerBlock); - printf("deviceOverlap= %i \n", prop.deviceOverlap); + printf("sharedMemPerBlock= %zu \n", prop.sharedMemPerBlock); + printf("asyncEngineCount= %i \n", prop.asyncEngineCount); printf("multiProcessorCount= %i \n", prop.multiProcessorCount); } else @@ -179,8 +188,8 @@ bool cuda_multicorrelator::init_cuda_integrated_resampler( printf("L2 Cache size= %u \n", prop.l2CacheSize); printf("maxThreadsPerBlock= %u \n", prop.maxThreadsPerBlock); printf("maxGridSize= %i \n", prop.maxGridSize[0]); - printf("sharedMemPerBlock= %lu \n", prop.sharedMemPerBlock); - printf("deviceOverlap= %i \n", prop.deviceOverlap); + printf("sharedMemPerBlock= %zu \n", prop.sharedMemPerBlock); + printf("asyncEngineCount= %i \n", prop.asyncEngineCount); printf("multiProcessorCount= %i \n", prop.multiProcessorCount); } @@ -319,9 +328,17 @@ bool cuda_multicorrelator::Carrier_wipeoff_multicorrelator_resampler_cuda( //launch the multitap correlator with integrated local code resampler! + Doppler_wipeoff_kernel<<>>( + d_sig_in, + d_sig_doppler_wiped, + signal_length_samples, + rem_carrier_phase_in_rad, + phase_step_rad); + gpuErrchk(cudaPeekAtLastError()); + + // Same stream: the correlator starts only after the wipe-off has completed Doppler_wippe_scalarProdGPUCPXxN_shifts_chips<<>>( d_corr_out, - d_sig_in, d_sig_doppler_wiped, d_local_codes_in, d_shifts_chips, @@ -329,9 +346,7 @@ bool cuda_multicorrelator::Carrier_wipeoff_multicorrelator_resampler_cuda( code_phase_step_chips, rem_code_phase_chips, n_correlators, - signal_length_samples, - rem_carrier_phase_in_rad, - phase_step_rad); + signal_length_samples); gpuErrchk(cudaPeekAtLastError()); gpuErrchk(cudaStreamSynchronize(stream1)); @@ -355,28 +370,41 @@ cuda_multicorrelator::cuda_multicorrelator() d_shifts_samples = NULL; d_shifts_chips = NULL; d_corr_out = NULL; + d_sig_in_cpu = NULL; + d_corr_out_cpu = NULL; + stream1 = 0; threadsPerBlock = 0; blocksPerGrid = 0; d_code_length_chips = 0; + selected_gps_device = 0; + num_gpu_devices = 0; + selected_device = 0; } bool cuda_multicorrelator::free_cuda() { - // Free device global memory - if (d_sig_in != NULL) cudaFree(d_sig_in); + cudaSetDevice(selected_gps_device); + if (stream1 != 0) cudaStreamSynchronize(stream1); + // Free device global memory. d_sig_in and d_corr_out are device aliases of + // host-mapped buffers owned by the caller (see set_input_output_vectors), so + // they must NOT be released here: the caller frees them with cudaFreeHost(). if (d_nco_in != NULL) cudaFree(d_nco_in); if (d_sig_doppler_wiped != NULL) cudaFree(d_sig_doppler_wiped); if (d_local_codes_in != NULL) cudaFree(d_local_codes_in); - if (d_corr_out != NULL) cudaFree(d_corr_out); if (d_shifts_samples != NULL) cudaFree(d_shifts_samples); if (d_shifts_chips != NULL) cudaFree(d_shifts_chips); - // Reset the device and exit - // cudaDeviceReset causes the driver to clean up all state. While - // not mandatory in normal operation, it is good practice. It is also - // needed to ensure correct operation when the application is being - // profiled. Calling cudaDeviceReset causes all profile data to be - // flushed before the application exits - cudaDeviceReset(); + d_sig_in = NULL; + d_corr_out = NULL; + d_nco_in = NULL; + d_sig_doppler_wiped = NULL; + d_local_codes_in = NULL; + d_shifts_samples = NULL; + d_shifts_chips = NULL; + if (stream1 != 0) cudaStreamDestroy(stream1); + stream1 = 0; + // Do NOT call cudaDeviceReset() here: this object is per-channel and a + // reset would tear down the context under every other channel that is + // still tracking. The process-wide reset is done once in main(). return true; } diff --git a/src/algorithms/tracking/libs/cuda_multicorrelator.h b/src/algorithms/tracking/libs/cuda_multicorrelator.h index 561adf2df..474ea5c54 100644 --- a/src/algorithms/tracking/libs/cuda_multicorrelator.h +++ b/src/algorithms/tracking/libs/cuda_multicorrelator.h @@ -48,8 +48,8 @@ struct GPU_Complex float i; CUDA_CALLABLE_MEMBER_DEVICE GPU_Complex() {}; CUDA_CALLABLE_MEMBER_DEVICE GPU_Complex(float a, float b) : r(a), i(b) {} - CUDA_CALLABLE_MEMBER_DEVICE float magnitude2(void) { return r * r + i * i; } - CUDA_CALLABLE_MEMBER_DEVICE GPU_Complex operator*(const GPU_Complex& a) + CUDA_CALLABLE_MEMBER_DEVICE float magnitude2(void) const { return r * r + i * i; } + CUDA_CALLABLE_MEMBER_DEVICE GPU_Complex operator*(const GPU_Complex& a) const { #ifdef __CUDACC__ return GPU_Complex(__fmul_rn(r, a.r) - __fmul_rn(i, a.i), __fmul_rn(i, a.r) + __fmul_rn(r, a.i)); @@ -57,7 +57,7 @@ struct GPU_Complex return GPU_Complex(r * a.r - i * a.i, i * a.r + r * a.i); #endif } - CUDA_CALLABLE_MEMBER_DEVICE GPU_Complex operator+(const GPU_Complex& a) + CUDA_CALLABLE_MEMBER_DEVICE GPU_Complex operator+(const GPU_Complex& a) const { return GPU_Complex(r + a.r, i + a.i); } @@ -90,15 +90,15 @@ struct GPU_Complex_Short float r; float i; CUDA_CALLABLE_MEMBER_DEVICE GPU_Complex_Short(short int a, short int b) : r(a), i(b) {} - CUDA_CALLABLE_MEMBER_DEVICE float magnitude2(void) + CUDA_CALLABLE_MEMBER_DEVICE float magnitude2(void) const { return r * r + i * i; } - CUDA_CALLABLE_MEMBER_DEVICE GPU_Complex_Short operator*(const GPU_Complex_Short& a) + CUDA_CALLABLE_MEMBER_DEVICE GPU_Complex_Short operator*(const GPU_Complex_Short& a) const { return GPU_Complex_Short(r * a.r - i * a.i, i * a.r + r * a.i); } - CUDA_CALLABLE_MEMBER_DEVICE GPU_Complex_Short operator+(const GPU_Complex_Short& a) + CUDA_CALLABLE_MEMBER_DEVICE GPU_Complex_Short operator+(const GPU_Complex_Short& a) const { return GPU_Complex_Short(r + a.r, i + a.i); } diff --git a/src/core/receiver/CMakeLists.txt b/src/core/receiver/CMakeLists.txt index a046ad8b0..7c7b35e6a 100644 --- a/src/core/receiver/CMakeLists.txt +++ b/src/core/receiver/CMakeLists.txt @@ -163,6 +163,7 @@ if(ENABLE_CUDA) target_include_directories(core_receiver PUBLIC ${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES} ) + target_link_libraries(core_receiver PUBLIC CUDA::cudart) else() target_link_libraries(core_receiver PUBLIC ${CUDA_LIBRARIES} diff --git a/src/main/CMakeLists.txt b/src/main/CMakeLists.txt index fed23b8ba..4a85d296e 100644 --- a/src/main/CMakeLists.txt +++ b/src/main/CMakeLists.txt @@ -41,6 +41,7 @@ if(ENABLE_CUDA) target_include_directories(gnss-sdr PUBLIC ${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES} ) + target_link_libraries(gnss-sdr PUBLIC CUDA::cudart) else() target_link_libraries(gnss-sdr PUBLIC ${CUDA_LIBRARIES} diff --git a/tests/benchmarks/CMakeLists.txt b/tests/benchmarks/CMakeLists.txt index 831b5dc19..94ae16f64 100644 --- a/tests/benchmarks/CMakeLists.txt +++ b/tests/benchmarks/CMakeLists.txt @@ -117,6 +117,8 @@ add_benchmark(benchmark_crypto core_libs Boost::headers ${EXTRA_BENCHMARK_DEPEND add_benchmark(benchmark_detector core_system_parameters ${EXTRA_BENCHMARK_DEPENDENCIES}) add_benchmark(benchmark_preamble core_system_parameters ${EXTRA_BENCHMARK_DEPENDENCIES}) add_benchmark(benchmark_reed_solomon core_system_parameters ${EXTRA_BENCHMARK_DEPENDENCIES}) +# CPU baseline of the PCPS acquisition grid; adds the CUDA engine when ENABLE_CUDA=ON +add_benchmark(benchmark_pcps_grid algorithms_libs acquisition_libs Gnuradio::fft Volk::volk Volkgnsssdr::volkgnsssdr ${EXTRA_BENCHMARK_DEPENDENCIES}) if(has_std_plus_void) target_compile_definitions(benchmark_detector PRIVATE -DCOMPILER_HAS_STD_PLUS_VOID=1) diff --git a/tests/benchmarks/benchmark_pcps_grid.cc b/tests/benchmarks/benchmark_pcps_grid.cc new file mode 100644 index 000000000..d95e2efa9 --- /dev/null +++ b/tests/benchmarks/benchmark_pcps_grid.cc @@ -0,0 +1,188 @@ +/*! + * \file benchmark_pcps_grid.cc + * \brief Benchmark of the PCPS acquisition grid: CPU (volk + gr::fft) baseline + * versus the CUDA engine, over a sweep of FFT sizes and Doppler bins. + * \author Phillip Vu, 2026. phillipvu(at)users.noreply.github.com + * + * Run e.g.: + * ./benchmark_pcps_grid --benchmark_counters_tabular=true + * ./benchmark_pcps_grid --benchmark_filter='cuda' --benchmark_repetitions=5 + * + * The "grid_cells/s" counter is (fft_size x bins) per second. It is the number of + * code-phase x Doppler hypotheses are evaluated per second. "dwells/s" is the + * number of complete search grids per second, which is the figure that decides + * how many channels can be in acquisition at once in real time. + * + * ----------------------------------------------------------------------------- + * + * GNSS-SDR is a Global Navigation Satellite System software-defined receiver. + * This file is part of GNSS-SDR. + * + * Copyright (C) 2010-2026 (see AUTHORS file for a list of contributors) + * SPDX-License-Identifier: GPL-3.0-or-later + * + * ----------------------------------------------------------------------------- + */ + +#include "gnss_sdr_fft.h" +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#if CUDA_GPU_ACCEL +#include "cuda_pcps_engine.h" +#endif + +namespace +{ +using cvec = volk_gnsssdr::vector>; +using fvec = volk_gnsssdr::vector; +constexpr double TWO_PI_D = 6.283185307179586; + +// Random complex vectors: the arithmetic cost does not depend on the content. +cvec random_cvec(size_t n, uint32_t seed) +{ + std::mt19937 rng(seed); + std::normal_distribution dist(0.0F, 1.0F); + cvec v(n); + for (auto& x : v) + { + x = std::complex(dist(rng), dist(rng)); + } + return v; +} + +std::vector make_wipeoffs(uint32_t fft_size, uint32_t bins, double fs) +{ + std::vector w(bins, cvec(fft_size)); + for (uint32_t k = 0; k < bins; k++) + { + const double doppler = -5000.0 + 10000.0 * static_cast(k) / static_cast(bins); + const float step = static_cast(TWO_PI_D * doppler / fs); + std::array phase{}; + volk_gnsssdr_s32f_sincos_32fc(w[k].data(), -step, phase.data(), fft_size); + } + return w; +} + +void set_counters(benchmark::State& state, uint32_t fft_size, uint32_t bins) +{ + state.counters["fft_size"] = fft_size; + state.counters["bins"] = bins; + state.counters["dwells/s"] = benchmark::Counter(1, benchmark::Counter::kIsIterationInvariantRate); + state.counters["grid_cells/s"] = benchmark::Counter(static_cast(fft_size) * bins, benchmark::Counter::kIsIterationInvariantRate); +} + + +// -------------------------------------------------------------------------- +// CPU baseline: same operation sequence as pcps_acquisition::doppler_grid_cpu() +// -------------------------------------------------------------------------- +void bm_pcps_grid_cpu(benchmark::State& state) +{ + const auto fft_size = static_cast(state.range(0)); + const auto bins = static_cast(state.range(1)); + const double fs = 4.0e6; + + auto fft_fwd = gnss_fft_fwd_make_unique(fft_size); + auto fft_rev = gnss_fft_rev_make_unique(fft_size); + const cvec in = random_cvec(fft_size, 1U); + const cvec fft_codes = random_cvec(fft_size, 2U); + const std::vector wipeoffs = make_wipeoffs(fft_size, bins, fs); + std::vector magnitude(bins, fvec(fft_size)); + + for (auto _ : state) + { + for (uint32_t k = 0; k < bins; k++) + { + volk_32fc_x2_multiply_32fc(fft_fwd->get_inbuf(), in.data(), wipeoffs[k].data(), fft_size); + fft_fwd->execute(); + volk_32fc_x2_multiply_32fc(fft_rev->get_inbuf(), fft_fwd->get_outbuf(), fft_codes.data(), fft_size); + fft_rev->execute(); + volk_32fc_magnitude_squared_32f(magnitude[k].data(), fft_rev->get_outbuf(), fft_size); + } + benchmark::DoNotOptimize(magnitude[0].data()); + benchmark::ClobberMemory(); + } + set_counters(state, fft_size, bins); +} + + +#if CUDA_GPU_ACCEL +// -------------------------------------------------------------------------- +// CUDA engine, including host<->device transfers (what the receiver sees) +// -------------------------------------------------------------------------- +void bm_pcps_grid_cuda(benchmark::State& state) +{ + const auto fft_size = static_cast(state.range(0)); + const auto bins = static_cast(state.range(1)); + const double fs = 4.0e6; + + CudaPcpsEngine gpu(fft_size, fft_size, bins); + if (!gpu.is_valid()) + { + state.SkipWithError(gpu.last_error().c_str()); + return; + } + const cvec in = random_cvec(fft_size, 1U); + const cvec fft_codes = random_cvec(fft_size, 2U); + const std::vector wipeoffs = make_wipeoffs(fft_size, bins, fs); + std::vector*> wipe_rows(bins); + for (uint32_t k = 0; k < bins; k++) wipe_rows[k] = wipeoffs[k].data(); + std::vector magnitude(bins, fvec(fft_size)); + std::vector out_rows(bins); + for (uint32_t k = 0; k < bins; k++) out_rows[k] = magnitude[k].data(); + + if (!gpu.set_doppler_wipeoffs(CudaPcpsEngine::MAIN_GRID, wipe_rows.data(), bins) || !gpu.set_fft_codes(fft_codes.data())) + { + state.SkipWithError(gpu.last_error().c_str()); + return; + } + // Warm-up: cuFFT plan creation and first-launch costs are one-off in the receiver + if (!gpu.compute_grid(in.data(), CudaPcpsEngine::MAIN_GRID, bins, 0, false, out_rows.data())) + { + state.SkipWithError(gpu.last_error().c_str()); + return; + } + + for (auto _ : state) + { + if (!gpu.compute_grid(in.data(), CudaPcpsEngine::MAIN_GRID, bins, 0, false, out_rows.data())) + { + state.SkipWithError(gpu.last_error().c_str()); + break; + } + benchmark::DoNotOptimize(magnitude[0].data()); + } + set_counters(state, fft_size, bins); + state.SetLabel(gpu.device_name()); +} +#endif +} // namespace + + +// 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}); + } + } +} + +BENCHMARK(bm_pcps_grid_cpu)->Apply(grid_args)->Unit(benchmark::kMicrosecond)->UseRealTime(); +#if CUDA_GPU_ACCEL +BENCHMARK(bm_pcps_grid_cuda)->Apply(grid_args)->Unit(benchmark::kMicrosecond)->UseRealTime(); +#endif + +BENCHMARK_MAIN(); diff --git a/tests/test_main.cc b/tests/test_main.cc index c2a7fc718..896c3b335 100644 --- a/tests/test_main.cc +++ b/tests/test_main.cc @@ -252,6 +252,8 @@ private: #endif #if CUDA_BLOCKS_TEST +#include "unit-tests/signal-processing-blocks/acquisition/cuda_pcps_engine_test.cc" +#include "unit-tests/signal-processing-blocks/acquisition/gps_l1_ca_pcps_acquisition_cuda_test.cc" #include "unit-tests/signal-processing-blocks/tracking/gpu_multicorrelator_test.cc" #endif 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 new file mode 100644 index 000000000..c885580ba --- /dev/null +++ b/tests/unit-tests/signal-processing-blocks/acquisition/cuda_pcps_engine_test.cc @@ -0,0 +1,353 @@ +/*! + * \file cuda_pcps_engine_test.cc + * \brief Checks that the CUDA PCPS grid matches the CPU (volk + gr::fft) grid. + * \author Phillip Vu, 2026. phillipvu(at)users.noreply.github.com + * + * ----------------------------------------------------------------------------- + * + * GNSS-SDR is a Global Navigation Satellite System software-defined receiver. + * This file is part of GNSS-SDR. + * + * Copyright (C) 2010-2026 (see AUTHORS file for a list of contributors) + * SPDX-License-Identifier: GPL-3.0-or-later + * + * ----------------------------------------------------------------------------- + */ + +#include "GPS_L1_CA.h" +#include "MATH_CONSTANTS.h" +#include "cuda_pcps_engine.h" +#include "gnss_sdr_fft.h" +#include "gps_sdr_signal_replica.h" +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + + +namespace +{ +using cvec = volk_gnsssdr::vector>; +using fvec = volk_gnsssdr::vector; + +struct PcpsScenario +{ + int64_t fs_in{4000000}; + uint32_t sampled_ms{1}; + bool bit_transition{false}; + int32_t doppler_max{5000}; + int32_t doppler_step{250}; + uint32_t prn{1}; + uint32_t delay_samples{524}; + float doppler_hz{1680.0F}; + float noise_sigma{0.5F}; + + uint32_t consumed_samples() const { return static_cast(fs_in / 1000) * sampled_ms * (bit_transition ? 2U : 1U); } + uint32_t fft_size() const { return consumed_samples(); } + uint32_t effective_fft_size() const { return bit_transition ? fft_size() / 2U : fft_size(); } + uint32_t num_bins() const { return static_cast(std::ceil(2.0 * doppler_max / doppler_step)); } + uint32_t offset() const { return bit_transition ? effective_fft_size() : 0U; } +}; + + +// Reference implementation: identical operation sequence to +// pcps_acquisition::doppler_grid_cpu(), but standalone. +class CpuPcpsReference +{ +public: + explicit CpuPcpsReference(const PcpsScenario& sc) + : sc_(sc), + N_(sc.fft_size()), + E_(sc.effective_fft_size()), + fft_fwd_(gnss_fft_fwd_make_unique(N_)), + fft_rev_(gnss_fft_rev_make_unique(N_)), + fft_codes_(N_), + tmp_(E_) + { + // Doppler wipe-off grid + wipeoffs_.resize(sc.num_bins(), cvec(N_)); + for (uint32_t k = 0; k < sc.num_bins(); k++) + { + const int32_t doppler = -sc.doppler_max + sc.doppler_step * static_cast(k); + const float phase_step_rad = static_cast(TWO_PI) * static_cast(doppler) / static_cast(sc.fs_in); + std::array phase{}; + volk_gnsssdr_s32f_sincos_32fc(wipeoffs_[k].data(), -phase_step_rad, phase.data(), N_); + } + // conj(FFT(code)), zero padded exactly like pcps_acquisition::set_local_code. + // The adapter tiles one code period sampled_ms times before handing it over. + const uint32_t samples_per_code = static_cast(sc.fs_in / 1000); + cvec period(samples_per_code); + gps_l1_ca_code_gen_complex_sampled(own::span>(period.data(), period.size()), sc.prn, static_cast(sc.fs_in), 0); + cvec code(samples_per_code * sc.sampled_ms); + for (uint32_t r = 0; r < sc.sampled_ms; r++) + { + std::copy(period.begin(), period.end(), code.begin() + static_cast(r * samples_per_code)); + } + if (sc.bit_transition) + { + const uint32_t half = N_ / 2; + std::fill_n(fft_fwd_->get_inbuf(), half, std::complex(0.0F, 0.0F)); + std::copy(code.data(), code.data() + half, fft_fwd_->get_inbuf() + half); + } + else + { + std::copy(code.data(), code.data() + N_, fft_fwd_->get_inbuf()); + } + fft_fwd_->execute(); + volk_32fc_conjugate_32fc(fft_codes_.data(), fft_fwd_->get_outbuf(), N_); + magnitude_.resize(sc.num_bins(), fvec(E_)); + } + + const cvec& fft_codes() const { return fft_codes_; } + const std::vector& wipeoffs() const { return wipeoffs_; } + std::vector& magnitude() { return magnitude_; } + + void doppler_grid(const std::complex* in, bool accumulate) + { + for (uint32_t k = 0; k < sc_.num_bins(); k++) + { + volk_32fc_x2_multiply_32fc(fft_fwd_->get_inbuf(), in, wipeoffs_[k].data(), N_); + fft_fwd_->execute(); + volk_32fc_x2_multiply_32fc(fft_rev_->get_inbuf(), fft_fwd_->get_outbuf(), fft_codes_.data(), N_); + fft_rev_->execute(); + if (!accumulate) + { + volk_32fc_magnitude_squared_32f(magnitude_[k].data(), fft_rev_->get_outbuf() + sc_.offset(), E_); + } + else + { + volk_32fc_magnitude_squared_32f(tmp_.data(), fft_rev_->get_outbuf() + sc_.offset(), E_); + volk_32f_x2_add_32f(magnitude_[k].data(), magnitude_[k].data(), tmp_.data(), E_); + } + } + } + +private: + PcpsScenario sc_; + uint32_t N_; + uint32_t E_; + std::unique_ptr fft_fwd_; + std::unique_ptr fft_rev_; + cvec fft_codes_; + fvec tmp_; + std::vector wipeoffs_; + std::vector magnitude_; +}; + + +// Synthetic GPS L1 C/A signal: delayed, Doppler-shifted code + AWGN +cvec make_signal(const PcpsScenario& sc, uint32_t seed) +{ + const uint32_t N = sc.fft_size(); + const uint32_t samples_per_code = static_cast(sc.fs_in / 1000); + cvec code(samples_per_code); + gps_l1_ca_code_gen_complex_sampled(own::span>(code.data(), code.size()), sc.prn, static_cast(sc.fs_in), 0); + + std::mt19937 rng(seed); + std::normal_distribution noise(0.0F, sc.noise_sigma); + cvec sig(N); + const float phase_step = static_cast(TWO_PI) * sc.doppler_hz / static_cast(sc.fs_in); + for (uint32_t n = 0; n < N; n++) + { + const uint32_t idx = (n + samples_per_code - (sc.delay_samples % samples_per_code)) % samples_per_code; + const std::complex carrier(std::cos(phase_step * static_cast(n)), std::sin(phase_step * static_cast(n))); + sig[n] = code[idx] * carrier + std::complex(noise(rng), noise(rng)); + } + return sig; +} + + +struct GridPeak +{ + uint32_t bin{0}; + uint32_t index{0}; + float value{0.0F}; +}; + +GridPeak find_peak(const std::vector& grid, uint32_t bins, uint32_t E) +{ + GridPeak p; + for (uint32_t k = 0; k < bins; k++) + { + for (uint32_t i = 0; i < E; i++) + { + if (grid[k][i] > p.value) + { + p.value = grid[k][i]; + p.bin = k; + p.index = i; + } + } + } + return p; +} + +// Max |gpu - cpu| over the grid, normalized by the grid maximum +double max_relative_error(const std::vector& a, const std::vector& b, uint32_t bins, uint32_t E) +{ + double max_abs = 0.0; + double max_val = 0.0; + for (uint32_t k = 0; k < bins; k++) + { + for (uint32_t i = 0; i < E; i++) + { + max_abs = std::max(max_abs, static_cast(std::fabs(a[k][i] - b[k][i]))); + max_val = std::max(max_val, static_cast(std::fabs(a[k][i]))); + } + } + return max_val > 0.0 ? max_abs / max_val : max_abs; +} + + +void run_parity_case(const PcpsScenario& sc, uint32_t dwells) +{ + const uint32_t N = sc.fft_size(); + const uint32_t E = sc.effective_fft_size(); + const uint32_t bins = sc.num_bins(); + + CpuPcpsReference cpu(sc); + CudaPcpsEngine gpu(N, E, bins); + ASSERT_TRUE(gpu.is_valid()) << gpu.last_error(); + + std::vector*> wipe_rows(bins); + for (uint32_t k = 0; k < bins; k++) + { + wipe_rows[k] = cpu.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(); + + std::vector gpu_grid(bins, fvec(E)); + std::vector out_rows(bins); + for (uint32_t k = 0; k < bins; k++) + { + out_rows[k] = gpu_grid[k].data(); + } + + for (uint32_t d = 0; d < dwells; d++) + { + const cvec sig = make_signal(sc, 1234U + d); + const bool accumulate = (d != 0); + cpu.doppler_grid(sig.data(), accumulate); + ASSERT_TRUE(gpu.compute_grid(sig.data(), CudaPcpsEngine::MAIN_GRID, bins, sc.offset(), accumulate, out_rows.data())) << gpu.last_error(); + } + + const double rel_err = max_relative_error(cpu.magnitude(), gpu_grid, bins, E); + EXPECT_LT(rel_err, 2e-4) << "GPU grid deviates from the CPU grid"; + + const GridPeak pc = find_peak(cpu.magnitude(), bins, E); + const GridPeak pg = find_peak(gpu_grid, bins, E); + EXPECT_EQ(pc.bin, pg.bin) << "Doppler bin of the peak differs"; + EXPECT_EQ(pc.index, pg.index) << "Code phase of the peak differs"; + + // And the peak must sit where the synthetic signal was put + const int32_t expected_bin = (sc.doppler_max + static_cast(std::lround(sc.doppler_hz))) / sc.doppler_step; + EXPECT_NEAR(static_cast(pg.bin), static_cast(expected_bin), 1.0); + const uint32_t samples_per_code = static_cast(sc.fs_in / 1000); + EXPECT_EQ(pg.index % samples_per_code, sc.delay_samples % samples_per_code); + + std::cout << " fft_size=" << N << " bins=" << bins << " dwells=" << dwells + << " max rel. error=" << rel_err << " peak@(bin " << pg.bin << ", idx " << pg.index << ")\n"; +} +} // namespace + + +TEST(CudaPcpsEngineTest, MatchesCpuReferenceSingleDwell) +{ + PcpsScenario sc; + run_parity_case(sc, 1); +} + + +TEST(CudaPcpsEngineTest, MatchesCpuReferenceNonCoherentAccumulation) +{ + PcpsScenario sc; + sc.noise_sigma = 2.0F; + run_parity_case(sc, 4); +} + + +TEST(CudaPcpsEngineTest, MatchesCpuReferenceBitTransitionMode) +{ + PcpsScenario sc; + sc.bit_transition = true; + run_parity_case(sc, 1); +} + + +TEST(CudaPcpsEngineTest, MatchesCpuReferenceLongCoherentIntegration) +{ + PcpsScenario sc; + sc.sampled_ms = 4; + sc.doppler_step = 100; + run_parity_case(sc, 1); +} + + +TEST(CudaPcpsEngineTest, ReportsInvalidUsage) +{ + PcpsScenario sc; + CudaPcpsEngine gpu(sc.fft_size(), sc.effective_fft_size(), sc.num_bins()); + ASSERT_TRUE(gpu.is_valid()) << gpu.last_error(); + cvec sig = make_signal(sc, 7U); + fvec row(sc.effective_fft_size()); + float* rows[1] = {row.data()}; + // No wipe-offs uploaded yet + EXPECT_FALSE(gpu.compute_grid(sig.data(), CudaPcpsEngine::MAIN_GRID, 1, 0, false, rows)); + EXPECT_FALSE(gpu.last_error().empty()); +} + + +TEST(CudaPcpsEngineTest, MeasureExecutionTime) +{ + // Quick CPU vs GPU timing on the default scenario. The Google Benchmark + // target benchmark_pcps_grid gives the full sweep. + PcpsScenario sc; + const uint32_t N = sc.fft_size(); + const uint32_t E = sc.effective_fft_size(); + const uint32_t bins = sc.num_bins(); + const int iterations = 200; + + CpuPcpsReference cpu(sc); + CudaPcpsEngine gpu(N, E, bins); + ASSERT_TRUE(gpu.is_valid()) << gpu.last_error(); + std::vector*> wipe_rows(bins); + for (uint32_t k = 0; k < bins; k++) wipe_rows[k] = cpu.wipeoffs()[k].data(); + ASSERT_TRUE(gpu.set_doppler_wipeoffs(CudaPcpsEngine::MAIN_GRID, wipe_rows.data(), bins)); + ASSERT_TRUE(gpu.set_fft_codes(cpu.fft_codes().data())); + std::vector gpu_grid(bins, fvec(E)); + std::vector out_rows(bins); + for (uint32_t k = 0; k < bins; k++) out_rows[k] = gpu_grid[k].data(); + const cvec sig = make_signal(sc, 99U); + + // warm-up (cuFFT plan creation, first-launch overhead) + ASSERT_TRUE(gpu.compute_grid(sig.data(), CudaPcpsEngine::MAIN_GRID, bins, sc.offset(), false, out_rows.data())); + cpu.doppler_grid(sig.data(), false); + + auto t0 = std::chrono::steady_clock::now(); + for (int i = 0; i < iterations; i++) + { + cpu.doppler_grid(sig.data(), false); + } + auto t1 = std::chrono::steady_clock::now(); + for (int i = 0; i < iterations; i++) + { + gpu.compute_grid(sig.data(), CudaPcpsEngine::MAIN_GRID, bins, sc.offset(), false, out_rows.data()); + } + auto t2 = std::chrono::steady_clock::now(); + + const double cpu_us = std::chrono::duration(t1 - t0).count() / iterations; + const double gpu_us = std::chrono::duration(t2 - t1).count() / iterations; + std::cout << "PCPS grid, fft_size=" << N << ", " << bins << " Doppler bins: CPU " << cpu_us + << " us/dwell, GPU (" << gpu.device_name() << ") " << gpu_us << " us/dwell, speedup x" + << cpu_us / gpu_us << "\n"; +} 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 new file mode 100644 index 000000000..127f05206 --- /dev/null +++ b/tests/unit-tests/signal-processing-blocks/acquisition/gps_l1_ca_pcps_acquisition_cuda_test.cc @@ -0,0 +1,256 @@ +/*! + * \file gps_l1_ca_pcps_acquisition_cuda_test.cc + * \brief End-to-end test of GPS_L1_CA_PCPS_Acquisition with use_cuda=true on a + * real capture, compared against the CPU path on the same data. + * \author Phillip Vu, 2026. phillipvu(at)users.noreply.github.com + * + * ----------------------------------------------------------------------------- + * + * GNSS-SDR is a Global Navigation Satellite System software-defined receiver. + * This file is part of GNSS-SDR. + * + * Copyright (C) 2010-2026 (see AUTHORS file for a list of contributors) + * SPDX-License-Identifier: GPL-3.0-or-later + * + * ----------------------------------------------------------------------------- + */ + +#include "GPS_L1_CA.h" +#include "concurrent_queue.h" +#include "gnss_block_interface.h" +#include "gnss_sdr_filesystem.h" +#include "gnss_synchro.h" +#include "in_memory_configuration.h" +#include "pcps_acquisition_adapter.h" +#include "test_flags.h" +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#if USE_GLOG_AND_GFLAGS +#include +#else +#include +#endif +#if HAS_GENERIC_LAMBDA +#else +#include +#endif +#if PMT_USES_BOOST_ANY +namespace wht = boost; +#else +namespace wht = std; +#endif + + +// ######## GNURADIO BLOCK MESSAGE RECEIVER ######### +class GpsL1CaPcpsAcquisitionCudaTest_msg_rx; + +using GpsL1CaPcpsAcquisitionCudaTest_msg_rx_sptr = gnss_shared_ptr; + +GpsL1CaPcpsAcquisitionCudaTest_msg_rx_sptr GpsL1CaPcpsAcquisitionCudaTest_msg_rx_make(); + +class GpsL1CaPcpsAcquisitionCudaTest_msg_rx : public gr::block +{ +private: + friend GpsL1CaPcpsAcquisitionCudaTest_msg_rx_sptr GpsL1CaPcpsAcquisitionCudaTest_msg_rx_make(); + void msg_handler_channel_events(const pmt::pmt_t &msg); + GpsL1CaPcpsAcquisitionCudaTest_msg_rx(); + +public: + int rx_message{0}; +}; + + +GpsL1CaPcpsAcquisitionCudaTest_msg_rx_sptr GpsL1CaPcpsAcquisitionCudaTest_msg_rx_make() +{ + return GpsL1CaPcpsAcquisitionCudaTest_msg_rx_sptr(new GpsL1CaPcpsAcquisitionCudaTest_msg_rx()); +} + + +void GpsL1CaPcpsAcquisitionCudaTest_msg_rx::msg_handler_channel_events(const pmt::pmt_t &msg) +{ + try + { + int64_t message = pmt::to_long(msg); + rx_message = message; + } + catch (const wht::bad_any_cast &e) + { + LOG(WARNING) << "msg_handler_channel_events Bad any_cast: " << e.what(); + rx_message = 0; + } +} + + +GpsL1CaPcpsAcquisitionCudaTest_msg_rx::GpsL1CaPcpsAcquisitionCudaTest_msg_rx() + : gr::block("GpsL1CaPcpsAcquisitionCudaTest_msg_rx", gr::io_signature::make(0, 0, 0), gr::io_signature::make(0, 0, 0)) +{ + this->message_port_register_in(pmt::mp("events")); + this->set_msg_handler(pmt::mp("events"), +#if HAS_GENERIC_LAMBDA + [this](auto &&PH1) { msg_handler_channel_events(std::forward(PH1)); }); +#else +#if USE_BOOST_BIND_PLACEHOLDERS + boost::bind(&GpsL1CaPcpsAcquisitionCudaTest_msg_rx::msg_handler_channel_events, this, boost::placeholders::_1)); +#else + boost::bind(&GpsL1CaPcpsAcquisitionCudaTest_msg_rx::msg_handler_channel_events, this, _1)); +#endif +#endif +} + + +// ########################################################### + +class GpsL1CaPcpsAcquisitionCudaTest : public ::testing::Test +{ +protected: + struct Outcome + { + int message{0}; + double doppler_hz{0.0}; + double delay_samples{0.0}; + double elapsed_us{0.0}; + }; + + std::shared_ptr make_config(bool use_cuda, bool two_steps) const; + Outcome run_once(bool use_cuda, bool two_steps = false); + + unsigned int doppler_max{5000}; + unsigned int doppler_step{100}; +}; + + +std::shared_ptr GpsL1CaPcpsAcquisitionCudaTest::make_config(bool use_cuda, bool two_steps) const +{ + auto config = std::make_shared(); + config->set_property("GNSS-SDR.internal_fs_sps", "4000000"); + 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.dump", "false"); + config->set_property("Acquisition_1C.threshold", "0.001"); + config->set_property("Acquisition_1C.doppler_max", std::to_string(doppler_max)); + config->set_property("Acquisition_1C.doppler_step", std::to_string(doppler_step)); + config->set_property("Acquisition_1C.repeat_satellite", "false"); + config->set_property("Acquisition_1C.use_cuda", use_cuda ? "true" : "false"); + if (two_steps) + { + config->set_property("Acquisition_1C.make_two_steps", "true"); + config->set_property("Acquisition_1C.second_nbins", "5"); + config->set_property("Acquisition_1C.second_doppler_step", "20"); + } + return config; +} + + +GpsL1CaPcpsAcquisitionCudaTest::Outcome GpsL1CaPcpsAcquisitionCudaTest::run_once(bool use_cuda, bool two_steps) +{ + Outcome out; + auto config = make_config(use_cuda, two_steps); + auto top_block = gr::make_top_block("Acquisition CUDA test"); + + Gnss_Synchro gnss_synchro = Gnss_Synchro(); + gnss_synchro.Channel_ID = 0; + gnss_synchro.System = 'G'; + std::string signal = "1C"; + signal.copy(gnss_synchro.Signal, 2, 0); + gnss_synchro.PRN = 1; + + auto acquisition = std::make_shared(config.get(), "Acquisition_1C", "GPS_L1_CA_PCPS_Acquisition", 1, 0, GPS_1C); + auto msg_rx = GpsL1CaPcpsAcquisitionCudaTest_msg_rx_make(); + + acquisition->set_channel(1); + acquisition->set_gnss_synchro(&gnss_synchro); + acquisition->connect(top_block); + + std::string path = std::string(TEST_PATH); + // The two-step search needs more than the 2 ms of the single-step capture + std::string file = path + (two_steps ? "signal_samples/GSoC_CTTC_capture_2012_07_26_4Msps_4ms.dat" : "signal_samples/GPS_L1_CA_ID_1_Fs_4Msps_2ms.dat"); + gr::blocks::file_source::sptr file_source = gr::blocks::file_source::make(sizeof(gr_complex), file.c_str(), false); + top_block->connect(file_source, 0, acquisition->get_left_block(), 0); + top_block->msg_connect(acquisition->get_right_block(), pmt::mp("events"), msg_rx, pmt::mp("events")); + + acquisition->set_local_code(); + acquisition->reset(); + + const auto start = std::chrono::steady_clock::now(); + top_block->run(); + const auto end = std::chrono::steady_clock::now(); + + out.message = msg_rx->rx_message; + out.doppler_hz = gnss_synchro.Acq_doppler_hz; + out.delay_samples = gnss_synchro.Acq_delay_samples; + out.elapsed_us = std::chrono::duration(end - start).count(); + return out; +} + + +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); +} + + +TEST_F(GpsL1CaPcpsAcquisitionCudaTest /*unused*/, ValidationOfResults /*unused*/) +{ + const double expected_delay_samples = 524; + const double expected_doppler_hz = 1680; + + Outcome gpu; + ASSERT_NO_THROW({ gpu = run_once(true); }) << "Failure running the top_block with use_cuda=true."; + std::cout << "CUDA acquisition: message=" << gpu.message << " Doppler=" << gpu.doppler_hz + << " Hz, delay=" << gpu.delay_samples << " samples, flowgraph time " << gpu.elapsed_us << " us\n"; + + ASSERT_EQ(1, gpu.message) << "Acquisition failure. Expected message: 1=ACQ SUCCESS."; + + const double delay_error_samples = std::abs(expected_delay_samples - gpu.delay_samples); + const auto delay_error_chips = static_cast(delay_error_samples * 1023 / 4000); + const double doppler_error_hz = std::abs(expected_doppler_hz - gpu.doppler_hz); + + EXPECT_LE(doppler_error_hz, 666) << "Doppler error exceeds the expected value: 666 Hz = 2/(3*integration period)"; + EXPECT_LT(delay_error_chips, 0.5) << "Delay error exceeds the expected value: 0.5 chips"; +} + + +TEST_F(GpsL1CaPcpsAcquisitionCudaTest /*unused*/, SameEstimateAsCpu /*unused*/) +{ + Outcome cpu; + Outcome gpu; + ASSERT_NO_THROW({ cpu = run_once(false); }) << "Failure running the top_block with use_cuda=false."; + ASSERT_NO_THROW({ gpu = run_once(true); }) << "Failure running the top_block with use_cuda=true."; + + std::cout << "CPU: Doppler=" << cpu.doppler_hz << " Hz, delay=" << cpu.delay_samples << " samples (" << cpu.elapsed_us << " us)\n"; + std::cout << "CUDA: Doppler=" << gpu.doppler_hz << " Hz, delay=" << gpu.delay_samples << " samples (" << gpu.elapsed_us << " us)\n"; + + ASSERT_EQ(cpu.message, gpu.message); + ASSERT_EQ(1, cpu.message); + // Same grid, same peak search: the estimates must coincide to the bin + EXPECT_NEAR(cpu.doppler_hz, gpu.doppler_hz, static_cast(doppler_step) / 2.0); + EXPECT_NEAR(cpu.delay_samples, gpu.delay_samples, 1.0); +} + + +TEST_F(GpsL1CaPcpsAcquisitionCudaTest /*unused*/, SameEstimateAsCpuMakeTwoStep /*unused*/) +{ + // Exercises the fine-Doppler (STEP2_GRID) path of the engine through the real block + Outcome cpu; + Outcome gpu; + ASSERT_NO_THROW({ cpu = run_once(false, true); }) << "Failure running the top_block with use_cuda=false."; + ASSERT_NO_THROW({ gpu = run_once(true, true); }) << "Failure running the top_block with use_cuda=true."; + + std::cout << "CPU (two steps): Doppler=" << cpu.doppler_hz << " Hz, delay=" << cpu.delay_samples << " samples (" << cpu.elapsed_us << " us)\n"; + std::cout << "CUDA (two steps): Doppler=" << gpu.doppler_hz << " Hz, delay=" << gpu.delay_samples << " samples (" << gpu.elapsed_us << " us)\n"; + + ASSERT_EQ(1, cpu.message) << "CPU acquisition failure. Expected message: 1=ACQ SUCCESS."; + ASSERT_EQ(1, gpu.message) << "CUDA acquisition failure. Expected message: 1=ACQ SUCCESS."; + EXPECT_NEAR(cpu.doppler_hz, gpu.doppler_hz, 20.0); // second_doppler_step + EXPECT_NEAR(cpu.delay_samples, gpu.delay_samples, 1.0); +} diff --git a/utils/scripts/jetson-build.sh b/utils/scripts/jetson-build.sh new file mode 100755 index 000000000..b83a7edb0 --- /dev/null +++ b/utils/scripts/jetson-build.sh @@ -0,0 +1,97 @@ +#!/usr/bin/env bash +# SPDX-License-Identifier: GPL-3.0-or-later +# SPDX-FileCopyrightText: 2026 Phillip Vu <36169227+phillipvu@users.noreply.github.com> +# +# Configure, build, and (optionally) test/benchmark GNSS-SDR with CUDA on an +# NVIDIA Jetson. Companion to docs/JETSON.md. Run from the source tree root: +# +# utils/scripts/jetson-build.sh [--deps] [--tests] [--bench] [--install] [-j N] +# +# --deps apt-get install the Ubuntu build dependencies (needs sudo) +# --tests build and run the CUDA unit tests (run_tests, gtest filter) +# --bench build and run benchmark_pcps_grid (CPU baseline vs GPU) +# --install cmake --install (needs sudo) +# -j N parallel jobs (default: nproc, capped at 6 on <=8 GB boards) + +set -euo pipefail + +SRC_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")/../.." && pwd)" +BUILD_DIR="${BUILD_DIR:-$SRC_DIR/build}" +DO_DEPS=0 +DO_TESTS=0 +DO_BENCH=0 +DO_INSTALL=0 +JOBS="" + +while [[ $# -gt 0 ]]; do + case "$1" in + --deps) DO_DEPS=1 ;; + --tests) DO_TESTS=1 ;; + --bench) DO_BENCH=1 ;; + --install) DO_INSTALL=1 ;; + -j) JOBS="$2"; shift ;; + -j*) JOBS="${1#-j}" ;; + -h|--help) sed -n '5,15p' "$0"; exit 0 ;; + *) echo "unknown option: $1" >&2; exit 2 ;; + esac + shift +done + +export PATH="/usr/local/cuda/bin:$PATH" +export LD_LIBRARY_PATH="/usr/local/cuda/lib64:${LD_LIBRARY_PATH:-}" + +if [[ -z "$JOBS" ]]; then + JOBS="$(nproc)" + MEM_GB=$(awk '/MemTotal/ {printf "%d", $2/1024/1024}' /proc/meminfo) + if [[ "$MEM_GB" -le 8 && "$JOBS" -gt 6 ]]; then JOBS=6; fi +fi + +echo "== GNSS-SDR CUDA build on $(tr -d '\0' < /proc/device-tree/model 2>/dev/null || hostname)" +echo " source: $SRC_DIR" +echo " build: $BUILD_DIR" +echo " jobs: $JOBS" +command -v nvcc >/dev/null || { echo "nvcc not found: is the CUDA toolkit installed under /usr/local/cuda?" >&2; exit 1; } +nvcc --version | tail -n 2 + +if [[ $DO_DEPS -eq 1 ]]; then + echo "== Installing build dependencies" + sudo apt-get update + sudo apt-get install -y build-essential cmake git pkg-config libboost-dev \ + libboost-date-time-dev libboost-system-dev libboost-filesystem-dev \ + libboost-thread-dev libboost-chrono-dev libboost-serialization-dev \ + libabsl-dev libad9361-dev libarmadillo-dev libblas-dev \ + libgnutls-openssl-dev libgnutls28-dev libgtest-dev liblapack-dev \ + libmatio-dev libpcap-dev libprotobuf-dev libpugixml-dev libssl-dev \ + libuhd-dev gnuradio-dev gr-osmosdr protobuf-compiler python3-mako +fi + +CMAKE_ARGS=(-DENABLE_CUDA=ON) +if [[ $DO_TESTS -eq 1 ]]; then CMAKE_ARGS+=(-DENABLE_UNIT_TESTING=ON); fi +if [[ $DO_BENCH -eq 1 ]]; then CMAKE_ARGS+=(-DENABLE_BENCHMARKS=ON); fi + +echo "== Configuring: cmake ${CMAKE_ARGS[*]}" +cmake -S "$SRC_DIR" -B "$BUILD_DIR" "${CMAKE_ARGS[@]}" + +echo "== Building" +cmake --build "$BUILD_DIR" -j"$JOBS" + +if [[ $DO_TESTS -eq 1 ]]; then + echo "== Running CUDA unit tests" + "$BUILD_DIR/src/tests/run_tests" \ + --gtest_filter='CudaPcpsEngineTest.*:GpsL1CaPcpsAcquisitionCudaTest.*:GpuMulticorrelatorTest.*' +fi + +if [[ $DO_BENCH -eq 1 ]]; then + echo "== Running benchmark_pcps_grid (CPU baseline vs GPU)" + if command -v nvpmodel >/dev/null; then + echo " (for repeatable numbers: sudo nvpmodel -m 0 && sudo jetson_clocks)" + fi + "$BUILD_DIR/tests/benchmarks/benchmark_pcps_grid" --benchmark_counters_tabular=true +fi + +if [[ $DO_INSTALL -eq 1 ]]; then + echo "== Installing" + sudo cmake --install "$BUILD_DIR" +fi + +echo "== Done"