mirrors/gnss-sdr
1
0
Fork 0
mirror of https://github.com/gnss-sdr/gnss-sdr synced 2025-10-27 13:37:38 +00:00
Code Issues Releases Wiki Activity

Merge branch 'next' of https://github.com/gnss-sdr/gnss-sdr into fpga

Browse Source
This commit is contained in:
Marc Majoral
2019-03-20 16:35:55 +01:00
parent a24d26f427 0be8c657de
commit 6da82535ba
171 changed files with 5547 additions and 1200 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
src/algorithms/tracking/adapters/CMakeLists.txt
View File

@@ -62,6 +62,7 @@ set(TRACKING_ADAPTER_SOURCES
glonass_l2_ca_dll_pll_tracking.cc
glonass_l2_ca_dll_pll_c_aid_tracking.cc
beidou_b1i_dll_pll_tracking.cc
beidou_b3i_dll_pll_tracking.cc
${OPT_TRACKING_ADAPTERS_SOURCES}
)
@@ -80,6 +81,7 @@ set(TRACKING_ADAPTER_HEADERS
glonass_l2_ca_dll_pll_tracking.h
glonass_l2_ca_dll_pll_c_aid_tracking.h
beidou_b1i_dll_pll_tracking.h
beidou_b3i_dll_pll_tracking.h
${OPT_TRACKING_ADAPTERS_HEADERS}
)

203
src/algorithms/tracking/adapters/beidou_b3i_dll_pll_tracking.cc Normal file
View File

@@ -0,0 +1,203 @@
/*!
* \file beidou_b3i_dll_pll_tracking.cc
* \brief Implementation of an adapter of a DLL+PLL tracking loop block
* for Beidou B3I to a TrackingInterface
* \author Damian Miralles, 2019. dmiralles2009(at)gmail.com
*
* Code DLL + carrier PLL according to the algorithms described in:
* K.Borre, D.M.Akos, N.Bertelsen, P.Rinder, and S.H.Jensen,
* A Software-Defined GPS and Galileo Receiver. A Single-Frequency
* Approach, Birkhauser, 2007
*
* -------------------------------------------------------------------------
*
* Copyright (C) 2010-2019 (see AUTHORS file for a list of contributors)
*
* GNSS-SDR is a software defined Global Navigation
* Satellite Systems receiver
*
* This file is part of GNSS-SDR.
*
* GNSS-SDR is free software: you can redistribute it and/or modify
* it under the terms of the GNU General Public License as published by
* the Free Software Foundation, either version 3 of the License, or
* (at your option) any later version.
*
* GNSS-SDR is distributed in the hope that it will be useful,
* but WITHOUT ANY WARRANTY; without even the implied warranty of
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
* GNU General Public License for more details.
*
* You should have received a copy of the GNU General Public License
* along with GNSS-SDR. If not, see <https://www.gnu.org/licenses/>.
*
* -------------------------------------------------------------------------
*/
#include "beidou_b3i_dll_pll_tracking.h"
#include "Beidou_B3I.h"
#include "configuration_interface.h"
#include "display.h"
#include "dll_pll_conf.h"
#include "gnss_sdr_flags.h"
#include <glog/logging.h>
using google::LogMessage;
BeidouB3iDllPllTracking::BeidouB3iDllPllTracking(
ConfigurationInterface* configuration, const std::string& role,
unsigned int in_streams, unsigned int out_streams) : role_(role), in_streams_(in_streams), out_streams_(out_streams)
{
Dll_Pll_Conf trk_param = Dll_Pll_Conf();
DLOG(INFO) << "role " << role;
//################# CONFIGURATION PARAMETERS ########################
std::string default_item_type = "gr_complex";
std::string item_type = configuration->property(role + ".item_type", default_item_type);
int fs_in_deprecated = configuration->property("GNSS-SDR.internal_fs_hz", 2048000);
int fs_in = configuration->property("GNSS-SDR.internal_fs_sps", fs_in_deprecated);
trk_param.fs_in = fs_in;
bool dump = configuration->property(role + ".dump", false);
trk_param.dump = dump;
float pll_bw_hz = configuration->property(role + ".pll_bw_hz", 50.0);
if (FLAGS_pll_bw_hz != 0.0) pll_bw_hz = static_cast<float>(FLAGS_pll_bw_hz);
trk_param.pll_bw_hz = pll_bw_hz;
float pll_bw_narrow_hz = configuration->property(role + ".pll_bw_narrow_hz", 20.0);
trk_param.pll_bw_narrow_hz = pll_bw_narrow_hz;
float dll_bw_narrow_hz = configuration->property(role + ".dll_bw_narrow_hz", 2.0);
trk_param.dll_bw_narrow_hz = dll_bw_narrow_hz;
float dll_bw_hz = configuration->property(role + ".dll_bw_hz", 2.0);
if (FLAGS_dll_bw_hz != 0.0) dll_bw_hz = static_cast<float>(FLAGS_dll_bw_hz);
trk_param.dll_bw_hz = dll_bw_hz;
float early_late_space_chips = configuration->property(role + ".early_late_space_chips", 0.5);
trk_param.early_late_space_chips = early_late_space_chips;
float early_late_space_narrow_chips = configuration->property(role + ".early_late_space_narrow_chips", 0.5);
trk_param.early_late_space_narrow_chips = early_late_space_narrow_chips;
std::string default_dump_filename = "./track_ch";
std::string dump_filename = configuration->property(role + ".dump_filename", default_dump_filename);
trk_param.dump_filename = dump_filename;
int vector_length = std::round(fs_in / (BEIDOU_B3I_CODE_RATE_HZ / BEIDOU_B3I_CODE_LENGTH_CHIPS));
trk_param.vector_length = vector_length;
int symbols_extended_correlator = configuration->property(role + ".extend_correlation_symbols", 1);
if (symbols_extended_correlator < 1)
{
symbols_extended_correlator = 1;
std::cout << TEXT_RED << "WARNING: BEIDOU B3I. extend_correlation_symbols must be bigger than 1. Coherent integration has been set to 1 symbol (1 ms)" << TEXT_RESET << std::endl;
}
else if (symbols_extended_correlator > 20)
{
symbols_extended_correlator = 20;
std::cout << TEXT_RED << "WARNING: BEIDOU B3I. extend_correlation_symbols must be lower than 21. Coherent integration has been set to 20 symbols (20 ms)" << TEXT_RESET << std::endl;
}
trk_param.extend_correlation_symbols = symbols_extended_correlator;
bool track_pilot = configuration->property(role + ".track_pilot", false);
if (track_pilot)
{
std::cout << TEXT_RED << "WARNING: BEIDOU B3I does not have pilot signal. Data tracking has been enabled" << TEXT_RESET << std::endl;
}
if ((symbols_extended_correlator > 1) and (pll_bw_narrow_hz > pll_bw_hz or dll_bw_narrow_hz > dll_bw_hz))
{
std::cout << TEXT_RED << "WARNING: BEIDOU B3I. PLL or DLL narrow tracking bandwidth is higher than wide tracking one" << TEXT_RESET << std::endl;
}
trk_param.very_early_late_space_chips = 0.0;
trk_param.very_early_late_space_narrow_chips = 0.0;
trk_param.track_pilot = false;
trk_param.system = 'C';
char sig_[3] = "B3";
std::memcpy(trk_param.signal, sig_, 3);
int cn0_samples = configuration->property(role + ".cn0_samples", 20);
if (FLAGS_cn0_samples != 20) cn0_samples = FLAGS_cn0_samples;
trk_param.cn0_samples = cn0_samples;
int cn0_min = configuration->property(role + ".cn0_min", 25);
if (FLAGS_cn0_min != 25) cn0_min = FLAGS_cn0_min;
trk_param.cn0_min = cn0_min;
int max_lock_fail = configuration->property(role + ".max_lock_fail", 50);
if (FLAGS_max_lock_fail != 50) max_lock_fail = FLAGS_max_lock_fail;
trk_param.max_lock_fail = max_lock_fail;
double carrier_lock_th = configuration->property(role + ".carrier_lock_th", 0.85);
if (FLAGS_carrier_lock_th != 0.85) carrier_lock_th = FLAGS_carrier_lock_th;
trk_param.carrier_lock_th = carrier_lock_th;
//################# MAKE TRACKING GNURadio object ###################
if (item_type == "gr_complex")
{
item_size_ = sizeof(gr_complex);
tracking_ = dll_pll_veml_make_tracking(trk_param);
}
else
{
item_size_ = sizeof(gr_complex);
LOG(WARNING) << item_type << " unknown tracking item type.";
}
channel_ = 0;
DLOG(INFO) << "tracking(" << tracking_->unique_id() << ")";
if (in_streams_ > 1)
{
LOG(ERROR) << "This implementation only supports one input stream";
}
if (out_streams_ > 1)
{
LOG(ERROR) << "This implementation only supports one output stream";
}
}
BeidouB3iDllPllTracking::~BeidouB3iDllPllTracking()
= default;
void BeidouB3iDllPllTracking::start_tracking()
{
tracking_->start_tracking();
}
void BeidouB3iDllPllTracking::stop_tracking()
{
tracking_->stop_tracking();
}
/*
* Set tracking channel unique ID
*/
void BeidouB3iDllPllTracking::set_channel(unsigned int channel)
{
channel_ = channel;
tracking_->set_channel(channel);
}
void BeidouB3iDllPllTracking::set_gnss_synchro(Gnss_Synchro* p_gnss_synchro)
{
tracking_->set_gnss_synchro(p_gnss_synchro);
}
void BeidouB3iDllPllTracking::connect(gr::top_block_sptr top_block)
{
if (top_block)
{ /* top_block is not null */
};
//nothing to connect, now the tracking uses gr_sync_decimator
}
void BeidouB3iDllPllTracking::disconnect(gr::top_block_sptr top_block)
{
if (top_block)
{ /* top_block is not null */
};
//nothing to disconnect, now the tracking uses gr_sync_decimator
}
gr::basic_block_sptr BeidouB3iDllPllTracking::get_left_block()
{
return tracking_;
}
gr::basic_block_sptr BeidouB3iDllPllTracking::get_right_block()
{
return tracking_;
}

106
src/algorithms/tracking/adapters/beidou_b3i_dll_pll_tracking.h Normal file
View File

@@ -0,0 +1,106 @@
/*!
* \file beidou_b3i_dll_pll_tracking.h
* \brief Interface of an adapter of a DLL+PLL tracking loop block
* for Beidou B3I to a TrackingInterface
* \author Damian Miralles, 2019. dmiralles2009(at)gmail.com
*
* Code DLL + carrier PLL according to the algorithms described in:
* K.Borre, D.M.Akos, N.Bertelsen, P.Rinder, and S.H.Jensen,
* A Software-Defined GPS and Galileo Receiver. A Single-Frequency
* Approach, Birkhauser, 2007
*
* -------------------------------------------------------------------------
*
* Copyright (C) 2010-2019 (see AUTHORS file for a list of contributors)
*
* GNSS-SDR is a software defined Global Navigation
* Satellite Systems receiver
*
* This file is part of GNSS-SDR.
*
* GNSS-SDR is free software: you can redistribute it and/or modify
* it under the terms of the GNU General Public License as published by
* the Free Software Foundation, either version 3 of the License, or
* (at your option) any later version.
*
* GNSS-SDR is distributed in the hope that it will be useful,
* but WITHOUT ANY WARRANTY; without even the implied warranty of
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
* GNU General Public License for more details.
*
* You should have received a copy of the GNU General Public License
* along with GNSS-SDR. If not, see <https://www.gnu.org/licenses/>.
*
* -------------------------------------------------------------------------
*/
#ifndef GNSS_SDR_BEIDOU_B3I_DLL_PLL_TRACKING_H_
#define GNSS_SDR_BEIDOU_B3I_DLL_PLL_TRACKING_H_
#include "dll_pll_veml_tracking.h"
#include "tracking_interface.h"
#include <string>
class ConfigurationInterface;
/*!
* \brief This class implements a code DLL + carrier PLL tracking loop
*/
class BeidouB3iDllPllTracking : public TrackingInterface
{
public:
BeidouB3iDllPllTracking(ConfigurationInterface* configuration,
const std::string& role,
unsigned int in_streams,
unsigned int out_streams);
virtual ~BeidouB3iDllPllTracking();
inline std::string role() override
{
return role_;
}
inline std::string implementation() override
{
return "BEIDOU_B3I_DLL_PLL_Tracking";
}
inline size_t item_size() override
{
return item_size_;
}
void connect(gr::top_block_sptr top_block) override;
void disconnect(gr::top_block_sptr top_block) override;
gr::basic_block_sptr get_left_block() override;
gr::basic_block_sptr get_right_block() override;
/*!
* \brief Set tracking channel unique ID
*/
void set_channel(unsigned int channel) override;
/*!
* \brief Set acquisition/tracking common Gnss_Synchro object pointer
* to efficiently exchange synchronization data between acquisition and tracking blocks
*/
void set_gnss_synchro(Gnss_Synchro* p_gnss_synchro) override;
void start_tracking() override;
/*!
* \brief Stop running tracking
*/
void stop_tracking() override;
private:
dll_pll_veml_tracking_sptr tracking_;
size_t item_size_;
unsigned int channel_;
std::string role_;
unsigned int in_streams_;
unsigned int out_streams_;
};
#endif // GNSS_SDR_BEIDOU_B3I_DLL_PLL_TRACKING_H_

232
src/algorithms/tracking/gnuradio_blocks/dll_pll_veml_tracking.cc
View File

@@ -36,6 +36,7 @@
#include "dll_pll_veml_tracking.h"
#include "Beidou_B1I.h"
#include "Beidou_B3I.h"
#include "GPS_L1_CA.h"
#include "GPS_L2C.h"
#include "GPS_L5.h"
@@ -43,9 +44,9 @@
#include "Galileo_E5a.h"
#include "MATH_CONSTANTS.h"
#include "beidou_b1i_signal_processing.h"
#include "beidou_b3i_signal_processing.h"
#include "galileo_e1_signal_processing.h"
#include "galileo_e5_signal_processing.h"
#include "gnss_satellite.h"
#include "gnss_sdr_create_directory.h"
#include "gnss_synchro.h"
#include "gps_l2c_signal.h"
@@ -62,11 +63,9 @@
#include <volk_gnsssdr/volk_gnsssdr.h>
#include <algorithm> // for fill_n
#include <cmath> // for fmod, round, floor
#include <complex> // for complex
#include <cstdlib> // for abs, size_t
#include <exception> // for exception
#include <iostream> // for cout, cerr
#include <map> // for map
#include <map>
dll_pll_veml_tracking_sptr dll_pll_veml_make_tracking(const Dll_Pll_Conf &conf_)
@@ -75,16 +74,6 @@ dll_pll_veml_tracking_sptr dll_pll_veml_make_tracking(const Dll_Pll_Conf &conf_)
}
void dll_pll_veml_tracking::forecast(int noutput_items,
gr_vector_int &ninput_items_required)
{
if (noutput_items != 0)
{
ninput_items_required[0] = static_cast<int32_t>(trk_parameters.vector_length) * 2;
}
}
dll_pll_veml_tracking::dll_pll_veml_tracking(const Dll_Pll_Conf &conf_) : gr::block("dll_pll_veml_tracking", gr::io_signature::make(1, 1, sizeof(gr_complex)),
gr::io_signature::make(1, 1, sizeof(Gnss_Synchro)))
{
@@ -93,12 +82,14 @@ dll_pll_veml_tracking::dll_pll_veml_tracking(const Dll_Pll_Conf &conf_) : gr::bl
this->message_port_register_out(pmt::mp("events"));
this->set_relative_rate(1.0 / static_cast<double>(trk_parameters.vector_length));
// Telemetry bit synchronization message port input (mainly for GPS L1 CA)
this->message_port_register_in(pmt::mp("preamble_samplestamp"));
// Telemetry message port input
this->message_port_register_in(pmt::mp("telemetry_to_trk"));
this->set_msg_handler(pmt::mp("telemetry_to_trk"), boost::bind(&dll_pll_veml_tracking::msg_handler_telemetry_to_trk, this, _1));
// initialize internal vars
d_veml = false;
d_cloop = true;
d_pull_in_transitory = true;
d_code_chip_rate = 0.0;
d_secondary_code_length = 0U;
d_secondary_code_string = nullptr;
@@ -115,6 +106,7 @@ dll_pll_veml_tracking::dll_pll_veml_tracking(const Dll_Pll_Conf &conf_) : gr::bl
map_signal_pretty_name["5X"] = "E5a";
map_signal_pretty_name["L5"] = "L5";
map_signal_pretty_name["B1"] = "B1I";
map_signal_pretty_name["B3"] = "B3I";
signal_pretty_name = map_signal_pretty_name[signal_type];
@@ -288,16 +280,32 @@ dll_pll_veml_tracking::dll_pll_veml_tracking(const Dll_Pll_Conf &conf_) : gr::bl
d_signal_carrier_freq = BEIDOU_B1I_FREQ_HZ;
d_code_period = BEIDOU_B1I_CODE_PERIOD;
d_code_chip_rate = BEIDOU_B1I_CODE_RATE_HZ;
d_code_length_chips = static_cast<unsigned int>(BEIDOU_B1I_CODE_LENGTH_CHIPS);
d_code_length_chips = static_cast<uint32_t>(BEIDOU_B1I_CODE_LENGTH_CHIPS);
d_symbols_per_bit = BEIDOU_B1I_TELEMETRY_SYMBOLS_PER_BIT;
d_correlation_length_ms = 1;
d_code_samples_per_chip = 1;
d_secondary = true;
trk_parameters.track_pilot = false;
interchange_iq = false;
d_secondary_code_length = static_cast<unsigned int>(BEIDOU_B1I_SECONDARY_CODE_LENGTH);
d_secondary_code_length = static_cast<uint32_t>(BEIDOU_B1I_SECONDARY_CODE_LENGTH);
d_secondary_code_string = const_cast<std::string *>(&BEIDOU_B1I_SECONDARY_CODE_STR);
}
else if (signal_type == "B3")
{
// GEO Satellites use different secondary code
d_signal_carrier_freq = BEIDOU_B3I_FREQ_HZ;
d_code_period = BEIDOU_B3I_CODE_PERIOD;
d_code_chip_rate = BEIDOU_B3I_CODE_RATE_HZ;
d_code_length_chips = static_cast<uint32_t>(BEIDOU_B3I_CODE_LENGTH_CHIPS);
d_symbols_per_bit = BEIDOU_B3I_TELEMETRY_SYMBOLS_PER_BIT;
d_correlation_length_ms = 1;
d_code_samples_per_chip = 1;
d_secondary = true;
trk_parameters.track_pilot = false;
interchange_iq = false;
d_secondary_code_length = static_cast<uint32_t>(BEIDOU_B3I_SECONDARY_CODE_LENGTH);
d_secondary_code_string = const_cast<std::string *>(&BEIDOU_B3I_SECONDARY_CODE_STR);
}
else
{
LOG(WARNING) << "Invalid Signal argument when instantiating tracking blocks";
@@ -332,10 +340,8 @@ dll_pll_veml_tracking::dll_pll_veml_tracking(const Dll_Pll_Conf &conf_) : gr::bl
K_blk_samples = 0.0;
// Initialize tracking ==========================================
d_code_loop_filter = Tracking_2nd_DLL_filter(static_cast<float>(d_code_period));
d_carrier_loop_filter = Tracking_2nd_PLL_filter(static_cast<float>(d_code_period));
d_code_loop_filter.set_DLL_BW(trk_parameters.dll_bw_hz);
d_carrier_loop_filter.set_PLL_BW(trk_parameters.pll_bw_hz);
d_code_loop_filter = Tracking_loop_filter(d_code_period, trk_parameters.dll_bw_hz, trk_parameters.dll_filter_order, false);
d_carrier_loop_filter.set_params(trk_parameters.fll_bw_hz, trk_parameters.pll_bw_hz, trk_parameters.pll_filter_order);
// Initialization of local code replica
// Get space for a vector with the sinboc(1,1) replica sampled 2x/chip
@@ -423,6 +429,7 @@ dll_pll_veml_tracking::dll_pll_veml_tracking(const Dll_Pll_Conf &conf_) : gr::bl
d_acq_sample_stamp = 0ULL;
d_current_prn_length_samples = static_cast<int32_t>(trk_parameters.vector_length);
d_current_correlation_time_s = 0.0;
// CN0 estimation and lock detector buffers
d_cn0_estimation_counter = 0;
@@ -446,7 +453,6 @@ dll_pll_veml_tracking::dll_pll_veml_tracking(const Dll_Pll_Conf &conf_) : gr::bl
d_carrier_phase_step_rad = 0.0;
d_carrier_phase_rate_step_rad = 0.0;
d_rem_code_phase_chips = 0.0;
d_last_prompt = gr_complex(0.0, 0.0);
d_state = 0; // initial state: standby
clear_tracking_vars();
if (trk_parameters.smoother_length > 0)
@@ -498,6 +504,48 @@ dll_pll_veml_tracking::dll_pll_veml_tracking(const Dll_Pll_Conf &conf_) : gr::bl
}
void dll_pll_veml_tracking::forecast(int noutput_items,
gr_vector_int &ninput_items_required)
{
if (noutput_items != 0)
{
ninput_items_required[0] = static_cast<int32_t>(trk_parameters.vector_length) * 2;
}
}
void dll_pll_veml_tracking::msg_handler_telemetry_to_trk(const pmt::pmt_t &msg)
{
try
{
if (pmt::any_ref(msg).type() == typeid(int))
{
int tlm_event;
tlm_event = boost::any_cast<int>(pmt::any_ref(msg));
switch (tlm_event)
{
case 1: //tlm fault in current channel
{
DLOG(INFO) << "Telemetry fault received in ch " << this->d_channel;
gr::thread::scoped_lock lock(d_setlock);
d_carrier_lock_fail_counter = 10000; //force loss-of-lock condition
break;
}
default:
{
break;
}
}
}
}
catch (boost::bad_any_cast &e)
{
LOG(WARNING) << "msg_handler_telemetry_to_trk Bad any cast!";
}
}
void dll_pll_veml_tracking::start_tracking()
{
gr::thread::scoped_lock l(d_setlock);
@@ -513,8 +561,8 @@ void dll_pll_veml_tracking::start_tracking()
d_carr_ph_history.clear();
d_code_ph_history.clear();
// DLL/PLL filter initialization
d_carrier_loop_filter.initialize(); // initialize the carrier filter
d_code_loop_filter.initialize(); // initialize the code filter
d_carrier_loop_filter.initialize(static_cast<float>(d_acq_carrier_doppler_hz)); // initialize the carrier filter
d_code_loop_filter.initialize(); // initialize the code filter
if (systemName == "GPS" and signal_type == "1C")
{
@@ -596,8 +644,8 @@ void dll_pll_veml_tracking::start_tracking()
d_preamble_length_symbols = 22;
d_preambles_symbols = static_cast<int32_t *>(volk_gnsssdr_malloc(22 * sizeof(int32_t), volk_gnsssdr_get_alignment()));
int32_t n = 0;
uint16_t preambles_bits[BEIDOU_B1I_PREAMBLE_LENGTH_BITS] = {1, 1, 1, 0, 0, 0, 1, 0, 0, 1, 0};
for (uint16_t preambles_bit : preambles_bits)
uint32_t preambles_bits[BEIDOU_B1I_PREAMBLE_LENGTH_BITS] = {1, 1, 1, 0, 0, 0, 1, 0, 0, 1, 0};
for (uint32_t preambles_bit : preambles_bits)
{
for (int32_t j = 0; j < d_symbols_per_bit; j++)
{
@@ -612,7 +660,47 @@ void dll_pll_veml_tracking::start_tracking()
n++;
}
}
d_symbol_history.set_capacity(22); // Change fixed buffer size
d_symbol_history.resize(22); // Change fixed buffer size
d_symbol_history.clear();
}
}
else if (systemName == "Beidou" and signal_type == "B3")
{
beidou_b3i_code_gen_float(d_tracking_code, d_acquisition_gnss_synchro->PRN, 0);
// Update secondary code settings for geo satellites
if (d_acquisition_gnss_synchro->PRN > 0 and d_acquisition_gnss_synchro->PRN < 6)
{
d_symbols_per_bit = 2;
d_correlation_length_ms = 1;
d_code_samples_per_chip = 1;
d_secondary = false;
trk_parameters.track_pilot = false;
interchange_iq = false;
d_secondary_code_length = 0;
d_secondary_code_string = const_cast<std::string *>(&BEIDOU_B3I_D2_SECONDARY_CODE_STR);
// preamble bits to sampled symbols
d_preamble_length_symbols = 22;
d_preambles_symbols = static_cast<int32_t *>(volk_gnsssdr_malloc(22 * sizeof(int32_t), volk_gnsssdr_get_alignment()));
int32_t n = 0;
uint32_t preambles_bits[BEIDOU_B3I_PREAMBLE_LENGTH_BITS] = {1, 1, 1, 0, 0, 0, 1, 0, 0, 1, 0};
for (uint32_t preambles_bit : preambles_bits)
{
for (int32_t j = 0; j < d_symbols_per_bit; j++)
{
if (preambles_bit == 1)
{
d_preambles_symbols[n] = 1;
}
else
{
d_preambles_symbols[n] = -1;
}
n++;
}
}
d_symbol_history.resize(22); // Change fixed buffer size
d_symbol_history.clear();
}
}
@@ -642,10 +730,10 @@ void dll_pll_veml_tracking::start_tracking()
d_local_code_shift_chips[2] = trk_parameters.early_late_space_chips * static_cast<float>(d_code_samples_per_chip);
}
d_code_loop_filter.set_DLL_BW(trk_parameters.dll_bw_hz);
d_carrier_loop_filter.set_PLL_BW(trk_parameters.pll_bw_hz);
d_carrier_loop_filter.set_pdi(static_cast<float>(d_code_period));
d_code_loop_filter.set_pdi(static_cast<float>(d_code_period));
d_current_correlation_time_s = d_code_period;
d_code_loop_filter.set_noise_bandwidth(trk_parameters.dll_bw_hz);
d_code_loop_filter.set_update_interval(d_code_period);
// DEBUG OUTPUT
std::cout << "Tracking of " << systemName << " " << signal_pretty_name << " signal started on channel " << d_channel << " for satellite " << Gnss_Satellite(systemName, d_acquisition_gnss_synchro->PRN) << std::endl;
@@ -654,8 +742,8 @@ void dll_pll_veml_tracking::start_tracking()
// enable tracking pull-in
d_state = 1;
d_cloop = true;
d_pull_in_transitory = true;
d_Prompt_circular_buffer.clear();
d_last_prompt = gr_complex(0.0, 0.0);
}
@@ -764,15 +852,18 @@ bool dll_pll_veml_tracking::cn0_and_tracking_lock_status(double coh_integration_
// Carrier lock indicator
d_carrier_lock_test = carrier_lock_detector(d_Prompt_buffer, trk_parameters.cn0_samples);
// Loss of lock detection
if (d_carrier_lock_test < d_carrier_lock_threshold or d_CN0_SNV_dB_Hz < trk_parameters.cn0_min)
if (!d_pull_in_transitory)
{
d_carrier_lock_fail_counter++;
}
else
{
if (d_carrier_lock_fail_counter > 0)
if (d_carrier_lock_test < d_carrier_lock_threshold or d_CN0_SNV_dB_Hz < trk_parameters.cn0_min)
{
d_carrier_lock_fail_counter--;
d_carrier_lock_fail_counter++;
}
else
{
if (d_carrier_lock_fail_counter > 0)
{
d_carrier_lock_fail_counter--;
}
}
}
if (d_carrier_lock_fail_counter > trk_parameters.max_lock_fail)
@@ -827,20 +918,43 @@ void dll_pll_veml_tracking::run_dll_pll()
if (d_cloop)
{
// Costas loop discriminator, insensitive to 180 deg phase transitions
d_carr_error_hz = pll_cloop_two_quadrant_atan(d_P_accu) / PI_2;
d_carr_phase_error_hz = pll_cloop_two_quadrant_atan(d_P_accu) / PI_2;
}
else
{
// Secondary code acquired. No symbols transition should be present in the signal
d_carr_error_hz = pll_four_quadrant_atan(d_P_accu) / PI_2;
d_carr_phase_error_hz = pll_four_quadrant_atan(d_P_accu) / PI_2;
}
if ((d_pull_in_transitory == true and trk_parameters.enable_fll_pull_in == true) or trk_parameters.enable_fll_steady_state)
{
// FLL discriminator
d_carr_freq_error_hz = fll_four_quadrant_atan(d_P_accu_old, d_P_accu, 0, d_current_correlation_time_s) / GPS_TWO_PI;
d_P_accu_old = d_P_accu;
//std::cout << "d_carr_freq_error_hz: " << d_carr_freq_error_hz << std::endl;
// Carrier discriminator filter
if ((d_pull_in_transitory == true and trk_parameters.enable_fll_pull_in == true))
{
//pure FLL, disable PLL
d_carr_error_filt_hz = d_carrier_loop_filter.get_carrier_error(d_carr_freq_error_hz, 0, d_current_correlation_time_s);
}
else
{
//FLL-aided PLL
d_carr_error_filt_hz = d_carrier_loop_filter.get_carrier_error(d_carr_freq_error_hz, d_carr_phase_error_hz, d_current_correlation_time_s);
}
}
else
{
// Carrier discriminator filter
d_carr_error_filt_hz = d_carrier_loop_filter.get_carrier_error(0, d_carr_phase_error_hz, d_current_correlation_time_s);
}
// Carrier discriminator filter
d_carr_error_filt_hz = d_carrier_loop_filter.get_carrier_nco(d_carr_error_hz);
// New carrier Doppler frequency estimation
d_carrier_doppler_hz = d_acq_carrier_doppler_hz + d_carr_error_filt_hz;
d_carrier_doppler_hz = d_carr_error_filt_hz;
// std::cout << "d_carrier_doppler_hz: " << d_carrier_doppler_hz << std::endl;
// std::cout << "d_CN0_SNV_dB_Hz: " << this->d_CN0_SNV_dB_Hz << std::endl;
// ################## DLL ##########################################################
// DLL discriminator
if (d_veml)
@@ -852,7 +966,7 @@ void dll_pll_veml_tracking::run_dll_pll()
d_code_error_chips = dll_nc_e_minus_l_normalized(d_E_accu, d_L_accu); // [chips/Ti]
}
// Code discriminator filter
d_code_error_filt_chips = d_code_loop_filter.get_code_nco(d_code_error_chips); // [chips/second]
d_code_error_filt_chips = d_code_loop_filter.apply(d_code_error_chips); // [chips/second]
// New code Doppler frequency estimation
d_code_freq_chips = (1.0 + (d_carrier_doppler_hz / d_signal_carrier_freq)) * d_code_chip_rate - d_code_error_filt_chips;
@@ -866,13 +980,14 @@ void dll_pll_veml_tracking::clear_tracking_vars()
{
d_Prompt_Data[0] = gr_complex(0.0, 0.0);
}
d_carr_error_hz = 0.0;
d_P_accu_old = gr_complex(0.0, 0.0);
d_carr_phase_error_hz = 0.0;
d_carr_freq_error_hz = 0.0;
d_carr_error_filt_hz = 0.0;
d_code_error_chips = 0.0;
d_code_error_filt_chips = 0.0;
d_current_symbol = 0;
d_Prompt_circular_buffer.clear();
d_last_prompt = gr_complex(0.0, 0.0);
d_carrier_phase_rate_step_rad = 0.0;
d_code_phase_rate_step_chips = 0.0;
d_carr_ph_history.clear();
@@ -1104,7 +1219,7 @@ void dll_pll_veml_tracking::log_data(bool integrating)
tmp_float = d_code_phase_rate_step_chips * trk_parameters.fs_in * trk_parameters.fs_in;
d_dump_file.write(reinterpret_cast<char *>(&tmp_float), sizeof(float));
// PLL commands
tmp_float = d_carr_error_hz;
tmp_float = d_carr_phase_error_hz;
d_dump_file.write(reinterpret_cast<char *>(&tmp_float), sizeof(float));
tmp_float = d_carr_error_filt_hz;
d_dump_file.write(reinterpret_cast<char *>(&tmp_float), sizeof(float));
@@ -1434,6 +1549,13 @@ int dll_pll_veml_tracking::general_work(int noutput_items __attribute__((unused)
auto **out = reinterpret_cast<Gnss_Synchro **>(&output_items[0]);
Gnss_Synchro current_synchro_data = Gnss_Synchro();
if (d_pull_in_transitory == true)
{
if (trk_parameters.pull_in_time_s < (d_sample_counter - d_acq_sample_stamp) / static_cast<int>(trk_parameters.fs_in))
{
d_pull_in_transitory = false;
}
}
switch (d_state)
{
case 0: // Standby - Consume samples at full throttle, do nothing
@@ -1609,7 +1731,6 @@ int dll_pll_veml_tracking::general_work(int noutput_items __attribute__((unused)
d_P_accu = gr_complex(0.0, 0.0);
d_L_accu = gr_complex(0.0, 0.0);
d_VL_accu = gr_complex(0.0, 0.0);
d_last_prompt = gr_complex(0.0, 0.0);
d_Prompt_circular_buffer.clear();
d_current_symbol = 0;
@@ -1617,9 +1738,7 @@ int dll_pll_veml_tracking::general_work(int noutput_items __attribute__((unused)
{
// UPDATE INTEGRATION TIME
d_extend_correlation_symbols_count = 0;
float new_correlation_time = static_cast<float>(trk_parameters.extend_correlation_symbols) * static_cast<float>(d_code_period);
d_carrier_loop_filter.set_pdi(new_correlation_time);
d_code_loop_filter.set_pdi(new_correlation_time);
d_current_correlation_time_s = static_cast<float>(trk_parameters.extend_correlation_symbols) * static_cast<float>(d_code_period);
d_state = 3; // next state is the extended correlator integrator
LOG(INFO) << "Enabled " << trk_parameters.extend_correlation_symbols * static_cast<int32_t>(d_code_period * 1000.0) << " ms extended correlator in channel "
<< d_channel
@@ -1628,8 +1747,9 @@ int dll_pll_veml_tracking::general_work(int noutput_items __attribute__((unused)
<< d_channel
<< " for satellite " << Gnss_Satellite(systemName, d_acquisition_gnss_synchro->PRN) << std::endl;
// Set narrow taps delay values [chips]
d_code_loop_filter.set_DLL_BW(trk_parameters.dll_bw_narrow_hz);
d_carrier_loop_filter.set_PLL_BW(trk_parameters.pll_bw_narrow_hz);
d_code_loop_filter.set_update_interval(d_current_correlation_time_s);
d_code_loop_filter.set_noise_bandwidth(trk_parameters.dll_bw_narrow_hz);
d_carrier_loop_filter.set_params(trk_parameters.fll_bw_hz, trk_parameters.pll_bw_narrow_hz, trk_parameters.pll_filter_order);
if (d_veml)
{
d_local_code_shift_chips[0] = -trk_parameters.very_early_late_space_narrow_chips * static_cast<float>(d_code_samples_per_chip);

21
src/algorithms/tracking/gnuradio_blocks/dll_pll_veml_tracking.h
View File

@@ -34,8 +34,8 @@
#include "cpu_multicorrelator_real_codes.h"
#include "dll_pll_conf.h"
#include "tracking_2nd_DLL_filter.h"
#include "tracking_2nd_PLL_filter.h"
#include "tracking_FLL_PLL_filter.h" // for PLL/FLL filter
#include "tracking_loop_filter.h" // for DLL filter
#include <boost/circular_buffer.hpp>
#include <boost/shared_ptr.hpp> // for boost::shared_ptr
#include <gnuradio/block.h> // for block
@@ -44,10 +44,8 @@
#include <pmt/pmt.h> // for pmt_t
#include <cstdint> // for int32_t
#include <fstream> // for string, ofstream
#include <string> // for string
#include <utility> // for pair
class Gnss_Synchro;
class dll_pll_veml_tracking;
@@ -75,9 +73,8 @@ public:
private:
friend dll_pll_veml_tracking_sptr dll_pll_veml_make_tracking(const Dll_Pll_Conf &conf_);
void msg_handler_telemetry_to_trk(const pmt::pmt_t &msg);
dll_pll_veml_tracking(const Dll_Pll_Conf &conf_);
void msg_handler_preamble_index(pmt::pmt_t msg);
bool cn0_and_tracking_lock_status(double coh_integration_time_s);
bool acquire_secondary();
@@ -147,9 +144,9 @@ private:
gr_complex d_VE_accu;
gr_complex d_E_accu;
gr_complex d_P_accu;
gr_complex d_P_accu_old;
gr_complex d_L_accu;
gr_complex d_VL_accu;
gr_complex d_last_prompt;
gr_complex *d_Prompt_Data;
@@ -163,16 +160,18 @@ private:
double d_rem_code_phase_samples;
float d_rem_carr_phase_rad;
// PLL and DLL filter library
Tracking_2nd_DLL_filter d_code_loop_filter;
Tracking_2nd_PLL_filter d_carrier_loop_filter;
Tracking_loop_filter d_code_loop_filter;
Tracking_FLL_PLL_filter d_carrier_loop_filter;
// acquisition
double d_acq_code_phase_samples;
double d_acq_carrier_doppler_hz;
// tracking vars
double d_carr_error_hz;
bool d_pull_in_transitory;
double d_current_correlation_time_s;
double d_carr_phase_error_hz;
double d_carr_freq_error_hz;
double d_carr_error_filt_hz;
double d_code_error_chips;
double d_code_error_filt_chips;

3
src/algorithms/tracking/gnuradio_blocks/dll_pll_veml_tracking_fpga.cc
View File

@@ -78,6 +78,9 @@ dll_pll_veml_tracking_fpga::dll_pll_veml_tracking_fpga(const Dll_Pll_Conf_Fpga &
// Telemetry bit synchronization message port input (mainly for GPS L1 CA)
this->message_port_register_in(pmt::mp("preamble_samplestamp"));
// Telemetry message port input
this->message_port_register_in(pmt::mp("telemetry_to_trk"));
//todo: Implement the telemetry_to_trk handler in the same way the software version of tracking
// initialize internal vars
d_veml = false;

2
src/algorithms/tracking/gnuradio_blocks/galileo_e1_tcp_connector_tracking_cc.cc
View File

@@ -95,6 +95,8 @@ Galileo_E1_Tcp_Connector_Tracking_cc::Galileo_E1_Tcp_Connector_Tracking_cc(
{
this->message_port_register_out(pmt::mp("events"));
this->set_relative_rate(1.0 / vector_length);
// Telemetry message port input
this->message_port_register_in(pmt::mp("telemetry_to_trk"));
// initialize internal vars
d_dump = dump;
d_fs_in = fs_in;

1
src/algorithms/tracking/gnuradio_blocks/glonass_l1_ca_dll_pll_c_aid_tracking_cc.cc
View File

@@ -120,6 +120,7 @@ glonass_l1_ca_dll_pll_c_aid_tracking_cc::glonass_l1_ca_dll_pll_c_aid_tracking_cc
boost::bind(&glonass_l1_ca_dll_pll_c_aid_tracking_cc::msg_handler_preamble_index, this, _1));
this->message_port_register_out(pmt::mp("events"));
this->message_port_register_in(pmt::mp("telemetry_to_trk"));
// initialize internal vars
d_dump = dump;
d_fs_in = fs_in;

1
src/algorithms/tracking/gnuradio_blocks/glonass_l1_ca_dll_pll_c_aid_tracking_sc.cc
View File

@@ -115,6 +115,7 @@ glonass_l1_ca_dll_pll_c_aid_tracking_sc::glonass_l1_ca_dll_pll_c_aid_tracking_sc
this->set_msg_handler(pmt::mp("preamble_timestamp_s"),
boost::bind(&glonass_l1_ca_dll_pll_c_aid_tracking_sc::msg_handler_preamble_index, this, _1));
this->message_port_register_out(pmt::mp("events"));
this->message_port_register_in(pmt::mp("telemetry_to_trk"));
// initialize internal vars
d_dump = dump;
d_fs_in = fs_in;

2
src/algorithms/tracking/gnuradio_blocks/glonass_l1_ca_dll_pll_tracking_cc.cc
View File

@@ -92,7 +92,7 @@ Glonass_L1_Ca_Dll_Pll_Tracking_cc::Glonass_L1_Ca_Dll_Pll_Tracking_cc(
gr::io_signature::make(1, 1, sizeof(Gnss_Synchro)))
{
this->message_port_register_out(pmt::mp("events"));
this->message_port_register_in(pmt::mp("telemetry_to_trk"));
// initialize internal vars
d_dump = dump;
d_fs_in = fs_in;

1
src/algorithms/tracking/gnuradio_blocks/glonass_l2_ca_dll_pll_c_aid_tracking_cc.cc
View File

@@ -117,6 +117,7 @@ glonass_l2_ca_dll_pll_c_aid_tracking_cc::glonass_l2_ca_dll_pll_c_aid_tracking_cc
boost::bind(&glonass_l2_ca_dll_pll_c_aid_tracking_cc::msg_handler_preamble_index, this, _1));
this->message_port_register_out(pmt::mp("events"));
this->message_port_register_in(pmt::mp("telemetry_to_trk"));
// initialize internal vars
d_dump = dump;
d_fs_in = fs_in;

1
src/algorithms/tracking/gnuradio_blocks/glonass_l2_ca_dll_pll_c_aid_tracking_sc.cc
View File

@@ -114,6 +114,7 @@ glonass_l2_ca_dll_pll_c_aid_tracking_sc::glonass_l2_ca_dll_pll_c_aid_tracking_sc
this->set_msg_handler(pmt::mp("preamble_timestamp_s"),
boost::bind(&glonass_l2_ca_dll_pll_c_aid_tracking_sc::msg_handler_preamble_index, this, _1));
this->message_port_register_out(pmt::mp("events"));
this->message_port_register_in(pmt::mp("telemetry_to_trk"));
// initialize internal vars
d_dump = dump;
d_fs_in = fs_in;

2
src/algorithms/tracking/gnuradio_blocks/glonass_l2_ca_dll_pll_tracking_cc.cc
View File

@@ -92,7 +92,7 @@ Glonass_L2_Ca_Dll_Pll_Tracking_cc::Glonass_L2_Ca_Dll_Pll_Tracking_cc(
gr::io_signature::make(1, 1, sizeof(Gnss_Synchro)))
{
this->message_port_register_out(pmt::mp("events"));
this->message_port_register_in(pmt::mp("telemetry_to_trk"));
// initialize internal vars
d_dump = dump;
d_fs_in = fs_in;

2
src/algorithms/tracking/gnuradio_blocks/gps_l1_ca_dll_pll_c_aid_tracking_cc.cc
View File

@@ -108,6 +108,8 @@ gps_l1_ca_dll_pll_c_aid_tracking_cc::gps_l1_ca_dll_pll_c_aid_tracking_cc(
boost::bind(&gps_l1_ca_dll_pll_c_aid_tracking_cc::msg_handler_preamble_index, this, _1));
this->message_port_register_out(pmt::mp("events"));
this->message_port_register_in(pmt::mp("telemetry_to_trk"));
// initialize internal vars
d_dump = dump;
d_fs_in = fs_in;

1
src/algorithms/tracking/gnuradio_blocks/gps_l1_ca_dll_pll_c_aid_tracking_sc.cc
View File

@@ -105,6 +105,7 @@ gps_l1_ca_dll_pll_c_aid_tracking_sc::gps_l1_ca_dll_pll_c_aid_tracking_sc(
this->set_msg_handler(pmt::mp("preamble_timestamp_s"),
boost::bind(&gps_l1_ca_dll_pll_c_aid_tracking_sc::msg_handler_preamble_index, this, _1));
this->message_port_register_out(pmt::mp("events"));
this->message_port_register_in(pmt::mp("telemetry_to_trk"));
// initialize internal vars
d_dump = dump;
d_fs_in = fs_in;

1
src/algorithms/tracking/gnuradio_blocks/gps_l1_ca_dll_pll_tracking_gpu_cc.cc
View File

@@ -82,6 +82,7 @@ Gps_L1_Ca_Dll_Pll_Tracking_GPU_cc::Gps_L1_Ca_Dll_Pll_Tracking_GPU_cc(
// Telemetry bit synchronization message port input
this->message_port_register_in(pmt::mp("preamble_timestamp_s"));
this->message_port_register_out(pmt::mp("events"));
this->message_port_register_in(pmt::mp("telemetry_to_trk"));
// initialize internal vars
d_dump = dump;
d_fs_in = fs_in;

2
src/algorithms/tracking/gnuradio_blocks/gps_l1_ca_kf_tracking_cc.cc
View File

@@ -106,7 +106,7 @@ Gps_L1_Ca_Kf_Tracking_cc::Gps_L1_Ca_Kf_Tracking_cc(
// Telemetry bit synchronization message port input
this->message_port_register_in(pmt::mp("preamble_timestamp_s"));
this->message_port_register_out(pmt::mp("events"));
this->message_port_register_in(pmt::mp("telemetry_to_trk"));
// initialize internal vars
d_order = order;
d_dump = dump;

1
src/algorithms/tracking/gnuradio_blocks/gps_l1_ca_tcp_connector_tracking_cc.cc
View File

@@ -87,6 +87,7 @@ Gps_L1_Ca_Tcp_Connector_Tracking_cc::Gps_L1_Ca_Tcp_Connector_Tracking_cc(
gr::io_signature::make(1, 1, sizeof(Gnss_Synchro)))
{
this->message_port_register_out(pmt::mp("events"));
this->message_port_register_in(pmt::mp("telemetry_to_trk"));
// initialize internal vars
d_dump = dump;
d_fs_in = fs_in;

4
src/algorithms/tracking/libs/cpu_multicorrelator.cc
View File

@@ -1,11 +1,11 @@
/*!
* \file cpu_multicorrelator.cc
* \brief High optimized CPU vector multiTAP correlator class
* \brief Highly optimized CPU vector multiTAP correlator class
* \authors <ul>
* <li> Javier Arribas, 2015. jarribas(at)cttc.es
* </ul>
*
* Class that implements a high optimized vector multiTAP correlator class for CPUs
* Class that implements a highly optimized vector multiTAP correlator class for CPUs
*
* -------------------------------------------------------------------------
*

2
src/algorithms/tracking/libs/cpu_multicorrelator.h
View File

@@ -65,4 +65,4 @@ private:
};
#endif /* CPU_MULTICORRELATOR_H_ */
#endif /* GNSS_SDR_CPU_MULTICORRELATOR_H_ */

4
src/algorithms/tracking/libs/cpu_multicorrelator_16sc.cc
View File

@@ -1,11 +1,11 @@
/*!
* \file cpu_multicorrelator_16sc.cc
* \brief High optimized CPU vector multiTAP correlator class
* \brief Highly optimized CPU vector multiTAP correlator class
* \authors <ul>
* <li> Javier Arribas, 2015. jarribas(at)cttc.es
* </ul>
*
* Class that implements a high optimized vector multiTAP correlator class for CPUs
* Class that implements a highly optimized vector multiTAP correlator class for CPUs
*
* -------------------------------------------------------------------------
*

4
src/algorithms/tracking/libs/cpu_multicorrelator_16sc.h
View File

@@ -1,11 +1,11 @@
/*!
* \file cpu_multicorrelator_16sc.h
* \brief High optimized CPU vector multiTAP correlator class for lv_16sc_t (short int complex)
* \brief Highly optimized CPU vector multiTAP correlator class for lv_16sc_t (short int complex)
* \authors <ul>
* <li> Javier Arribas, 2016. jarribas(at)cttc.es
* </ul>
*
* Class that implements a high optimized vector multiTAP correlator class for CPUs
* Class that implements a highly optimized vector multiTAP correlator class for CPUs
*
* -------------------------------------------------------------------------
*

7
src/algorithms/tracking/libs/cpu_multicorrelator_real_codes.cc
View File

@@ -6,7 +6,7 @@
* <li> Cillian O'Driscoll, 2017. cillian.odriscoll(at)gmail.com
* </ul>
*
* Class that implements a high optimized vector multiTAP correlator class for CPUs
* Class that implements a highly optimized vector multiTAP correlator class for CPUs
*
* -------------------------------------------------------------------------
*
@@ -125,7 +125,7 @@ void Cpu_Multicorrelator_Real_Codes::update_local_code(int correlator_length_sam
}
}
// Overload Carrier_wipeoff_multicorrelator_resampler to ensure back compatibility
bool Cpu_Multicorrelator_Real_Codes::Carrier_wipeoff_multicorrelator_resampler(
float rem_carrier_phase_in_rad,
float phase_step_rad,
@@ -150,7 +150,8 @@ bool Cpu_Multicorrelator_Real_Codes::Carrier_wipeoff_multicorrelator_resampler(
}
return true;
}
// Overload Carrier_wipeoff_multicorrelator_resampler to ensure back compatibility
bool Cpu_Multicorrelator_Real_Codes::Carrier_wipeoff_multicorrelator_resampler(
float rem_carrier_phase_in_rad,
float phase_step_rad,

5
src/algorithms/tracking/libs/cpu_multicorrelator_real_codes.h
View File

@@ -6,7 +6,7 @@
* <li> Cillian O'Driscoll, 2017, cillian.odriscoll(at)gmail.com
* </ul>
*
* Class that implements a high optimized vector multiTAP correlator class for CPUs
* Class that implements a highly optimized vector multiTAP correlator class for CPUs
*
* -------------------------------------------------------------------------
*
@@ -52,7 +52,6 @@ public:
bool set_local_code_and_taps(int code_length_chips, const float *local_code_in, float *shifts_chips);
bool set_input_output_vectors(std::complex<float> *corr_out, const std::complex<float> *sig_in);
void update_local_code(int correlator_length_samples, float rem_code_phase_chips, float code_phase_step_chips, float code_phase_rate_step_chips = 0.0);
// Overload Carrier_wipeoff_multicorrelator_resampler to ensure back compatibility
bool Carrier_wipeoff_multicorrelator_resampler(float rem_carrier_phase_in_rad, float phase_step_rad, float phase_rate_step_rad, float rem_code_phase_chips, float code_phase_step_chips, float code_phase_rate_step_chips, int signal_length_samples);
bool Carrier_wipeoff_multicorrelator_resampler(float rem_carrier_phase_in_rad, float phase_step_rad, float rem_code_phase_chips, float code_phase_step_chips, float code_phase_rate_step_chips, int signal_length_samples);
bool free();
@@ -70,4 +69,4 @@ private:
};
#endif /* CPU_MULTICORRELATOR_REAL_CODES_H_ */
#endif /* GNSS_SDR_CPU_MULTICORRELATOR_REAL_CODES_H_ */

481
src/algorithms/tracking/libs/cuda_multicorrelator.cu
View File

@@ -1,11 +1,11 @@
/*!
* \file cuda_multicorrelator.cu
* \brief High optimized CUDA GPU vector multiTAP correlator class
* \brief Highly optimized CUDA GPU vector multiTAP correlator class
* \authors <ul>
* <li> Javier Arribas, 2015. jarribas(at)cttc.es
* </ul>
*
* Class that implements a high optimized vector multiTAP correlator class for NVIDIA CUDA GPUs
* Class that implements a highly optimized vector multiTAP correlator class for NVIDIA CUDA GPUs
*
* -------------------------------------------------------------------------
*
@@ -33,9 +33,8 @@
*/
#include "cuda_multicorrelator.h"
#include <stdio.h>
#include <iostream>
#include <stdio.h>
// For the CUDA runtime routines (prefixed with "cuda_")
#include <cuda_runtime.h>
@@ -53,22 +52,21 @@ __global__ void Doppler_wippe_scalarProdGPUCPXxN_shifts_chips(
int vectorN,
int elementN,
float rem_carrier_phase_in_rad,
float phase_step_rad
)
float phase_step_rad)
{
//Accumulators cache
__shared__ GPU_Complex accumResult[ACCUM_N];
// CUDA version of floating point NCO and vector dot product integrated
// CUDA version of floating point NCO and vector dot product integrated
float sin;
float cos;
for (int i = blockIdx.x * blockDim.x + threadIdx.x;
i < elementN;
i += blockDim.x * gridDim.x)
{
__sincosf(rem_carrier_phase_in_rad + i*phase_step_rad, &sin, &cos);
d_sig_wiped[i] = d_sig_in[i] * GPU_Complex(cos,-sin);
}
{
__sincosf(rem_carrier_phase_in_rad + i * phase_step_rad, &sin, &cos);
d_sig_wiped[i] = d_sig_in[i] * GPU_Complex(cos, -sin);
}
__syncthreads();
////////////////////////////////////////////////////////////////////////////
@@ -77,273 +75,279 @@ __global__ void Doppler_wippe_scalarProdGPUCPXxN_shifts_chips(
// from total number of thread blocks
////////////////////////////////////////////////////////////////////////////
for (int vec = blockIdx.x; vec < vectorN; vec += gridDim.x)
{
//int vectorBase = IMUL(elementN, vec);
//int vectorEnd = elementN;
////////////////////////////////////////////////////////////////////////
// Each accumulator cycles through vectors with
// stride equal to number of total number of accumulators ACCUM_N
// At this stage ACCUM_N is only preferred be a multiple of warp size
// to meet memory coalescing alignment constraints.
////////////////////////////////////////////////////////////////////////
for (int iAccum = threadIdx.x; iAccum < ACCUM_N; iAccum += blockDim.x)
{
GPU_Complex sum = GPU_Complex(0,0);
float local_code_chip_index=0.0;;
//float code_phase;
for (int pos = iAccum; pos < elementN; pos += ACCUM_N)
{
//original sample code
//sum = sum + d_sig_in[pos-vectorBase] * d_nco_in[pos-vectorBase] * d_local_codes_in[pos];
//sum = sum + d_sig_in[pos-vectorBase] * d_local_codes_in[pos];
//sum.multiply_acc(d_sig_in[pos],d_local_codes_in[pos+d_shifts_samples[vec]]);
//int vectorBase = IMUL(elementN, vec);
//int vectorEnd = elementN;
//custom code for multitap correlator
// 1.resample local code for the current shift
////////////////////////////////////////////////////////////////////////
// Each accumulator cycles through vectors with
// stride equal to number of total number of accumulators ACCUM_N
// At this stage ACCUM_N is only preferred be a multiple of warp size
// to meet memory coalescing alignment constraints.
////////////////////////////////////////////////////////////////////////
for (int iAccum = threadIdx.x; iAccum < ACCUM_N; iAccum += blockDim.x)
{
GPU_Complex sum = GPU_Complex(0, 0);
float local_code_chip_index = 0.0;
;
//float code_phase;
for (int pos = iAccum; pos < elementN; pos += ACCUM_N)
{
//original sample code
//sum = sum + d_sig_in[pos-vectorBase] * d_nco_in[pos-vectorBase] * d_local_codes_in[pos];
//sum = sum + d_sig_in[pos-vectorBase] * d_local_codes_in[pos];
//sum.multiply_acc(d_sig_in[pos],d_local_codes_in[pos+d_shifts_samples[vec]]);
local_code_chip_index= fmodf(code_phase_step_chips*__int2float_rd(pos)+ d_shifts_chips[vec] - rem_code_phase_chips, code_length_chips);
//custom code for multitap correlator
// 1.resample local code for the current shift
//Take into account that in multitap correlators, the shifts can be negative!
if (local_code_chip_index<0.0) local_code_chip_index+=(code_length_chips-1);
//printf("vec= %i, pos %i, chip_idx=%i chip_shift=%f \r\n",vec, pos,__float2int_rd(local_code_chip_index),local_code_chip_index);
// 2.correlate
sum.multiply_acc(d_sig_wiped[pos],d_local_code_in[__float2int_rd(local_code_chip_index)]);
local_code_chip_index = fmodf(code_phase_step_chips * __int2float_rd(pos) + d_shifts_chips[vec] - rem_code_phase_chips, code_length_chips);
}
accumResult[iAccum] = sum;
//Take into account that in multitap correlators, the shifts can be negative!
if (local_code_chip_index < 0.0) local_code_chip_index += (code_length_chips - 1);
//printf("vec= %i, pos %i, chip_idx=%i chip_shift=%f \r\n",vec, pos,__float2int_rd(local_code_chip_index),local_code_chip_index);
// 2.correlate
sum.multiply_acc(d_sig_wiped[pos], d_local_code_in[__float2int_rd(local_code_chip_index)]);
}
accumResult[iAccum] = sum;
}
////////////////////////////////////////////////////////////////////////
// Perform tree-like reduction of accumulators' results.
// ACCUM_N has to be power of two at this stage
////////////////////////////////////////////////////////////////////////
for (int stride = ACCUM_N / 2; stride > 0; stride >>= 1)
{
__syncthreads();
for (int iAccum = threadIdx.x; iAccum < stride; iAccum += blockDim.x)
{
accumResult[iAccum] += accumResult[stride + iAccum];
}
}
if (threadIdx.x == 0)
{
d_corr_out[vec] = accumResult[0];
}
}
////////////////////////////////////////////////////////////////////////
// Perform tree-like reduction of accumulators' results.
// ACCUM_N has to be power of two at this stage
////////////////////////////////////////////////////////////////////////
for (int stride = ACCUM_N / 2; stride > 0; stride >>= 1)
{
__syncthreads();
for (int iAccum = threadIdx.x; iAccum < stride; iAccum += blockDim.x)
{
accumResult[iAccum] += accumResult[stride + iAccum];
}
}
if (threadIdx.x == 0)
{
d_corr_out[vec] = accumResult[0];
}
}
}
bool cuda_multicorrelator::init_cuda_integrated_resampler(
int signal_length_samples,
int code_length_chips,
int n_correlators
)
int signal_length_samples,
int code_length_chips,
int n_correlators)
{
// use command-line specified CUDA device, otherwise use device with highest Gflops/s
// findCudaDevice(argc, (const char **)argv);
cudaDeviceProp prop;
// use command-line specified CUDA device, otherwise use device with highest Gflops/s
// findCudaDevice(argc, (const char **)argv);
cudaDeviceProp prop;
int num_devices, device;
cudaGetDeviceCount(&num_devices);
if (num_devices > 1) {
int max_multiprocessors = 0, max_device = 0;
for (device = 0; device < num_devices; device++) {
cudaDeviceProp properties;
cudaGetDeviceProperties(&properties, device);
if (max_multiprocessors < properties.multiProcessorCount) {
max_multiprocessors = properties.multiProcessorCount;
max_device = device;
}
printf("Found GPU device # %i\n",device);
}
//cudaSetDevice(max_device);
if (num_devices > 1)
{
int max_multiprocessors = 0, max_device = 0;
for (device = 0; device < num_devices; device++)
{
cudaDeviceProp properties;
cudaGetDeviceProperties(&properties, device);
if (max_multiprocessors < properties.multiProcessorCount)
{
max_multiprocessors = properties.multiProcessorCount;
max_device = device;
}
printf("Found GPU device # %i\n", device);
}
//cudaSetDevice(max_device);
//set random device!
selected_gps_device=rand() % num_devices;//generates a random number between 0 and num_devices to split the threads between GPUs
cudaSetDevice(selected_gps_device);
//set random device!
selected_gps_device = rand() % num_devices; //generates a random number between 0 and num_devices to split the threads between GPUs
cudaSetDevice(selected_gps_device);
cudaGetDeviceProperties( &prop, max_device );
//debug code
if (prop.canMapHostMemory != 1) {
printf( "Device can not map memory.\n" );
}
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("multiProcessorCount= %i \n",prop.multiProcessorCount);
}else{
cudaGetDevice( &selected_gps_device);
cudaGetDeviceProperties( &prop, selected_gps_device );
//debug code
if (prop.canMapHostMemory != 1) {
printf( "Device can not map memory.\n" );
}
cudaGetDeviceProperties(&prop, max_device);
//debug code
if (prop.canMapHostMemory != 1)
{
printf("Device can not map memory.\n");
}
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("multiProcessorCount= %i \n", prop.multiProcessorCount);
}
else
{
cudaGetDevice(&selected_gps_device);
cudaGetDeviceProperties(&prop, selected_gps_device);
//debug code
if (prop.canMapHostMemory != 1)
{
printf("Device can not map memory.\n");
}
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("multiProcessorCount= %i \n",prop.multiProcessorCount);
}
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("multiProcessorCount= %i \n", prop.multiProcessorCount);
}
// (cudaFuncSetCacheConfig(CUDA_32fc_x2_multiply_x2_dot_prod_32fc_, cudaFuncCachePreferShared));
// (cudaFuncSetCacheConfig(CUDA_32fc_x2_multiply_x2_dot_prod_32fc_, cudaFuncCachePreferShared));
// ALLOCATE GPU MEMORY FOR INPUT/OUTPUT and INTERNAL vectors
size_t size = signal_length_samples * sizeof(GPU_Complex);
//********* ZERO COPY VERSION ************
// Set flag to enable zero copy access
//********* ZERO COPY VERSION ************
// Set flag to enable zero copy access
// Optimal in shared memory devices (like Jetson K1)
//cudaSetDeviceFlags(cudaDeviceMapHost);
//cudaSetDeviceFlags(cudaDeviceMapHost);
//******** CudaMalloc version ***********
//******** CudaMalloc version ***********
// input signal GPU memory (can be mapped to CPU memory in shared memory devices!)
// cudaMalloc((void **)&d_sig_in, size);
// cudaMemset(d_sig_in,0,size);
// input signal GPU memory (can be mapped to CPU memory in shared memory devices!)
// cudaMalloc((void **)&d_sig_in, size);
// cudaMemset(d_sig_in,0,size);
// Doppler-free signal (internal GPU memory)
cudaMalloc((void **)&d_sig_doppler_wiped, size);
cudaMemset(d_sig_doppler_wiped,0,size);
// Doppler-free signal (internal GPU memory)
cudaMalloc((void **)&d_sig_doppler_wiped, size);
cudaMemset(d_sig_doppler_wiped, 0, size);
// Local code GPU memory (can be mapped to CPU memory in shared memory devices!)
cudaMalloc((void **)&d_local_codes_in, sizeof(std::complex<float>)*code_length_chips);
cudaMemset(d_local_codes_in,0,sizeof(std::complex<float>)*code_length_chips);
// Local code GPU memory (can be mapped to CPU memory in shared memory devices!)
cudaMalloc((void **)&d_local_codes_in, sizeof(std::complex<float>) * code_length_chips);
cudaMemset(d_local_codes_in, 0, sizeof(std::complex<float>) * code_length_chips);
d_code_length_chips=code_length_chips;
d_code_length_chips = code_length_chips;
// Vector with the chip shifts for each correlator tap
// Vector with the chip shifts for each correlator tap
//GPU memory (can be mapped to CPU memory in shared memory devices!)
cudaMalloc((void **)&d_shifts_chips, sizeof(float)*n_correlators);
cudaMemset(d_shifts_chips,0,sizeof(float)*n_correlators);
cudaMalloc((void **)&d_shifts_chips, sizeof(float) * n_correlators);
cudaMemset(d_shifts_chips, 0, sizeof(float) * n_correlators);
//scalars
//cudaMalloc((void **)&d_corr_out, sizeof(std::complex<float>)*n_correlators);
//cudaMemset(d_corr_out,0,sizeof(std::complex<float>)*n_correlators);
//scalars
//cudaMalloc((void **)&d_corr_out, sizeof(std::complex<float>)*n_correlators);
//cudaMemset(d_corr_out,0,sizeof(std::complex<float>)*n_correlators);
// Launch the Vector Add CUDA Kernel
// TODO: write a smart load balance using device info!
threadsPerBlock = 64;
blocksPerGrid = 128;//(int)(signal_length_samples+threadsPerBlock-1)/threadsPerBlock;
threadsPerBlock = 64;
blocksPerGrid = 128; //(int)(signal_length_samples+threadsPerBlock-1)/threadsPerBlock;
cudaStreamCreate (&stream1) ;
//cudaStreamCreate (&stream2) ;
return true;
cudaStreamCreate(&stream1);
//cudaStreamCreate (&stream2) ;
return true;
}
bool cuda_multicorrelator::set_local_code_and_taps(
int code_length_chips,
const std::complex<float>* local_codes_in,
float *shifts_chips,
int n_correlators
)
int code_length_chips,
const std::complex<float> *local_codes_in,
float *shifts_chips,
int n_correlators)
{
cudaSetDevice(selected_gps_device);
//********* ZERO COPY VERSION ************
// // Get device pointer from host memory. No allocation or memcpy
// cudaError_t code;
// // local code CPU -> GPU copy memory
// code=cudaHostGetDevicePointer((void **)&d_local_codes_in, (void *) local_codes_in, 0);
// if (code!=cudaSuccess)
// {
// printf("cuda cudaHostGetDevicePointer error in set_local_code_and_taps \r\n");
// }
// // Correlator shifts vector CPU -> GPU copy memory (fractional chip shifts are allowed!)
// code=cudaHostGetDevicePointer((void **)&d_shifts_chips, (void *) shifts_chips, 0);
// if (code!=cudaSuccess)
// {
// printf("cuda cudaHostGetDevicePointer error in set_local_code_and_taps \r\n");
// }
cudaSetDevice(selected_gps_device);
//********* ZERO COPY VERSION ************
// // Get device pointer from host memory. No allocation or memcpy
// cudaError_t code;
// // local code CPU -> GPU copy memory
// code=cudaHostGetDevicePointer((void **)&d_local_codes_in, (void *) local_codes_in, 0);
// if (code!=cudaSuccess)
// {
// printf("cuda cudaHostGetDevicePointer error in set_local_code_and_taps \r\n");
// }
// // Correlator shifts vector CPU -> GPU copy memory (fractional chip shifts are allowed!)
// code=cudaHostGetDevicePointer((void **)&d_shifts_chips, (void *) shifts_chips, 0);
// if (code!=cudaSuccess)
// {
// printf("cuda cudaHostGetDevicePointer error in set_local_code_and_taps \r\n");
// }
//******** CudaMalloc version ***********
//******** CudaMalloc version ***********
//local code CPU -> GPU copy memory
cudaMemcpyAsync(d_local_codes_in, local_codes_in, sizeof(GPU_Complex)*code_length_chips, cudaMemcpyHostToDevice,stream1);
d_code_length_chips=code_length_chips;
cudaMemcpyAsync(d_local_codes_in, local_codes_in, sizeof(GPU_Complex) * code_length_chips, cudaMemcpyHostToDevice, stream1);
d_code_length_chips = code_length_chips;
//Correlator shifts vector CPU -> GPU copy memory (fractional chip shifts are allowed!)
cudaMemcpyAsync(d_shifts_chips, shifts_chips, sizeof(float)*n_correlators,
cudaMemcpyHostToDevice,stream1);
cudaMemcpyAsync(d_shifts_chips, shifts_chips, sizeof(float) * n_correlators,
cudaMemcpyHostToDevice, stream1);
return true;
return true;
}
bool cuda_multicorrelator::set_input_output_vectors(
std::complex<float>* corr_out,
std::complex<float>* sig_in
)
std::complex<float> *corr_out,
std::complex<float> *sig_in)
{
cudaSetDevice(selected_gps_device);
// Save CPU pointers
d_sig_in_cpu = sig_in;
d_corr_out_cpu = corr_out;
cudaSetDevice(selected_gps_device);
// Save CPU pointers
d_sig_in_cpu =sig_in;
d_corr_out_cpu = corr_out;
// Zero Copy version
// Get device pointer from host memory. No allocation or memcpy
cudaError_t code;
code=cudaHostGetDevicePointer((void **)&d_sig_in, (void *) sig_in, 0);
code=cudaHostGetDevicePointer((void **)&d_corr_out, (void *) corr_out, 0);
if (code!=cudaSuccess)
{
printf("cuda cudaHostGetDevicePointer error \r\n");
}
return true;
// Zero Copy version
// Get device pointer from host memory. No allocation or memcpy
cudaError_t code;
code = cudaHostGetDevicePointer((void **)&d_sig_in, (void *)sig_in, 0);
code = cudaHostGetDevicePointer((void **)&d_corr_out, (void *)corr_out, 0);
if (code != cudaSuccess)
{
printf("cuda cudaHostGetDevicePointer error \r\n");
}
return true;
}
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
#define gpuErrchk(ans) \
{ \
gpuAssert((ans), __FILE__, __LINE__); \
}
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort = true)
{
if (code != cudaSuccess)
{
fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
if (abort) exit(code);
}
if (code != cudaSuccess)
{
fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
if (abort) exit(code);
}
}
bool cuda_multicorrelator::Carrier_wipeoff_multicorrelator_resampler_cuda(
float rem_carrier_phase_in_rad,
float phase_step_rad,
float code_phase_step_chips,
float rem_code_phase_chips,
int signal_length_samples,
int n_correlators)
{
cudaSetDevice(selected_gps_device);
// cudaMemCpy version
//size_t memSize = signal_length_samples * sizeof(std::complex<float>);
// input signal CPU -> GPU copy memory
float rem_carrier_phase_in_rad,
float phase_step_rad,
float code_phase_step_chips,
float rem_code_phase_chips,
int signal_length_samples,
int n_correlators)
{
cudaSetDevice(selected_gps_device);
// cudaMemCpy version
//size_t memSize = signal_length_samples * sizeof(std::complex<float>);
// input signal CPU -> GPU copy memory
//cudaMemcpyAsync(d_sig_in, d_sig_in_cpu, memSize,
// cudaMemcpyHostToDevice, stream2);
//***** NOTICE: NCO is computed on-the-fly, not need to copy NCO into GPU! ****
//launch the multitap correlator with integrated local code resampler!
Doppler_wippe_scalarProdGPUCPXxN_shifts_chips<<<blocksPerGrid, threadsPerBlock,0 ,stream1>>>(
d_corr_out,
d_sig_in,
d_sig_doppler_wiped,
d_local_codes_in,
d_shifts_chips,
d_code_length_chips,
code_phase_step_chips,
rem_code_phase_chips,
n_correlators,
signal_length_samples,
rem_carrier_phase_in_rad,
phase_step_rad
);
Doppler_wippe_scalarProdGPUCPXxN_shifts_chips<<<blocksPerGrid, threadsPerBlock, 0, stream1>>>(
d_corr_out,
d_sig_in,
d_sig_doppler_wiped,
d_local_codes_in,
d_shifts_chips,
d_code_length_chips,
code_phase_step_chips,
rem_code_phase_chips,
n_correlators,
signal_length_samples,
rem_carrier_phase_in_rad,
phase_step_rad);
gpuErrchk( cudaPeekAtLastError() );
gpuErrchk( cudaStreamSynchronize(stream1));
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaStreamSynchronize(stream1));
// cudaMemCpy version
// cudaMemCpy version
// Copy the device result vector in device memory to the host result vector
// in host memory.
//scalar products (correlators outputs)
@@ -352,37 +356,38 @@ bool cuda_multicorrelator::Carrier_wipeoff_multicorrelator_resampler_cuda(
return true;
}
cuda_multicorrelator::cuda_multicorrelator()
{
d_sig_in=NULL;
d_nco_in=NULL;
d_sig_doppler_wiped=NULL;
d_local_codes_in=NULL;
d_shifts_samples=NULL;
d_shifts_chips=NULL;
d_corr_out=NULL;
threadsPerBlock=0;
blocksPerGrid=0;
d_code_length_chips=0;
d_sig_in = NULL;
d_nco_in = NULL;
d_sig_doppler_wiped = NULL;
d_local_codes_in = NULL;
d_shifts_samples = NULL;
d_shifts_chips = NULL;
d_corr_out = NULL;
threadsPerBlock = 0;
blocksPerGrid = 0;
d_code_length_chips = 0;
}
bool cuda_multicorrelator::free_cuda()
{
// Free device global memory
if (d_sig_in!=NULL) cudaFree(d_sig_in);
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);
// Free device global memory
if (d_sig_in != NULL) cudaFree(d_sig_in);
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();
return true;
cudaDeviceReset();
return true;
}

6
src/algorithms/tracking/libs/cuda_multicorrelator.h
View File

@@ -1,11 +1,11 @@
/*!
* \file cuda_multicorrelator.h
* \brief High optimized CUDA GPU vector multiTAP correlator class
* \brief Highly optimized CUDA GPU vector multiTAP correlator class
* \authors <ul>
* <li> Javier Arribas, 2015. jarribas(at)cttc.es
* </ul>
*
* Class that implements a high optimized vector multiTAP correlator class for NVIDIA CUDA GPUs
* Class that implements a highly optimized vector multiTAP correlator class for NVIDIA CUDA GPUs
*
* -------------------------------------------------------------------------
*
@@ -92,6 +92,7 @@ struct GPU_Complex
}
};
struct GPU_Complex_Short
{
float r;
@@ -149,7 +150,6 @@ private:
GPU_Complex* d_local_codes_in;
GPU_Complex* d_corr_out;
//
std::complex<float>* d_sig_in_cpu;
std::complex<float>* d_corr_out_cpu;

7
src/algorithms/tracking/libs/dll_pll_conf.cc
View File

@@ -43,6 +43,13 @@ Dll_Pll_Conf::Dll_Pll_Conf()
dump = false;
dump_mat = true;
dump_filename = std::string("./dll_pll_dump.dat");
enable_fll_pull_in = false;
enable_fll_steady_state = false;
pull_in_time_s = 2;
fll_filter_order = 1;
pll_filter_order = 3;
dll_filter_order = 2;
fll_bw_hz = 35.0;
pll_pull_in_bw_hz = 50.0;
dll_pull_in_bw_hz = 3.0;
pll_bw_hz = 35.0;

9
src/algorithms/tracking/libs/dll_pll_conf.h
View File

@@ -40,6 +40,14 @@ class Dll_Pll_Conf
{
public:
/* DLL/PLL tracking configuration */
int fll_filter_order;
bool enable_fll_pull_in;
bool enable_fll_steady_state;
unsigned int pull_in_time_s;
int pll_filter_order;
int dll_filter_order;
double fs_in;
uint32_t vector_length;
bool dump;
@@ -47,6 +55,7 @@ public:
std::string dump_filename;
float pll_pull_in_bw_hz;
float dll_pull_in_bw_hz;
float fll_bw_hz;
float pll_bw_hz;
float dll_bw_hz;
float pll_bw_narrow_hz;

51
src/algorithms/tracking/libs/tracking_FLL_PLL_filter.cc
View File

@@ -32,7 +32,26 @@
*/
#include "tracking_FLL_PLL_filter.h"
#include <iostream>
Tracking_FLL_PLL_filter::Tracking_FLL_PLL_filter()
{
d_order = 0;
d_pll_w = 0.0;
d_pll_w0p3 = 0.0;
d_pll_w0f2 = 0.0;
d_pll_x = 0.0;
d_pll_a2 = 0.0;
d_pll_w0f = 0.0;
d_pll_a3 = 0.0;
d_pll_w0p2 = 0.0;
d_pll_b3 = 0.0;
d_pll_w0p = 0.0;
}
Tracking_FLL_PLL_filter::~Tracking_FLL_PLL_filter() = default;
void Tracking_FLL_PLL_filter::set_params(float fll_bw_hz, float pll_bw_hz, int order)
{
@@ -89,7 +108,7 @@ float Tracking_FLL_PLL_filter::get_carrier_error(float FLL_discriminator, float
if (d_order == 3)
{
/*
* 3rd order PLL with 2nd order FLL assist
* 3rd order PLL with 2nd order FLL assist
*/
d_pll_w = d_pll_w + correlation_time_s * (d_pll_w0p3 * PLL_discriminator + d_pll_w0f2 * FLL_discriminator);
d_pll_x = d_pll_x + correlation_time_s * (0.5 * d_pll_w + d_pll_a2 * d_pll_w0f * FLL_discriminator + d_pll_a3 * d_pll_w0p2 * PLL_discriminator);
@@ -104,31 +123,11 @@ float Tracking_FLL_PLL_filter::get_carrier_error(float FLL_discriminator, float
pll_w_new = d_pll_w + PLL_discriminator * d_pll_w0p2 * correlation_time_s + FLL_discriminator * d_pll_w0f * correlation_time_s;
carrier_error_hz = 0.5 * (pll_w_new + d_pll_w) + d_pll_a2 * d_pll_w0p * PLL_discriminator;
d_pll_w = pll_w_new;
/*std::cout<<" d_pll_w = "<<carrier_error_hz<<
", pll_w_new = "<<pll_w_new
<<", PLL_discriminator=" <<PLL_discriminator
<<" FLL_discriminator ="<<FLL_discriminator
<<" correlation_time_s = "<<correlation_time_s<<"\r\n";*/
/* std::cout << " d_pll_w = " << carrier_error_hz << ", pll_w_new = " << pll_w_new
<< ", PLL_discriminator=" << PLL_discriminator
<< " FLL_discriminator =" << FLL_discriminator
<< " correlation_time_s = " << correlation_time_s << "\r\n"; */
}
return carrier_error_hz;
}
Tracking_FLL_PLL_filter::Tracking_FLL_PLL_filter()
{
d_order = 0;
d_pll_w = 0;
d_pll_w0p3 = 0;
d_pll_w0f2 = 0;
d_pll_x = 0;
d_pll_a2 = 0;
d_pll_w0f = 0;
d_pll_a3 = 0;
d_pll_w0p2 = 0;
d_pll_b3 = 0;
d_pll_w0p = 0;
}
Tracking_FLL_PLL_filter::~Tracking_FLL_PLL_filter() = default;

3
src/algorithms/tracking/libs/tracking_discriminators.cc
View File

@@ -36,6 +36,7 @@
#include <cmath>
// All the outputs are in RADIANS
/*
* FLL four quadrant arctan discriminator:
* \f{equation}
@@ -45,7 +46,6 @@
* \f$I_{PS1},Q_{PS1}\f$ are the inphase and quadrature prompt correlator outputs respectively at sample time \f$t_1\f$, and
* \f$I_{PS2},Q_{PS2}\f$ are the inphase and quadrature prompt correlator outputs respectively at sample time \f$t_2\f$. The output is in [radians/second].
*/
double fll_four_quadrant_atan(gr_complex prompt_s1, gr_complex prompt_s2, double t1, double t2)
{
double cross, dot;
@@ -105,6 +105,7 @@ double dll_nc_e_minus_l_normalized(gr_complex early_s1, gr_complex late_s1)
return 0.5 * (P_early - P_late) / (P_early + P_late);
}
/*
* DLL Noncoherent Very Early Minus Late Power (VEMLP) normalized discriminator, using the outputs
* of four correlators, Very Early (VE), Early (E), Late (L) and Very Late (VL):

18
src/algorithms/tracking/libs/tracking_loop_filter.cc
View File

@@ -4,7 +4,7 @@
* \author Cillian O'Driscoll, 2015. cillian.odriscoll(at)gmail.com
*
* Class implementing a generic 1st, 2nd or 3rd order loop filter. Based
* on the bilinear transform of the standard Weiner filter.
* on the bilinear transform of the standard Wiener filter.
*
* -------------------------------------------------------------------------
*
@@ -36,6 +36,8 @@
#include <glog/logging.h>
#include <cmath>
const int MAX_LOOP_ORDER = 3;
const int MAX_LOOP_HISTORY_LENGTH = 4;
Tracking_loop_filter::Tracking_loop_filter(float update_interval,
float noise_bandwidth,
@@ -74,7 +76,7 @@ float Tracking_loop_filter::apply(float current_input)
// Now apply the filter coefficients:
float result = 0.0;
// Hanlde the old outputs first:
// Handle the old outputs first:
for (unsigned int ii = 0; ii < d_output_coefficients.size(); ++ii)
{
result += d_output_coefficients[ii] * d_outputs[(d_current_index + ii) % MAX_LOOP_HISTORY_LENGTH];
@@ -95,16 +97,13 @@ float Tracking_loop_filter::apply(float current_input)
d_inputs[d_current_index] = current_input;
for (unsigned int ii = 0; ii < d_input_coefficients.size(); ++ii)
{
result += d_input_coefficients[ii] * d_inputs[(d_current_index + ii) % MAX_LOOP_HISTORY_LENGTH];
}
d_outputs[d_current_index] = result;
return result;
}
@@ -179,7 +178,6 @@ void Tracking_loop_filter::update_coefficients(void)
d_output_coefficients[0] = 1.0;
}
break;
case 3:
wn = d_noise_bandwidth / 0.7845; // From Kaplan
float a3 = 1.1;
@@ -208,7 +206,6 @@ void Tracking_loop_filter::update_coefficients(void)
d_input_coefficients[1] = g1 * T * T / 2.0 - 2.0 * g3;
d_input_coefficients[2] = g3 + T / 2.0 * (-g2 + T / 2.0 * g1);
d_output_coefficients.resize(2);
d_output_coefficients[0] = 2.0;
d_output_coefficients[1] = -1.0;
@@ -260,10 +257,9 @@ void Tracking_loop_filter::set_order(int loop_order)
{
if (loop_order < 1 or loop_order > MAX_LOOP_ORDER)
{
LOG(ERROR) << "Ignoring attempt to set loop order to " << loop_order
<< ". Maximum allowed order is: " << MAX_LOOP_ORDER
<< ". Not changing current value of " << d_loop_order;
LOG(WARNING) << "Ignoring attempt to set loop order to " << loop_order
<< ". Maximum allowed order is: " << MAX_LOOP_ORDER
<< ". Not changing current value of " << d_loop_order;
return;
}

5
src/algorithms/tracking/libs/tracking_loop_filter.h
View File

@@ -4,7 +4,7 @@
* \author Cillian O'Driscoll, 2015. cillian.odriscoll(at)gmail.com
*
* Class implementing a generic 1st, 2nd or 3rd order loop filter. Based
* on the bilinear transform of the standard Weiner filter.
* on the bilinear transform of the standard Wiener filter.
*
* -------------------------------------------------------------------------
*
@@ -33,8 +33,6 @@
#ifndef GNSS_SDR_TRACKING_LOOP_FILTER_H_
#define GNSS_SDR_TRACKING_LOOP_FILTER_H_
#define MAX_LOOP_ORDER 3
#define MAX_LOOP_HISTORY_LENGTH 4
#include <vector>
@@ -74,7 +72,6 @@ private:
// Compute the filter coefficients:
void update_coefficients(void);
public:
float get_noise_bandwidth(void) const;
float get_update_interval(void) const;