mirror of
https://github.com/gnss-sdr/gnss-sdr
synced 2026-10-15 04:01:49 +00:00
Use American spelling and plain punctuation in CUDA docs and comments
Signed-off-by: Phillip Vu <36169227+phillipvu@users.noreply.github.com>
This commit is contained in:
+3
-3
@@ -184,7 +184,7 @@ All notable changes to GNSS-SDR will be documented in this file.
|
||||
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 initialised. Added
|
||||
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.
|
||||
@@ -435,7 +435,7 @@ All notable changes to GNSS-SDR will be documented in this file.
|
||||
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 honoured and detected
|
||||
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
|
||||
@@ -464,7 +464,7 @@ All notable changes to GNSS-SDR will be documented in this file.
|
||||
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, synchronised only with `__syncthreads()`), fixed 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
|
||||
|
||||
+7
-7
@@ -22,7 +22,7 @@ Ubuntu 22.04, CUDA 12.x) differs only in package versions.
|
||||
|
||||
| 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 initialised. |
|
||||
| **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`
|
||||
@@ -154,8 +154,8 @@ 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
|
||||
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:
|
||||
|
||||
@@ -213,7 +213,7 @@ 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 characterised.
|
||||
throughput; concurrent multi-channel GPU throughput has not been characterized.
|
||||
|
||||
Run the sweep yourself with:
|
||||
|
||||
@@ -243,7 +243,7 @@ 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 initialised the block logs a `WARNING` and continues on
|
||||
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`).
|
||||
|
||||
@@ -272,7 +272,7 @@ Doppler range.
|
||||
|
||||
## 6. Notes on the implementation
|
||||
|
||||
- `src/algorithms/acquisition/libs/cuda_pcps_engine.{h,cu}` — the engine. One
|
||||
- `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
|
||||
@@ -294,6 +294,6 @@ Doppler range.
|
||||
| ----------------------------------------------------------------------------- | -------------------------------------------------------------------------------------------------------------------------------------------- |
|
||||
| `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 initialised (... cudaErrorNoDevice ...)` | The process cannot see the GPU. Check `nvidia-smi`/`tegrastats`, and that the user is in the `video` group on Jetson. |
|
||||
| `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). |
|
||||
|
||||
@@ -224,7 +224,7 @@ void pcps_acquisition::init_cuda_engine()
|
||||
auto engine = std::make_unique<CudaPcpsEngine>(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 initialised (" << engine->last_error()
|
||||
LOG(WARNING) << "CUDA acquisition engine could not be initialized (" << engine->last_error()
|
||||
<< "). Falling back to the CPU implementation.";
|
||||
return;
|
||||
}
|
||||
|
||||
@@ -434,7 +434,7 @@ bool CudaPcpsEngine::set_doppler_wipeoffs(GridId grid, const std::complex<float>
|
||||
}
|
||||
// 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 centre, or the two-step refinement), not per dwell.
|
||||
// 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++)
|
||||
{
|
||||
@@ -481,7 +481,7 @@ bool CudaPcpsEngine::run_pipeline(int grid, uint32_t bins, uint32_t offset, bool
|
||||
e = cudaGetLastError();
|
||||
if (e != cudaSuccess) return p->fail("k_mult_codes", e);
|
||||
|
||||
// Batched inverse FFT (in place, unnormalised like FFTW/gr::fft)
|
||||
// Batched inverse FFT (in place, unnormalized like FFTW/gr::fft)
|
||||
r = cufftExecC2C(p->plan[grid], reinterpret_cast<cufftComplex*>(p->d_batch), reinterpret_cast<cufftComplex*>(p->d_batch), CUFFT_INVERSE);
|
||||
if (r != CUFFT_SUCCESS) return p->fail("cufftExecC2C(inverse)", r);
|
||||
|
||||
@@ -530,7 +530,7 @@ bool CudaPcpsEngine::compute_grid(const std::complex<float>* in, GridId grid, ui
|
||||
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 (synchronises the stream)
|
||||
// Wipe-off, FFT, code multiply, IFFT, magnitude (synchronizes the stream)
|
||||
if (!run_pipeline(grid, bins, offset, accumulate))
|
||||
{
|
||||
return false;
|
||||
|
||||
@@ -8,7 +8,7 @@
|
||||
* ./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, i.e. how many
|
||||
* 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.
|
||||
|
||||
@@ -191,7 +191,7 @@ GridPeak find_peak(const std::vector<fvec>& grid, uint32_t bins, uint32_t E)
|
||||
return p;
|
||||
}
|
||||
|
||||
// Max |gpu - cpu| over the grid, normalised by the grid maximum
|
||||
// Max |gpu - cpu| over the grid, normalized by the grid maximum
|
||||
double max_relative_error(const std::vector<fvec>& a, const std::vector<fvec>& b, uint32_t bins, uint32_t E)
|
||||
{
|
||||
double max_abs = 0.0;
|
||||
|
||||
Reference in New Issue
Block a user