Ускорение генерации спутниковых сигналов: как мы перенесли GNSS-симулятор на GPU и достигли 5-кратного ускорения
История рефакторинга SignalSim — от CPU-референса до производительного CUDA-ускорителя с битовой верификацией
Введение: зачем вообще генерировать «искусственные» спутники?
Современные системы навигации и посадки (включая ГЛОНАСС/ГАЛС) требуют стресс-тестирования в условиях, которые невозможно воспроизвести в полёте: многолучевое распространение, радиопомехи, отказы спутников. Для этого используются программные симуляторы сигналов навигационных систем (GNSS) — такие как SignalSim. Однако классическая реализация на CPU упирается в производительность: генерация 1 секунды сигнала с частотой дискретизации 10 МГц требует обработки 10 млн комплексных отсчётов с модуляцией, несущей и шумом. На практике это означает минуты расчётов вместо реального времени.
В этой статье я расскажу, как мы форкнули SignalSim, создали проект «Имитатор» и перенесли критические пути на GPU с сохранением битовой точности и поддержкой сложных модуляций (BOC, L2C).
Этап 1. Базовая производительность
Поскольку это проект про ускорение работы, сначала нужно измерить скорость CPU версии без ускорения. За исходную точку взят коммит SignalSim bebde2533. Для этого эксперимента и всех последующих генерировался сигнал для 8 спутников GLONASS по данным эфемерид. Результаты замеров оказались следующими: IDE: Visual Studio 2022
| Конфигурация сборки | Скорость создания файла iq сигнала |
|---|---|
| x64/Release, OpenMP включен | 17 Мб/cек |
| x64/Debug, OpenMP выключен | 6.29 Мб/cек |
Этап 2. Инструментация и «золотой эталон»
Первый шаг — не оптимизация, а валидация. Чтобы доверять результатам GPU, нужен надёжный референс. Мы модифицировали класс CSatIfSignal, добавив бинарное логирование промежуточных состояний:
// Магические маркеры для надёжного парсинга
constexpr uint32_t HEADER_MAGIC_INIT = 0xC0DE0006; // Инициализация: длины PRN, атрибуты
constexpr uint32_t HEADER_MAGIC_GETIF = 0xC0DE0007; // Параметры сегмента: амплитуда, шаг кода/фазы
constexpr uint32_t HEADER_MAGIC_SAMPLE = 0xC0DE0008; // Пересчёт на каждом отсчёте
Логгер записывал:
- Исходные последовательности PRN (включая пилотные каналы)
- Текущую фазу несущей в 32-битном формате с фиксированной точкой (
CurIntPhase) - Промежуточные значения:
prn_value,rotate_value, финальный отсчёт - Моменты перехода через границу периода кода (когда
((int)CurChip % DataLength) < DataChip)
Это позволило создать валидатор, который читает трейс и воспроизводит вычисления на GPU, сравнивая результаты с эталоном с точностью до последнего бита.
Этап 3. Архитектурное решение: сегментация вместо пересчёта на каждом отсчёте
Ключевая проблема GNSS-генерации — неоднородность сигнала. Каждые 1 мс (для GPS L1 C/A) или 10 мс (для ГЛОНАСС) происходит:
- Смена навигационного бита
- Сброс фазы кода (возврат к началу псевдослучайной последовательности)
Если обрабатывать каждый отсчёт отдельно, на GPU это приведёт к ветвлению и падению производительности. Мы применили технику гомогенных сегментов:
- До запуска ядра находим границы сегментов, где параметры постоянны:
// Аналитический расчёт границ (не перебор!) double remainder = fmod(current_chip, DataLength); double to_boundary = DataLength - remainder; int steps = ceil(to_boundary / CodeStep - 1e-12); // эпсилон против ошибок округления - Каждый сегмент обрабатывается отдельным запуском ядра с фиксированными
data_signal/pilot_signal - На границе сегмента обновляем сигнал через
SatelliteSignal.GetSatelliteSignal()
Результат: вместо 10 млн условных переходов — 100–200 однородных блоков.
Производительность на данном этапе оказалась следующей:
| Конфигурация сборки | OpenMP | CPU/CUDA | Скорость создания файла iq сигнала | Комментарий |
|---|---|---|---|---|
| x64/Debug | однопоточный | CPU | 6.29 Мб/cек | бейзлайн |
| x64/Debug | однопоточный | CUDA | 3.84 Мб/cек | загрузка GPU 12% |
| x64/Release | многопоточный | CPU | 17 Мб/cек | безйлайн |
| x64/Release | многопоточный | CUDA | 10 Мб/cек |
Итак, первая версия ускорения не ускорила, а замедлила базовую скорость ! Попробуем разобраться в чем причина: запускаем профайлер Visual Studio, видим что по-прежнему много времени уходит на запуск CUDA ядра. А почему ? Запускаем Nsight Systems и видим, что основное время уходит на cudaMalloc/free и компирование данных. Само ядро работает в лучшем случае 1% от всего времени запуска.
Сделаем рефакторинг кода: выделим память для GPU под вектор результата конструкторе генератора сигнала для спутника. Ещё один момент: начиная с CUDA 7 стало возможно делать неблокирующие операции на каждом потоке CPU в дефолтном потоке управления CUDA. Для этого следует включить эту функцию флагом компилятора nvcc --default-stream per-thread. Вместе с этим заменим синхронизацию для всего устройства CUDA на синхронизацию на каждом потоке CPU:
cudaStreamSynchronize(cudaStreamPerThread);
Померяли скорость - почти не изменилась. Оказывается память снова копируется обратно на хост далее по стеку SignalSim. В идеале копироваться обратно на хост должен уже полностью сформированный сигнал. Продолжим избавление от cudaAlloc/cudaMemcpy.
Этап 4. Единый тип комплексного числа для CPU и GPU
Изначально в кодовой базе существовал дублирующий тип complex_number_dev для устройства. Мы устранили это:
class __align__(16) complex_number {
public:
double real, imag;
__host__ __device__ complex_number operator*(const complex_number o) const {
return { real * o.real - imag * o.imag, real * o.imag + imag * o.real };
}
// ... остальные операторы с __host__ __device__
};
Это упростило код, устранило копирование между complex_number ↔ complex_number_dev, а выравнивание __align__(16) обеспечило оптимальный доступ к памяти на GPU.
На этом этапе ускорение составило:
| Конфигурация сборки | OpenMP | CPU/CUDA | Скорость создания файла iq сигнала | Комментарий |
|---|---|---|---|---|
| x64/Release | многопоточный | CPU | 17 Мб/cек | безйлайн |
| x64/Release | многопоточный | CUDA | 16.78 Мб/cек |
Ура, мы достигли исходной скорости бейзлайна !
Этап 5. Шум и накопление каналов на GPU
Раньше шум генерировался на CPU в цикле:
for (i = 0; i < SampleFreq; i++)
NoiseArray[i] = GenerateNoise(1.0);
Мы заменили это на параллельную генерацию через cuRAND:
__global__ void generate_noise_kernel(curandState* states, complex_number* output, int n, float sigma) {
int idx = threadIdx.x + blockIdx.x * blockDim.x;
curandState local_state = states[idx];
float real = curand_normal(&local_state) * sigma;
float imag = curand_normal(&local_state) * sigma;
output[idx] = complex_number(real, imag);
states[idx] = local_state; // сохраняем состояние
}
А накопление сигналов 32 спутников — через отдельное ядро accumulate_channels_kernel, где каждый поток суммирует один отсчёт по всем каналам.
Теперь все основные расчеты происходят на GPU без копирования на хост и обратно. Результат:
| Конфигурация сборки | OpenMP | CPU/CUDA | Скорость создания файла iq сигнала | Комментарий |
|---|---|---|---|---|
| x64/Release | многопоточный | CPU | 17 Мб/cек | безйлайн |
| x64/Release | многопоточный | CUDA | 50 Мб/cек | GPU 40% |
Скорость выросла в 3 раза, но загрузка GPU всего 40%. Отлично! Но еще есть куда расти.
Этап 5. Ускорение расчета границ сегментов
Результаты профилирования показали, что расчет границ сегментов происходит неоптимально, простым перебором от 1 до 10000 на каждой миллисекунде генерации:
std::vector<int> segment_starts;
segment_starts.push_back(0);
double temp_chip = CurChip;
for (int idx = 0; idx < SampleNumber - 1; ++idx) {
int chip_before = static_cast<int>(temp_chip) % DataLength;
double next_chip = temp_chip + CodeStep;
int chip_after = static_cast<int>(next_chip) % DataLength;
if (chip_after < chip_before) {
segment_starts.push_back(idx + 1);
}
temp_chip = next_chip;
}
segment_starts.push_back(SampleNumber);
Но что, если можно вычислить границы, не прибегая к перебору? Да, это возможно! После оптимизации этого алгоритма получили финальную скорость:
| Конфигурация сборки | OpenMP | CPU/CUDA | Скорость создания файла iq сигнала | Комментарий |
|---|---|---|---|---|
| x64/Release | многопоточный | CPU | 17 Мб/cек | безйлайн |
| x64/Release | многопоточный | CUDA | 83 Мб/cек | GPU 63% |
Интересно, что добавление спутников GPS, Beidou, и Galileo не замедлили генерацию: скорость осталась той же. Это означает, что теперь дальнейшее повышение производительности ограничивается неоптимальной работой с GPU, а не CPU.
Результаты и выводы
| Этап | Пропускная способность | Комментарий |
|---|---|---|
| Исходный CPU-код | ~17 Мбайт/с | Бейзлайн |
| Конечный результат | ~83 Мбайт/с |
Ключевые уроки:
- Валидация важнее скорости. Без бинарных трейсов невозможно отладить численные расхождения между CPU/GPU.
- Сегментация побеждает ветвление. Гомогенные блоки — основа производительности на GPU для задач с естественными границами.
- Битовая точность достижима. При копировании логики (включая битовые сдвиги в индексации таблиц) расхождения отсутствуют даже после 10⁹ отсчётов.
- Аллокация и копирование памяти. Главная причина медлительности первой версии ускорения была в частых аллокации и копировании памяти с GPU на хост.
В результате данной работы был получен прототип имитатора для сравнения возможностей ускорения на GPU с другими методами оптимизации. При этом потолок скорости не достигнут: загрузка GPU всего 63%, что, в прочем, более чем достаточно для генерации сигнала в реальном времени.
Алексей Парфёнов, инженер-разработчик систем радионавигации
Челябинск, февраль 2026
