mirrors/gnss-sdr
1
0
Fork 0
mirror of https://github.com/gnss-sdr/gnss-sdr synced 2024-12-13 11:40:33 +00:00
Code Issues Releases Wiki Activity

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

Browse Source
This commit is contained in:
Marc Majoral 2021-12-20 10:51:15 +01:00
commit 39fd3f1ffa
37 changed files with 953 additions and 1029 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

409
.clang-tidy
View File

@ -56,6 +56,7 @@ Checks: '-*,
clang-analyzer-optin.performance.*,
clang-analyzer-optin.portability.UnixAPI,
clang-analyzer-security.*,
cppcoreguidelines-prefer-member-initializer,
cppcoreguidelines-pro-type-cstyle-cast,
cppcoreguidelines-pro-type-static-cast-downcast,
cppcoreguidelines-slicing,
@ -120,70 +121,16 @@ HeaderFilterRegex: ''
AnalyzeTemporaryDtors: false
FormatStyle: 'file'
CheckOptions:
- key: bugprone-argument-comment.CommentBoolLiterals
value: '0'
- key: bugprone-argument-comment.CommentCharacterLiterals
value: '0'
- key: bugprone-argument-comment.CommentFloatLiterals
value: '0'
- key: bugprone-argument-comment.CommentIntegerLiterals
value: '0'
- key: bugprone-argument-comment.CommentNullPtrs
value: '0'
- key: bugprone-argument-comment.CommentStringLiterals
value: '0'
- key: bugprone-argument-comment.CommentUserDefinedLiterals
value: '0'
- key: bugprone-argument-comment.IgnoreSingleArgument
value: '0'
- key: bugprone-argument-comment.StrictMode
value: '0'
- key: bugprone-assert-side-effect.AssertMacros
value: assert
- key: bugprone-assert-side-effect.CheckFunctionCalls
value: 'false'
- key: bugprone-dangling-handle.HandleClasses
value: 'std::basic_string_view;std::experimental::basic_string_view'
- key: bugprone-exception-escape.FunctionsThatShouldNotThrow
value: ''
- key: bugprone-exception-escape.IgnoredExceptions
value: ''
- key: bugprone-misplaced-widening-cast.CheckImplicitCasts
value: 'false'
- key: bugprone-reserved-identifier.AggressiveDependentMemberLookup
value: 'false'
- key: bugprone-reserved-identifier.AllowedIdentifiers
- key: performance-unnecessary-copy-initialization.ExcludedContainerTypes
value: ''
- key: modernize-replace-auto-ptr.IncludeStyle
value: llvm
- key: performance-move-const-arg.CheckTriviallyCopyableMove
value: 'true'
- key: modernize-use-auto.MinTypeNameLength
value: '5'
- key: bugprone-reserved-identifier.Invert
value: 'false'
- key: bugprone-sizeof-expression.WarnOnSizeOfCompareToConstant
value: 'true'
- key: bugprone-sizeof-expression.WarnOnSizeOfConstant
value: 'true'
- key: bugprone-sizeof-expression.WarnOnSizeOfIntegerExpression
value: 'false'
- key: bugprone-sizeof-expression.WarnOnSizeOfThis
value: 'true'
- key: bugprone-string-constructor.LargeLengthThreshold
value: '8388608'
- key: bugprone-string-constructor.WarnOnLargeLength
value: 'true'
- key: bugprone-suspicious-enum-usage.StrictMode
value: 'false'
- key: bugprone-suspicious-missing-comma.MaxConcatenatedTokens
value: '5'
- key: bugprone-suspicious-missing-comma.RatioThreshold
value: '0.200000'
- key: bugprone-suspicious-missing-comma.SizeThreshold
value: '5'
- key: bugprone-suspicious-string-compare.StringCompareLikeFunctions
value: ''
- key: bugprone-suspicious-string-compare.WarnOnImplicitComparison
value: 'true'
- key: bugprone-suspicious-string-compare.WarnOnLogicalNotComparison
value: 'false'
- key: bugprone-too-small-loop-variable.MagnitudeBitsUpperLimit
value: '16'
- key: bugprone-unused-return-value.CheckedFunctions
value: 'std::async;
std::launder;
@ -276,148 +223,218 @@ std::vector::at;
::ttyname'
- key: cert-dcl16-c.NewSuffixes
value: 'L;LL;LU;LLU'
- key: cert-msc51-cpp.DisallowedSeedTypes
value: 'time_t,std::time_t'
- key: cert-oop54-cpp.WarnOnlyIfThisHasSuspiciousField
value: '0'
- key: cert-str34-c.DiagnoseSignedUnsignedCharComparisons
value: '0'
- key: cppcoreguidelines-explicit-virtual-functions.IgnoreDestructors
value: '1'
- key: cppcoreguidelines-non-private-member-variables-in-classes.IgnoreClassesWithAllMemberVariablesBeingPublic
value: '1'
- key: cppcoreguidelines-special-member-functions.AllowMissingMoveFunctions
value: 'false'
- key: cppcoreguidelines-special-member-functions.AllowMissingMoveFunctionsWhenCopyIsDeleted
value: 'false'
- key: cppcoreguidelines-special-member-functions.AllowSoleDefaultDtor
value: 'false'
- key: google-build-namespaces.HeaderFileExtensions
value: ';h;hh;hpp;hxx'
- key: google-global-names-in-headers.HeaderFileExtensions
value: ';h;hh;hpp;hxx'
- key: google-readability-braces-around-statements.ShortStatementLines
value: '1'
- key: google-readability-function-size.StatementThreshold
value: '800'
- key: google-readability-namespace-comments.ShortNamespaceLines
value: '10'
- key: google-readability-namespace-comments.SpacesBeforeComments
value: '2'
- key: google-runtime-int.SignedTypePrefix
value: int
- key: google-runtime-int.TypeSuffix
value: _t
- key: google-runtime-int.UnsignedTypePrefix
value: uint
- key: llvm-else-after-return.WarnOnConditionVariables
value: '0'
- key: llvm-else-after-return.WarnOnUnfixable
value: '0'
- key: llvm-qualified-auto.AddConstToQualified
value: '0'
- key: misc-throw-by-value-catch-by-reference.CheckThrowTemporaries
value: 'true'
- key: misc-throw-by-value-catch-by-reference.MaxSize
value: '64'
- key: misc-throw-by-value-catch-by-reference.WarnOnLargeObjects
value: 'false'
- key: modernize-avoid-bind.PermissiveParameterList
value: 'false'
- key: modernize-loop-convert.MaxCopySize
value: '16'
- key: modernize-loop-convert.MinConfidence
value: reasonable
- key: modernize-loop-convert.NamingStyle
value: CamelCase
- key: modernize-pass-by-value.IncludeStyle
value: llvm
- key: modernize-raw-string-literal.DelimiterStem
value: lit
- key: modernize-raw-string-literal.ReplaceShorterLiterals
value: 'false'
- key: modernize-replace-auto-ptr.IncludeStyle
value: llvm
- key: modernize-use-auto.MinTypeNameLength
value: '5'
- key: modernize-use-auto.RemoveStars
value: 'false'
- key: modernize-use-bool-literals.IgnoreMacros
value: 'true'
- key: modernize-use-default-member-init.IgnoreMacros
value: 'true'
- key: modernize-use-default-member-init.UseAssignment
value: 'false'
- key: modernize-use-emplace.ContainersWithPushBack
value: 'std::vector;std::list;std::deque'
- key: modernize-use-emplace.IgnoreImplicitConstructors
value: 'false'
- key: modernize-use-emplace.SmartPointers
value: 'std::shared_ptr;std::unique_ptr;std::auto_ptr;std::weak_ptr'
- key: modernize-use-emplace.TupleMakeFunctions
value: 'std::make_pair;std::make_tuple'
- key: modernize-use-emplace.TupleTypes
value: 'std::pair;std::tuple'
- key: modernize-use-equals-default.IgnoreMacros
value: 'true'
- key: modernize-use-equals-delete.IgnoreMacros
value: 'true'
- key: modernize-use-noexcept.ReplacementString
value: ''
- key: modernize-use-noexcept.UseNoexceptFalse
value: 'true'
- key: modernize-use-nullptr.NullMacros
value: 'NULL'
- key: modernize-use-override.AllowOverrideAndFinal
value: 'false'
- key: modernize-use-override.FinalSpelling
value: final
- key: modernize-use-override.IgnoreDestructors
value: 'false'
- key: modernize-use-override.OverrideSpelling
value: override
- key: performance-faster-string-find.StringLikeClasses
value: 'std::basic_string;std::basic_string_view'
- key: performance-for-range-copy.AllowedTypes
value: ''
- key: performance-for-range-copy.WarnOnAllAutoCopies
value: 'false'
- key: performance-inefficient-string-concatenation.StrictMode
value: 'false'
- key: performance-inefficient-vector-operation.EnableProto
value: 'false'
- key: performance-inefficient-vector-operation.VectorLikeClasses
value: 'std::vector'
- key: performance-move-const-arg.CheckTriviallyCopyableMove
value: 'true'
- key: performance-move-constructor-init.IncludeStyle
value: llvm
- key: performance-type-promotion-in-math-fn.IncludeStyle
value: llvm
- key: performance-unnecessary-copy-initialization.AllowedTypes
value: ''
- key: performance-unnecessary-value-param.AllowedTypes
value: ''
- key: performance-unnecessary-value-param.IncludeStyle
value: llvm
- key: readability-braces-around-statements.ShortStatementLines
value: '0'
- key: readability-identifier-naming.AggressiveDependentMemberLookup
value: 'false'
- key: readability-identifier-naming.IgnoreFailedSplit
value: 'false'
- key: readability-identifier-naming.IgnoreMainLikeFunctions
value: 'false'
- key: readability-inconsistent-declaration-parameter-name.IgnoreMacros
- key: readability-identifier-naming.GetConfigPerFile
value: 'true'
- key: readability-inconsistent-declaration-parameter-name.Strict
value: 'false'
- key: modernize-use-default-member-init.UseAssignment
value: 'false'
- key: modernize-use-override.AllowOverrideAndFinal
value: 'false'
- key: modernize-loop-convert.IncludeStyle
value: llvm
- key: cert-str34-c.DiagnoseSignedUnsignedCharComparisons
value: 'false'
- key: misc-uniqueptr-reset-release.IncludeStyle
value: llvm
- key: bugprone-suspicious-string-compare.WarnOnLogicalNotComparison
value: 'false'
- key: readability-identifier-naming.AggressiveDependentMemberLookup
value: 'false'
- key: bugprone-suspicious-string-compare.WarnOnImplicitComparison
value: 'true'
- key: modernize-use-emplace.TupleTypes
value: 'std::pair;std::tuple'
- key: modernize-use-emplace.TupleMakeFunctions
value: 'std::make_pair;std::make_tuple'
- key: bugprone-argument-comment.CommentNullPtrs
value: '0'
- key: bugprone-argument-comment.StrictMode
value: '0'
- key: modernize-loop-convert.MakeReverseRangeHeader
value: ''
- key: modernize-use-bool-literals.IgnoreMacros
value: 'true'
- key: google-readability-namespace-comments.ShortNamespaceLines
value: '10'
- key: bugprone-suspicious-string-compare.StringCompareLikeFunctions
value: ''
- key: modernize-avoid-bind.PermissiveParameterList
value: 'false'
- key: modernize-use-override.FinalSpelling
value: final
- key: modernize-loop-convert.UseCxx20ReverseRanges
value: 'true'
- key: modernize-use-noexcept.ReplacementString
value: ''
- key: cppcoreguidelines-prefer-member-initializer.UseAssignment
value: 'false'
- key: performance-type-promotion-in-math-fn.IncludeStyle
value: llvm
- key: modernize-loop-convert.NamingStyle
value: CamelCase
- key: bugprone-suspicious-missing-comma.SizeThreshold
value: '5'
- key: readability-inconsistent-declaration-parameter-name.IgnoreMacros
value: 'true'
- key: performance-for-range-copy.WarnOnAllAutoCopies
value: 'false'
- key: bugprone-argument-comment.CommentIntegerLiterals
value: '0'
- key: google-runtime-int.UnsignedTypePrefix
value: uint
- key: modernize-loop-convert.MakeReverseRangeFunction
value: ''
- key: readability-identifier-naming.IgnoreFailedSplit
value: 'false'
- key: modernize-pass-by-value.IncludeStyle
value: llvm
- key: readability-qualified-auto.AddConstToQualified
value: 'true'
- key: readability-redundant-member-init.IgnoreBaseInCopyConstructors
value: 'false'
- key: bugprone-sizeof-expression.WarnOnSizeOfThis
value: 'true'
- key: bugprone-string-constructor.WarnOnLargeLength
value: 'true'
- key: bugprone-argument-comment.CommentCharacterLiterals
value: '0'
- key: bugprone-too-small-loop-variable.MagnitudeBitsUpperLimit
value: '16'
- key: bugprone-argument-comment.CommentFloatLiterals
value: '0'
- key: google-global-names-in-headers.HeaderFileExtensions
value: ';h;hh;hpp;hxx'
- key: readability-uppercase-literal-suffix.IgnoreMacros
value: 'true'
- key: modernize-use-nullptr.NullMacros
value: 'NULL'
- key: bugprone-suspicious-enum-usage.StrictMode
value: 'false'
- key: performance-unnecessary-copy-initialization.AllowedTypes
value: ''
- key: bugprone-suspicious-missing-comma.MaxConcatenatedTokens
value: '5'
- key: misc-throw-by-value-catch-by-reference.CheckThrowTemporaries
value: 'true'
- key: bugprone-string-constructor.LargeLengthThreshold
value: '8388608'
- key: cppcoreguidelines-special-member-functions.AllowMissingMoveFunctions
value: 'false'
- key: cert-oop54-cpp.WarnOnlyIfThisHasSuspiciousField
value: 'false'
- key: bugprone-exception-escape.FunctionsThatShouldNotThrow
value: ''
- key: performance-inefficient-vector-operation.EnableProto
value: 'false'
- key: bugprone-argument-comment.CommentStringLiterals
value: '0'
- key: modernize-loop-convert.MaxCopySize
value: '16'
- key: google-build-namespaces.HeaderFileExtensions
value: ';h;hh;hpp;hxx'
- key: cppcoreguidelines-explicit-virtual-functions.IgnoreDestructors
value: 'true'
- key: performance-for-range-copy.AllowedTypes
value: ''
- key: modernize-use-override.IgnoreDestructors
value: 'false'
- key: bugprone-sizeof-expression.WarnOnSizeOfConstant
value: 'true'
- key: modernize-use-emplace.ContainersWithPushBack
value: 'std::vector;std::list;std::deque'
- key: bugprone-argument-comment.CommentBoolLiterals
value: '0'
- key: readability-braces-around-statements.ShortStatementLines
value: '0'
- key: bugprone-argument-comment.CommentUserDefinedLiterals
value: '0'
- key: modernize-use-override.OverrideSpelling
value: override
- key: performance-inefficient-string-concatenation.StrictMode
value: 'false'
- key: google-readability-braces-around-statements.ShortStatementLines
value: '1'
- key: google-runtime-int.TypeSuffix
value: _t
- key: bugprone-reserved-identifier.AllowedIdentifiers
value: ''
- key: modernize-use-emplace.IgnoreImplicitConstructors
value: 'false'
- key: modernize-use-equals-delete.IgnoreMacros
value: 'true'
- key: bugprone-misplaced-widening-cast.CheckImplicitCasts
value: 'false'
- key: readability-uppercase-literal-suffix.NewSuffixes
value: ''
- key: modernize-loop-convert.MinConfidence
value: reasonable
- key: performance-unnecessary-value-param.AllowedTypes
value: ''
- key: bugprone-suspicious-missing-comma.RatioThreshold
value: '0.200000'
- key: cppcoreguidelines-special-member-functions.AllowMissingMoveFunctionsWhenCopyIsDeleted
value: 'false'
- key: modernize-use-noexcept.UseNoexceptFalse
value: 'true'
- key: google-readability-namespace-comments.SpacesBeforeComments
value: '2'
- key: misc-throw-by-value-catch-by-reference.MaxSize
value: '64'
- key: cppcoreguidelines-non-private-member-variables-in-classes.IgnoreClassesWithAllMemberVariablesBeingPublic
value: 'true'
- key: bugprone-argument-comment.IgnoreSingleArgument
value: '0'
- key: bugprone-sizeof-expression.WarnOnSizeOfIntegerExpression
value: 'false'
- key: performance-faster-string-find.StringLikeClasses
value: 'std::basic_string;std::basic_string_view'
- key: bugprone-assert-side-effect.CheckFunctionCalls
value: 'false'
- key: bugprone-string-constructor.StringNames
value: '::std::basic_string;::std::basic_string_view'
- key: bugprone-assert-side-effect.AssertMacros
value: assert
- key: bugprone-exception-escape.IgnoredExceptions
value: ''
- key: modernize-use-default-member-init.IgnoreMacros
value: 'true'
- key: llvm-qualified-auto.AddConstToQualified
value: 'false'
- key: readability-identifier-naming.IgnoreMainLikeFunctions
value: 'false'
- key: google-runtime-int.SignedTypePrefix
value: int
- key: google-readability-function-size.StatementThreshold
value: '800'
- key: llvm-else-after-return.WarnOnConditionVariables
value: 'false'
- key: cert-msc51-cpp.DisallowedSeedTypes
value: 'time_t,std::time_t'
- key: bugprone-sizeof-expression.WarnOnSizeOfCompareToConstant
value: 'true'
- key: bugprone-reserved-identifier.AggressiveDependentMemberLookup
value: 'false'
- key: modernize-raw-string-literal.DelimiterStem
value: lit
- key: modernize-use-equals-default.IgnoreMacros
value: 'true'
- key: misc-throw-by-value-catch-by-reference.WarnOnLargeObjects
value: 'false'
- key: cppcoreguidelines-special-member-functions.AllowSoleDefaultDtor
value: 'false'
- key: modernize-raw-string-literal.ReplaceShorterLiterals
value: 'false'
- key: modernize-use-emplace.SmartPointers
value: 'std::shared_ptr;std::unique_ptr;std::auto_ptr;std::weak_ptr'
- key: performance-inefficient-vector-operation.VectorLikeClasses
value: 'std::vector'
- key: modernize-use-auto.RemoveStars
value: 'false'
- key: bugprone-dangling-handle.HandleClasses
value: 'std::basic_string_view;std::experimental::basic_string_view'
- key: performance-unnecessary-value-param.IncludeStyle
value: llvm
- key: readability-redundant-member-init.IgnoreBaseInCopyConstructors
value: 'false'
- key: llvm-else-after-return.WarnOnUnfixable
value: 'false'
...

36
CMakeLists.txt
View File

@ -74,6 +74,8 @@ option(ENABLE_ARMA_NO_DEBUG OFF)
option(ENABLE_STRIP "Create stripped binaries without debugging symbols (in Release build mode only)" OFF)
option(Boost_USE_STATIC_LIBS "Use Boost static libs" OFF)
if(ENABLE_PACKAGING)
set(ENABLE_GENERIC_ARCH ON)
set(ENABLE_ARMA_NO_DEBUG ON)
@ -671,19 +673,19 @@ if(NOT (GNURADIO_VERSION VERSION_LESS 3.8) AND (LOG4CPP_READY_FOR_CXX17 OR GNURA
endif()
if(FILESYSTEM_FOUND)
set(CMAKE_CXX_STANDARD 17)
if(CMAKE_VERSION VERSION_GREATER 3.13)
# if(CMAKE_VERSION VERSION_GREATER 3.13)
# UHD 3.15.0.0-5 does not support C++20
# GNU Radio 3.10.0.git does not support C++20
if(((NOT UHD_FOUND) OR (UHD_FOUND AND ("${UHD_VERSION}" VERSION_LESS 3.14.99))) AND (GNURADIO_VERSION VERSION_LESS 3.9.99))
# set(CMAKE_CXX_STANDARD 20)
if(CMAKE_VERSION VERSION_GREATER 3.20.99)
if(((CMAKE_CXX_COMPILER_ID STREQUAL "GNU") AND NOT (CMAKE_CXX_COMPILER_VERSION VERSION_LESS "11.0.0")) OR
((CMAKE_CXX_COMPILER_ID STREQUAL "Clang") AND NOT (CMAKE_CXX_COMPILER_VERSION VERSION_LESS "12.0")))
# set(CMAKE_CXX_STANDARD 23)
endif()
endif()
endif()
endif()
# if(((NOT UHD_FOUND) OR (UHD_FOUND AND ("${UHD_VERSION}" VERSION_LESS 3.14.99))) AND (GNURADIO_VERSION VERSION_LESS 3.9.99))
# set(CMAKE_CXX_STANDARD 20)
# if(CMAKE_VERSION VERSION_GREATER 3.20.99)
# if(((CMAKE_CXX_COMPILER_ID STREQUAL "GNU") AND NOT (CMAKE_CXX_COMPILER_VERSION VERSION_LESS "11.0.0")) OR
# ((CMAKE_CXX_COMPILER_ID STREQUAL "Clang") AND NOT (CMAKE_CXX_COMPILER_VERSION VERSION_LESS "12.0")))
# set(CMAKE_CXX_STANDARD 23)
# endif()
# endif()
# endif()
# endif()
set(CMAKE_CXX_STANDARD_REQUIRED ON)
endif()
endif()
@ -711,9 +713,6 @@ set(Boost_ADDITIONAL_VERSIONS
"1.70.0" "1.70" "1.71.0" "1.71"
)
set(Boost_USE_MULTITHREAD ON)
#set(Boost_USE_STATIC_LIBS OFF)
option(Boost_USE_STATIC_LIBS "Use Boost static libs" OFF)
set(BOOST_COMPONENTS atomic chrono date_time serialization system thread)
if(NOT ${FILESYSTEM_FOUND})
set(BOOST_COMPONENTS ${BOOST_COMPONENTS} filesystem)
@ -1807,6 +1806,8 @@ if(NOT TARGET BLAS::BLAS)
)
endif()
################################################################################
# Check that LAPACK (Linear Algebra PACKage) is found in the system
# See https://www.netlib.org/lapack/
@ -1841,6 +1842,7 @@ if(NOT TARGET LAPACK::LAPACK)
endif()
################################################################################
# Armadillo - http://arma.sourceforge.net/
################################################################################
@ -2981,6 +2983,7 @@ else()
endif()
##########################################
# gr-limesdr - OPTIONAL
# https://github.com/myriadrf/gr-limesdr
@ -3008,6 +3011,7 @@ else()
endif()
##############################################
# gr-iio - OPTIONAL
# IIO blocks for GNU Radio
@ -3054,6 +3058,7 @@ if(ENABLE_PLUTOSDR OR ENABLE_FMCOMMS2)
endif()
#####################################################################
# libiio - OPTIONAL
# A library for interfacing with local and remote Linux IIO devices
@ -3081,6 +3086,7 @@ if(ENABLE_AD9361 OR ENABLE_FMCOMMS2)
endif()
##############################################
# TELEORBIT FLEXIBAND FRONTEND - OPTIONAL
##############################################
@ -3112,6 +3118,7 @@ if(ENABLE_FLEXIBAND)
endif()
#######################################################
# CTTC's digital array beamformer prototype - OPTIONAL
#######################################################
@ -3342,6 +3349,7 @@ add_feature_info(ENABLE_INSTALL_TESTS ENABLE_INSTALL_TESTS "Install test binarie
add_feature_info(ENABLE_BENCHMARKS ENABLE_BENCHMARKS "Enables building of code snippet benchmarks.")
add_feature_info(ENABLE_EXTERNAL_MATHJAX ENABLE_EXTERNAL_MATHJAX "Use MathJax from an external CDN in HTML docs when doing '${CMAKE_MAKE_PROGRAM_PRETTY_NAME} doc'.")
add_feature_info(ENABLE_OWN_CPUFEATURES ENABLE_OWN_CPUFEATURES "Force building own local version of the cpu_features library, even if it is already installed.")
add_feature_info(Boost_USE_STATIC_LIBS Boost_USE_STATIC_LIBS "Use Boost static libraries.")
message(STATUS "")
message(STATUS "***************************************")

2
cmake/Toolchains/arm_cortex_a53_hardfp_native.cmake
View File

@ -5,7 +5,7 @@
# SPDX-License-Identifier: BSD-3-Clause
########################################################################
# Toolchain file for building native on a ARM Cortex A72 w/ NEON
# Toolchain file for building native on a ARM Cortex A53 w/ NEON
# Usage: cmake -DCMAKE_TOOLCHAIN_FILE=<this file> <source directory>
########################################################################
set(CMAKE_CXX_COMPILER g++)

9
docs/CHANGELOG.md
View File

@ -42,7 +42,10 @@ All notable changes to GNSS-SDR will be documented in this file.
readability of the code, could potentially increase performance, and allows
for easier detection of unused data members (see
https://github.com/isocpp/CppCoreGuidelines/blob/master/CppCoreGuidelines.md/#Rc-initialize).
Added the `cppcoreguidelines-prefer-member-initializer` clang-tidy check to
enforce this policy.
- Non-functional change: Fixed formatting defects detected by clang-format 13.0.
- Non-functional change: Simplified flow graph disconnection.
- Updated GSL implementation to v0.40.0. See
https://github.com/gsl-lite/gsl-lite/releases/tag/v0.40.0
- CI - `cpplint` job on GitHub: Added the `build/include_what_you_use` filter
@ -95,6 +98,12 @@ All notable changes to GNSS-SDR will be documented in this file.
- Added a new output parameter `Flag_PLL_180_deg_phase_locked` in the monitor
output that indicates if the PLL got locked at 180 degrees, so the symbol sign
is reversed.
- Fix bug in the satellite selection algorithm for configurations with a large
number of channels. The maximum number of channels per signal is now limited
to the number of available satellites per system minus one. The number of
channels performing concurrent acquisition, `Channels.in_acquisition`, cannot
be larger than the total number of channels. The program will stop if those
requirements are not met in the configuration file.
See the definitions of concepts and metrics at
https://gnss-sdr.org/design-forces/

2
src/algorithms/acquisition/adapters/gps_l1_ca_pcps_opencl_acquisition.cc
View File

@ -224,7 +224,7 @@ void GpsL1CaPcpsOpenClAcquisition::reset()
}
float GpsL1CaPcpsOpenClAcquisition::calculate_threshold(float pfa)
float GpsL1CaPcpsOpenClAcquisition::calculate_threshold(float pfa) const
{
// Calculate the threshold
unsigned int frequency_bins = 0;

2
src/algorithms/acquisition/adapters/gps_l1_ca_pcps_opencl_acquisition.h
View File

@ -146,7 +146,7 @@ public:
}
private:
float calculate_threshold(float pfa);
float calculate_threshold(float pfa) const;
const ConfigurationInterface* configuration_;
pcps_opencl_acquisition_cc_sptr acquisition_cc_;
gr::blocks::stream_to_vector::sptr stream_to_vector_;

4
src/algorithms/channel/adapters/channel.cc
View File

@ -89,8 +89,8 @@ Channel::Channel(const ConfigurationInterface* configuration,
acq_->set_threshold(threshold);
acq_->init();
repeat_ = configuration->property("Acquisition_" + signal_str + std::to_string(channel_) + ".repeat_satellite", false);
repeat_ = configuration->property("Acquisition_" + signal_str + ".repeat_satellite", false);
repeat_ = configuration->property("Acquisition_" + signal_str + std::to_string(channel_) + ".repeat_satellite", repeat_);
DLOG(INFO) << "Channel " << channel_ << " satellite repeat = " << repeat_;
channel_fsm_->set_acquisition(acq_);

66
src/algorithms/libs/opencl/fft_execute.cc
View File

@ -28,13 +28,16 @@ allocateTemporaryBufferInterleaved(cl_fft_plan *plan, cl_uint batchSize)
size_t tmpLength = plan->n.x * plan->n.y * plan->n.z * batchSize * 2 * sizeof(cl_float);
if (plan->tempmemobj)
clReleaseMemObject(plan->tempmemobj);
{
clReleaseMemObject(plan->tempmemobj);
}
plan->tempmemobj = clCreateBuffer(plan->context, CL_MEM_READ_WRITE, tmpLength, nullptr, &err);
}
return err;
}
static cl_int
allocateTemporaryBufferPlannar(cl_fft_plan *plan, cl_uint batchSize)
{
@ -46,10 +49,14 @@ allocateTemporaryBufferPlannar(cl_fft_plan *plan, cl_uint batchSize)
size_t tmpLength = plan->n.x * plan->n.y * plan->n.z * batchSize * sizeof(cl_float);
if (plan->tempmemobj_real)
clReleaseMemObject(plan->tempmemobj_real);
{
clReleaseMemObject(plan->tempmemobj_real);
}
if (plan->tempmemobj_imag)
clReleaseMemObject(plan->tempmemobj_imag);
{
clReleaseMemObject(plan->tempmemobj_imag);
}
plan->tempmemobj_real = clCreateBuffer(plan->context, CL_MEM_READ_WRITE, tmpLength, nullptr, &err);
plan->tempmemobj_imag = clCreateBuffer(plan->context, CL_MEM_READ_WRITE, tmpLength, nullptr, &terr);
@ -58,6 +65,7 @@ allocateTemporaryBufferPlannar(cl_fft_plan *plan, cl_uint batchSize)
return err;
}
void getKernelWorkDimensions(cl_fft_plan *plan, cl_fft_kernel_info *kernelInfo, cl_int *batchSize, size_t *gWorkItems, size_t *lWorkItems)
{
*lWorkItems = kernelInfo->num_workitems_per_workgroup;
@ -83,6 +91,7 @@ void getKernelWorkDimensions(cl_fft_plan *plan, cl_fft_kernel_info *kernelInfo,
*gWorkItems = numWorkGroups * *lWorkItems;
}
cl_int
clFFT_ExecuteInterleaved(cl_command_queue queue, clFFT_Plan Plan, cl_int batchSize, clFFT_Direction dir,
cl_mem data_in, cl_mem data_out,
@ -91,7 +100,9 @@ clFFT_ExecuteInterleaved(cl_command_queue queue, clFFT_Plan Plan, cl_int batchSi
int s;
auto *plan = (cl_fft_plan *)Plan;
if (plan->format != clFFT_InterleavedComplexFormat)
return CL_INVALID_VALUE;
{
return CL_INVALID_VALUE;
}
cl_int err;
size_t gWorkItems;
@ -101,7 +112,9 @@ clFFT_ExecuteInterleaved(cl_command_queue queue, clFFT_Plan Plan, cl_int batchSi
cl_int isInPlace = data_in == data_out ? 1 : 0;
if ((err = allocateTemporaryBufferInterleaved(plan, batchSize)) != CL_SUCCESS)
return err;
{
return err;
}
cl_mem memObj[3];
memObj[0] = data_in;
@ -146,7 +159,9 @@ clFFT_ExecuteInterleaved(cl_command_queue queue, clFFT_Plan Plan, cl_int batchSi
err |= clEnqueueNDRangeKernel(queue, kernelInfo->kernel, 1, nullptr, &gWorkItems, &lWorkItems, 0, nullptr, nullptr);
if (err)
return err;
{
return err;
}
currRead = (currWrite == 1) ? 1 : 2;
currWrite = (currWrite == 1) ? 2 : 1;
@ -169,7 +184,9 @@ clFFT_ExecuteInterleaved(cl_command_queue queue, clFFT_Plan Plan, cl_int batchSi
err |= clEnqueueNDRangeKernel(queue, kernelInfo->kernel, 1, nullptr, &gWorkItems, &lWorkItems, 0, nullptr, nullptr);
if (err)
return err;
{
return err;
}
currRead = 1;
currWrite = 1;
@ -181,6 +198,7 @@ clFFT_ExecuteInterleaved(cl_command_queue queue, clFFT_Plan Plan, cl_int batchSi
return err;
}
cl_int
clFFT_ExecutePlannar(cl_command_queue queue, clFFT_Plan Plan, cl_int batchSize, clFFT_Direction dir,
cl_mem data_in_real, cl_mem data_in_imag, cl_mem data_out_real, cl_mem data_out_imag,
@ -190,7 +208,9 @@ clFFT_ExecutePlannar(cl_command_queue queue, clFFT_Plan Plan, cl_int batchSize,
auto *plan = (cl_fft_plan *)Plan;
if (plan->format != clFFT_SplitComplexFormat)
return CL_INVALID_VALUE;
{
return CL_INVALID_VALUE;
}
cl_int err;
size_t gWorkItems;
@ -200,7 +220,9 @@ clFFT_ExecutePlannar(cl_command_queue queue, clFFT_Plan Plan, cl_int batchSize,
cl_int isInPlace = ((data_in_real == data_out_real) && (data_in_imag == data_out_imag)) ? 1 : 0;
if ((err = allocateTemporaryBufferPlannar(plan, batchSize)) != CL_SUCCESS)
return err;
{
return err;
}
cl_mem memObj_real[3];
cl_mem memObj_imag[3];
@ -252,7 +274,9 @@ clFFT_ExecutePlannar(cl_command_queue queue, clFFT_Plan Plan, cl_int batchSize,
err |= clEnqueueNDRangeKernel(queue, kernelInfo->kernel, 1, nullptr, &gWorkItems, &lWorkItems, 0, nullptr, nullptr);
if (err)
return err;
{
return err;
}
currRead = (currWrite == 1) ? 1 : 2;
currWrite = (currWrite == 1) ? 2 : 1;
@ -276,7 +300,9 @@ clFFT_ExecutePlannar(cl_command_queue queue, clFFT_Plan Plan, cl_int batchSize,
err |= clEnqueueNDRangeKernel(queue, kernelInfo->kernel, 1, nullptr, &gWorkItems, &lWorkItems, 0, nullptr, nullptr);
if (err)
return err;
{
return err;
}
currRead = 1;
currWrite = 1;
@ -288,6 +314,7 @@ clFFT_ExecutePlannar(cl_command_queue queue, clFFT_Plan Plan, cl_int batchSize,
return err;
}
cl_int
clFFT_1DTwistInterleaved(clFFT_Plan Plan, cl_command_queue queue, cl_mem array,
unsigned numRows, unsigned numCols, unsigned startRow, unsigned rowsToProcess, clFFT_Direction dir)
@ -304,12 +331,16 @@ clFFT_1DTwistInterleaved(clFFT_Plan Plan, cl_command_queue queue, cl_mem array,
cl_device_id device_id;
err = clGetCommandQueueInfo(queue, CL_QUEUE_DEVICE, sizeof(cl_device_id), &device_id, nullptr);
if (err)
return err;
{
return err;
}
size_t gSize;
err = clGetKernelWorkGroupInfo(plan->twist_kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &gSize, nullptr);
if (err)
return err;
{
return err;
}
gSize = min(128, gSize);
size_t numGlobalThreads[1] = {max(numCols / gSize, 1) * gSize};
@ -327,6 +358,7 @@ clFFT_1DTwistInterleaved(clFFT_Plan Plan, cl_command_queue queue, cl_mem array,
return err;
}
cl_int
clFFT_1DTwistPlannar(clFFT_Plan Plan, cl_command_queue queue, cl_mem array_real, cl_mem array_imag,
unsigned numRows, unsigned numCols, unsigned startRow, unsigned rowsToProcess, clFFT_Direction dir)
@ -343,12 +375,16 @@ clFFT_1DTwistPlannar(clFFT_Plan Plan, cl_command_queue queue, cl_mem array_real,
cl_device_id device_id;
err = clGetCommandQueueInfo(queue, CL_QUEUE_DEVICE, sizeof(cl_device_id), &device_id, nullptr);
if (err)
return err;
{
return err;
}
size_t gSize;
err = clGetKernelWorkGroupInfo(plan->twist_kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &gSize, nullptr);
if (err)
return err;
{
return err;
}
gSize = min(128, gSize);
size_t numGlobalThreads[1] = {max(numCols / gSize, 1) * gSize};

386
src/algorithms/libs/opencl/fft_kernelstring.cc
View File

@ -30,7 +30,7 @@ num2str(int num)
{
char temp[200];
snprintf(temp, sizeof(temp), "%d", num);
return string(temp);
return {temp};
}
// For any n, this function decomposes n into factors for loacal memory tranpose
@ -155,15 +155,21 @@ getRadixArray(unsigned int n, unsigned int *radixArray, unsigned int *numRadices
}
}
static void
insertHeader(string &kernelString, string &kernelName, clFFT_DataFormat dataFormat)
{
if (dataFormat == clFFT_SplitComplexFormat)
kernelString += string("__kernel void ") + kernelName + string("(__global float *in_real, __global float *in_imag, __global float *out_real, __global float *out_imag, int dir, int S)\n");
{
kernelString += string("__kernel void ") + kernelName + string("(__global float *in_real, __global float *in_imag, __global float *out_real, __global float *out_imag, int dir, int S)\n");
}
else
kernelString += string("__kernel void ") + kernelName + string("(__global float2 *in, __global float2 *out, int dir, int S)\n");
{
kernelString += string("__kernel void ") + kernelName + string("(__global float2 *in, __global float2 *out, int dir, int S)\n");
}
}
static void
insertVariables(string &kStream, int maxRadix)
{
@ -177,11 +183,14 @@ insertVariables(string &kStream, int maxRadix)
kStream += string(" int groupId = get_group_id( 0 );\n");
}
static void
formattedLoad(string &kernelString, int aIndex, int gIndex, clFFT_DataFormat dataFormat)
{
if (dataFormat == clFFT_InterleavedComplexFormat)
kernelString += string(" a[") + num2str(aIndex) + string("] = in[") + num2str(gIndex) + string("];\n");
{
kernelString += string(" a[") + num2str(aIndex) + string("] = in[") + num2str(gIndex) + string("];\n");
}
else
{
kernelString += string(" a[") + num2str(aIndex) + string("].x = in_real[") + num2str(gIndex) + string("];\n");
@ -189,11 +198,14 @@ formattedLoad(string &kernelString, int aIndex, int gIndex, clFFT_DataFormat dat
}
}
static void
formattedStore(string &kernelString, int aIndex, int gIndex, clFFT_DataFormat dataFormat)
{
if (dataFormat == clFFT_InterleavedComplexFormat)
kernelString += string(" out[") + num2str(gIndex) + string("] = a[") + num2str(aIndex) + string("];\n");
{
kernelString += string(" out[") + num2str(gIndex) + string("] = a[") + num2str(aIndex) + string("];\n");
}
else
{
kernelString += string(" out_real[") + num2str(gIndex) + string("] = a[") + num2str(aIndex) + string("].x;\n");
@ -201,6 +213,7 @@ formattedStore(string &kernelString, int aIndex, int gIndex, clFFT_DataFormat da
}
}
static int
insertGlobalLoadsAndTranspose(string &kernelString, int N, int numWorkItemsPerXForm, int numXFormsPerWG, int R0, int mem_coalesce_width, clFFT_DataFormat dataFormat)
{
@ -211,7 +224,9 @@ insertGlobalLoadsAndTranspose(string &kernelString, int N, int numWorkItemsPerXF
int lMemSize = 0;
if (numXFormsPerWG > 1)
kernelString += string(" s = S & ") + num2str(numXFormsPerWG - 1) + string(";\n");
{
kernelString += string(" s = S & ") + num2str(numXFormsPerWG - 1) + string(";\n");
}
if (numWorkItemsPerXForm >= mem_coalesce_width)
{
@ -234,7 +249,9 @@ insertGlobalLoadsAndTranspose(string &kernelString, int N, int numWorkItemsPerXF
kernelString += string(" out_imag += offset;\n");
}
for (i = 0; i < R0; i++)
formattedLoad(kernelString, i, i * numWorkItemsPerXForm, dataFormat);
{
formattedLoad(kernelString, i, i * numWorkItemsPerXForm, dataFormat);
}
kernelString += string(" }\n");
}
else
@ -255,7 +272,9 @@ insertGlobalLoadsAndTranspose(string &kernelString, int N, int numWorkItemsPerXF
kernelString += string(" out_imag += offset;\n");
}
for (i = 0; i < R0; i++)
formattedLoad(kernelString, i, i * numWorkItemsPerXForm, dataFormat);
{
formattedLoad(kernelString, i, i * numWorkItemsPerXForm, dataFormat);
}
}
}
else if (N >= mem_coalesce_width)
@ -286,17 +305,23 @@ insertGlobalLoadsAndTranspose(string &kernelString, int N, int numWorkItemsPerXF
{
kernelString += string(" if( jj < s ) {\n");
for (j = 0; j < numInnerIter; j++)
formattedLoad(kernelString, i * numInnerIter + j, j * mem_coalesce_width + i * (groupSize / mem_coalesce_width) * N, dataFormat);
{
formattedLoad(kernelString, i * numInnerIter + j, j * mem_coalesce_width + i * (groupSize / mem_coalesce_width) * N, dataFormat);
}
kernelString += string(" }\n");
if (i != numOuterIter - 1)
kernelString += string(" jj += ") + num2str(groupSize / mem_coalesce_width) + string(";\n");
{
kernelString += string(" jj += ") + num2str(groupSize / mem_coalesce_width) + string(";\n");
}
}
kernelString += string("}\n ");
kernelString += string("else {\n");
for (i = 0; i < numOuterIter; i++)
{
for (j = 0; j < numInnerIter; j++)
formattedLoad(kernelString, i * numInnerIter + j, j * mem_coalesce_width + i * (groupSize / mem_coalesce_width) * N, dataFormat);
{
formattedLoad(kernelString, i * numInnerIter + j, j * mem_coalesce_width + i * (groupSize / mem_coalesce_width) * N, dataFormat);
}
}
kernelString += string("}\n");
@ -315,7 +340,9 @@ insertGlobalLoadsAndTranspose(string &kernelString, int N, int numWorkItemsPerXF
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
for (i = 0; i < R0; i++)
kernelString += string(" a[") + num2str(i) + string("].x = lMemLoad[") + num2str(i * numWorkItemsPerXForm) + string("];\n");
{
kernelString += string(" a[") + num2str(i) + string("].x = lMemLoad[") + num2str(i * numWorkItemsPerXForm) + string("];\n");
}
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
for (i = 0; i < numOuterIter; i++)
@ -329,7 +356,9 @@ insertGlobalLoadsAndTranspose(string &kernelString, int N, int numWorkItemsPerXF
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
for (i = 0; i < R0; i++)
kernelString += string(" a[") + num2str(i) + string("].y = lMemLoad[") + num2str(i * numWorkItemsPerXForm) + string("];\n");
{
kernelString += string(" a[") + num2str(i) + string("].y = lMemLoad[") + num2str(i * numWorkItemsPerXForm) + string("];\n");
}
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
lMemSize = (N + numWorkItemsPerXForm) * numXFormsPerWG;
@ -360,7 +389,9 @@ insertGlobalLoadsAndTranspose(string &kernelString, int N, int numWorkItemsPerXF
kernelString += string(" if(jj < s )\n");
formattedLoad(kernelString, i, i * groupSize, dataFormat);
if (i != R0 - 1)
kernelString += string(" jj += ") + num2str(groupSize / N) + string(";\n");
{
kernelString += string(" jj += ") + num2str(groupSize / N) + string(";\n");
}
}
kernelString += string("}\n");
kernelString += string("else {\n");
@ -385,19 +416,27 @@ insertGlobalLoadsAndTranspose(string &kernelString, int N, int numWorkItemsPerXF
for (i = 0; i < R0; i++)
kernelString += string(" lMemStore[") + num2str(i * (groupSize / N) * (N + numWorkItemsPerXForm)) + string("] = a[") + num2str(i) + string("].x;\n");
{
kernelString += string(" lMemStore[") + num2str(i * (groupSize / N) * (N + numWorkItemsPerXForm)) + string("] = a[") + num2str(i) + string("].x;\n");
}
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
for (i = 0; i < R0; i++)
kernelString += string(" a[") + num2str(i) + string("].x = lMemLoad[") + num2str(i * numWorkItemsPerXForm) + string("];\n");
{
kernelString += string(" a[") + num2str(i) + string("].x = lMemLoad[") + num2str(i * numWorkItemsPerXForm) + string("];\n");
}
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
for (i = 0; i < R0; i++)
kernelString += string(" lMemStore[") + num2str(i * (groupSize / N) * (N + numWorkItemsPerXForm)) + string("] = a[") + num2str(i) + string("].y;\n");
{
kernelString += string(" lMemStore[") + num2str(i * (groupSize / N) * (N + numWorkItemsPerXForm)) + string("] = a[") + num2str(i) + string("].y;\n");
}
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
for (i = 0; i < R0; i++)
kernelString += string(" a[") + num2str(i) + string("].y = lMemLoad[") + num2str(i * numWorkItemsPerXForm) + string("];\n");
{
kernelString += string(" a[") + num2str(i) + string("].y = lMemLoad[") + num2str(i * numWorkItemsPerXForm) + string("];\n");
}
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
lMemSize = (N + numWorkItemsPerXForm) * numXFormsPerWG;
@ -406,6 +445,7 @@ insertGlobalLoadsAndTranspose(string &kernelString, int N, int numWorkItemsPerXF
return lMemSize;
}
static int
insertGlobalStoresAndTranspose(string &kernelString, int N, int maxRadix, int Nr, int numWorkItemsPerXForm, int numXFormsPerWG, int mem_coalesce_width, clFFT_DataFormat dataFormat)
{
@ -433,7 +473,9 @@ insertGlobalStoresAndTranspose(string &kernelString, int N, int maxRadix, int Nr
formattedStore(kernelString, ind, i * numWorkItemsPerXForm, dataFormat);
}
if (numXFormsPerWG > 1)
kernelString += string(" }\n");
{
kernelString += string(" }\n");
}
}
else if (N >= mem_coalesce_width)
{
@ -455,8 +497,12 @@ insertGlobalStoresAndTranspose(string &kernelString, int N, int maxRadix, int Nr
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
for (i = 0; i < numOuterIter; i++)
for (j = 0; j < numInnerIter; j++)
kernelString += string(" a[") + num2str(i * numInnerIter + j) + string("].x = lMemStore[") + num2str(j * mem_coalesce_width + i * (groupSize / mem_coalesce_width) * (N + numWorkItemsPerXForm)) + string("];\n");
{
for (j = 0; j < numInnerIter; j++)
{
kernelString += string(" a[") + num2str(i * numInnerIter + j) + string("].x = lMemStore[") + num2str(j * mem_coalesce_width + i * (groupSize / mem_coalesce_width) * (N + numWorkItemsPerXForm)) + string("];\n");
}
}
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
for (i = 0; i < maxRadix; i++)
@ -469,8 +515,12 @@ insertGlobalStoresAndTranspose(string &kernelString, int N, int maxRadix, int Nr
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
for (i = 0; i < numOuterIter; i++)
for (j = 0; j < numInnerIter; j++)
kernelString += string(" a[") + num2str(i * numInnerIter + j) + string("].y = lMemStore[") + num2str(j * mem_coalesce_width + i * (groupSize / mem_coalesce_width) * (N + numWorkItemsPerXForm)) + string("];\n");
{
for (j = 0; j < numInnerIter; j++)
{
kernelString += string(" a[") + num2str(i * numInnerIter + j) + string("].y = lMemStore[") + num2str(j * mem_coalesce_width + i * (groupSize / mem_coalesce_width) * (N + numWorkItemsPerXForm)) + string("];\n");
}
}
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
kernelString += string("if((groupId == get_num_groups(0)-1) && s) {\n");
@ -478,17 +528,23 @@ insertGlobalStoresAndTranspose(string &kernelString, int N, int maxRadix, int Nr
{
kernelString += string(" if( jj < s ) {\n");
for (j = 0; j < numInnerIter; j++)
formattedStore(kernelString, i * numInnerIter + j, j * mem_coalesce_width + i * (groupSize / mem_coalesce_width) * N, dataFormat);
{
formattedStore(kernelString, i * numInnerIter + j, j * mem_coalesce_width + i * (groupSize / mem_coalesce_width) * N, dataFormat);
}
kernelString += string(" }\n");
if (i != numOuterIter - 1)
kernelString += string(" jj += ") + num2str(groupSize / mem_coalesce_width) + string(";\n");
{
kernelString += string(" jj += ") + num2str(groupSize / mem_coalesce_width) + string(";\n");
}
}
kernelString += string("}\n");
kernelString += string("else {\n");
for (i = 0; i < numOuterIter; i++)
{
for (j = 0; j < numInnerIter; j++)
formattedStore(kernelString, i * numInnerIter + j, j * mem_coalesce_width + i * (groupSize / mem_coalesce_width) * N, dataFormat);
{
formattedStore(kernelString, i * numInnerIter + j, j * mem_coalesce_width + i * (groupSize / mem_coalesce_width) * N, dataFormat);
}
}
kernelString += string("}\n");
@ -512,7 +568,9 @@ insertGlobalStoresAndTranspose(string &kernelString, int N, int maxRadix, int Nr
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
for (i = 0; i < maxRadix; i++)
kernelString += string(" a[") + num2str(i) + string("].x = lMemStore[") + num2str(i * (groupSize / N) * (N + numWorkItemsPerXForm)) + string("];\n");
{
kernelString += string(" a[") + num2str(i) + string("].x = lMemStore[") + num2str(i * (groupSize / N) * (N + numWorkItemsPerXForm)) + string("];\n");
}
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
for (i = 0; i < maxRadix; i++)
@ -525,7 +583,9 @@ insertGlobalStoresAndTranspose(string &kernelString, int N, int maxRadix, int Nr
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
for (i = 0; i < maxRadix; i++)
kernelString += string(" a[") + num2str(i) + string("].y = lMemStore[") + num2str(i * (groupSize / N) * (N + numWorkItemsPerXForm)) + string("];\n");
{
kernelString += string(" a[") + num2str(i) + string("].y = lMemStore[") + num2str(i * (groupSize / N) * (N + numWorkItemsPerXForm)) + string("];\n");
}
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
kernelString += string("if((groupId == get_num_groups(0)-1) && s) {\n");
@ -535,7 +595,9 @@ insertGlobalStoresAndTranspose(string &kernelString, int N, int maxRadix, int Nr
formattedStore(kernelString, i, i * groupSize, dataFormat);
kernelString += string(" }\n");
if (i != maxRadix - 1)
kernelString += string(" jj +=") + num2str(groupSize / N) + string(";\n");
{
kernelString += string(" jj +=") + num2str(groupSize / N) + string(";\n");
}
}
kernelString += string("}\n");
kernelString += string("else {\n");
@ -551,6 +613,7 @@ insertGlobalStoresAndTranspose(string &kernelString, int N, int maxRadix, int Nr
return lMemSize;
}
static void
insertfftKernel(string &kernelString, int Nr, int numIter)
{
@ -561,6 +624,7 @@ insertfftKernel(string &kernelString, int Nr, int numIter)
}
}
static void
insertTwiddleKernel(string &kernelString, int Nr, int numIter, int Nprev, int len, int numWorkItemsPerXForm)
{
@ -573,16 +637,24 @@ insertTwiddleKernel(string &kernelString, int Nr, int numIter, int Nprev, int le
if (z == 0)
{
if (Nprev > 1)
kernelString += string(" angf = (float) (ii >> ") + num2str(logNPrev) + string(");\n");
{
kernelString += string(" angf = (float) (ii >> ") + num2str(logNPrev) + string(");\n");
}
else
kernelString += string(" angf = (float) ii;\n");
{
kernelString += string(" angf = (float) ii;\n");
}
}
else
{
if (Nprev > 1)
kernelString += string(" angf = (float) ((") + num2str(z * numWorkItemsPerXForm) + string(" + ii) >>") + num2str(logNPrev) + string(");\n");
{
kernelString += string(" angf = (float) ((") + num2str(z * numWorkItemsPerXForm) + string(" + ii) >>") + num2str(logNPrev) + string(");\n");
}
else
kernelString += string(" angf = (float) (") + num2str(z * numWorkItemsPerXForm) + string(" + ii);\n");
{
kernelString += string(" angf = (float) (") + num2str(z * numWorkItemsPerXForm) + string(" + ii);\n");
}
}
for (k = 1; k < Nr; k++)
@ -596,30 +668,41 @@ insertTwiddleKernel(string &kernelString, int Nr, int numIter, int Nprev, int le
}
}
static int
getPadding(int numWorkItemsPerXForm, int Nprev, int numWorkItemsReq, int numXFormsPerWG, int Nr, int numBanks, int *offset, int *midPad)
{
if ((numWorkItemsPerXForm <= Nprev) || (Nprev >= numBanks))
*offset = 0;
{
*offset = 0;
}
else
{
int numRowsReq = ((numWorkItemsPerXForm < numBanks) ? numWorkItemsPerXForm : numBanks) / Nprev;
int numColsReq = 1;
if (numRowsReq > Nr)
numColsReq = numRowsReq / Nr;
{
numColsReq = numRowsReq / Nr;
}
numColsReq = Nprev * numColsReq;
*offset = numColsReq;
}
if (numWorkItemsPerXForm >= numBanks || numXFormsPerWG == 1)
*midPad = 0;
{
*midPad = 0;
}
else
{
int bankNum = ((numWorkItemsReq + *offset) * Nr) & (numBanks - 1);
if (bankNum >= numWorkItemsPerXForm)
*midPad = 0;
{
*midPad = 0;
}
else
*midPad = numWorkItemsPerXForm - bankNum;
{
*midPad = numWorkItemsPerXForm - bankNum;
}
}
int lMemSize = (numWorkItemsReq + *offset) * Nr * numXFormsPerWG + *midPad * (numXFormsPerWG - 1);
@ -644,6 +727,7 @@ insertLocalStores(string &kernelString, int numIter, int Nr, int numWorkItemsPer
kernelString += string(" barrier(CLK_LOCAL_MEM_FENCE);\n");
}
static void
insertLocalLoads(string &kernelString, int n, int Nr, int Nrn, int Nprev, int Ncurr, int numWorkItemsPerXForm, int numWorkItemsReq, int offset, string &comp)
{
@ -676,6 +760,7 @@ insertLocalLoads(string &kernelString, int n, int Nr, int Nrn, int Nprev, int Nc
kernelString += string(" barrier(CLK_LOCAL_MEM_FENCE);\n");
}
static void
insertLocalLoadIndexArithmatic(string &kernelString, int Nprev, int Nr, int numWorkItemsReq, int numWorkItemsPerXForm, int numXFormsPerWG, int offset, int midPad)
{
@ -687,33 +772,52 @@ insertLocalLoadIndexArithmatic(string &kernelString, int Nprev, int Nr, int numW
if (Ncurr < numWorkItemsPerXForm)
{
if (Nprev == 1)
kernelString += string(" j = ii & ") + num2str(Ncurr - 1) + string(";\n");
{
kernelString += string(" j = ii & ") + num2str(Ncurr - 1) + string(";\n");
}
else
kernelString += string(" j = (ii & ") + num2str(Ncurr - 1) + string(") >> ") + num2str(logNprev) + string(";\n");
{
kernelString += string(" j = (ii & ") + num2str(Ncurr - 1) + string(") >> ") + num2str(logNprev) + string(";\n");
}
if (Nprev == 1)
kernelString += string(" i = ii >> ") + num2str(logNcurr) + string(";\n");
{
kernelString += string(" i = ii >> ") + num2str(logNcurr) + string(";\n");
}
else
kernelString += string(" i = mad24(ii >> ") + num2str(logNcurr) + string(", ") + num2str(Nprev) + string(", ii & ") + num2str(Nprev - 1) + string(");\n");
{
kernelString += string(" i = mad24(ii >> ") + num2str(logNcurr) + string(", ") + num2str(Nprev) + string(", ii & ") + num2str(Nprev - 1) + string(");\n");
}
}
else
{
if (Nprev == 1)
kernelString += string(" j = ii;\n");
{
kernelString += string(" j = ii;\n");
}
else
kernelString += string(" j = ii >> ") + num2str(logNprev) + string(";\n");
{
kernelString += string(" j = ii >> ") + num2str(logNprev) + string(";\n");
}
if (Nprev == 1)
kernelString += string(" i = 0;\n");
{
kernelString += string(" i = 0;\n");
}
else
kernelString += string(" i = ii & ") + num2str(Nprev - 1) + string(";\n");
{
kernelString += string(" i = ii & ") + num2str(Nprev - 1) + string(";\n");
}
}
if (numXFormsPerWG > 1)
kernelString += string(" i = mad24(jj, ") + num2str(incr) + string(", i);\n");
{
kernelString += string(" i = mad24(jj, ") + num2str(incr) + string(", i);\n");
}
kernelString += string(" lMemLoad = sMem + mad24(j, ") + num2str(numWorkItemsReq + offset) + string(", i);\n");
}
static void
insertLocalStoreIndexArithmatic(string &kernelString, int numWorkItemsReq, int numXFormsPerWG, int Nr, int offset, int midPad)
{
@ -742,7 +846,9 @@ createLocalMemfftKernelString(cl_fft_plan *plan)
assert(numRadix > 0 && "no radix array supplied\n");
if (n / radixArray[0] > plan->max_work_item_per_workgroup)
getRadixArray(n, radixArray, &numRadix, plan->max_radix);
{
getRadixArray(n, radixArray, &numRadix, plan->max_radix);
}
assert(radixArray[0] <= plan->max_radix && "max radix choosen is greater than allowed\n");
assert(n / radixArray[0] <= plan->max_work_item_per_workgroup && "required work items per xform greater than maximum work items allowed per work group for local mem fft\n");
@ -839,11 +945,14 @@ createLocalMemfftKernelString(cl_fft_plan *plan)
insertHeader(*kernelString, kernelName, dataFormat);
*kernelString += string("{\n");
if ((*kInfo)->lmem_size)
*kernelString += string(" __local float sMem[") + num2str((*kInfo)->lmem_size) + string("];\n");
{
*kernelString += string(" __local float sMem[") + num2str((*kInfo)->lmem_size) + string("];\n");
}
*kernelString += localString;
*kernelString += string("}\n");
}
// For n larger than what can be computed using local memory fft, global transposes
// multiple kernel launces is needed. For these sizes, n can be decomposed using
// much larger base radices i.e. say n = 262144 = 128 x 64 x 32. Thus three kernel
@ -864,7 +973,6 @@ createLocalMemfftKernelString(cl_fft_plan *plan)
// in this example. Users can play with difference base radices and difference
// decompositions of base radices to generates different kernels and see which gives
// best performance. Following function is just fixed to use 128 as base radix
void getGlobalRadixInfo(int n, int *radix, int *R1, int *R2, int *numRadices)
{
int baseRadix = min(n, 128);
@ -878,7 +986,9 @@ void getGlobalRadixInfo(int n, int *radix, int *R1, int *R2, int *numRadices)
}
for (int i = 0; i < numR; i++)
radix[i] = baseRadix;
{
radix[i] = baseRadix;
}
radix[numR] = N;
numR++;
@ -906,6 +1016,7 @@ void getGlobalRadixInfo(int n, int *radix, int *R1, int *R2, int *numRadices)
}
}
static void
createGlobalFFTKernelString(cl_fft_plan *plan, int n, int BS, cl_fft_kernel_dir dir, int vertBS)
{
@ -960,12 +1071,18 @@ createGlobalFFTKernelString(cl_fft_plan *plan, int n, int BS, cl_fft_kernel_dir
int strideI = Rinit;
for (i = 0; i < numPasses; i++)
if (i != passNum)
strideI *= radixArr[i];
{
if (i != passNum)
{
strideI *= radixArr[i];
}
}
int strideO = Rinit;
for (i = 0; i < passNum; i++)
strideO *= radixArr[i];
{
strideO *= radixArr[i];
}
int threadsPerXForm = R2;
batchSize = R2 == 1 ? plan->max_work_item_per_workgroup : batchSize;
@ -986,30 +1103,44 @@ createGlobalFFTKernelString(cl_fft_plan *plan, int n, int BS, cl_fft_kernel_dir
int numBlocksPerXForm = strideI / batchSize;
int numBlocks = numBlocksPerXForm;
if (!vertical)
numBlocks *= BS;
{
numBlocks *= BS;
}
else
numBlocks *= vertBS;
{
numBlocks *= vertBS;
}
kernelName = string("fft") + num2str(kCount);
*kInfo = (cl_fft_kernel_info *)malloc(sizeof(cl_fft_kernel_info));
(*kInfo)->kernel = nullptr;
if (R2 == 1)
(*kInfo)->lmem_size = 0;
{
(*kInfo)->lmem_size = 0;
}
else
{
if (strideO == 1)
(*kInfo)->lmem_size = (radix + 1) * batchSize;
{
(*kInfo)->lmem_size = (radix + 1) * batchSize;
}
else
(*kInfo)->lmem_size = threadsPerBlock * R1;
{
(*kInfo)->lmem_size = threadsPerBlock * R1;
}
}
(*kInfo)->num_workgroups = numBlocks;
(*kInfo)->num_xforms_per_workgroup = 1;
(*kInfo)->num_workitems_per_workgroup = threadsPerBlock;
(*kInfo)->dir = dir;
if ((passNum == (numPasses - 1)) && (numPasses & 1))
(*kInfo)->in_place_possible = 1;
{
(*kInfo)->in_place_possible = 1;
}
else
(*kInfo)->in_place_possible = 0;
{
(*kInfo)->in_place_possible = 0;
}
(*kInfo)->next = nullptr;
(*kInfo)->kernel_name = (char *)malloc(sizeof(char) * (kernelName.size() + 1));
snprintf((*kInfo)->kernel_name, sizeof((*kInfo)->kernel_name), kernelName.c_str());
@ -1026,7 +1157,9 @@ createGlobalFFTKernelString(cl_fft_plan *plan, int n, int BS, cl_fft_kernel_dir
localString += string("j = tid & ") + num2str(strideO - 1) + string(";\n");
int stride = radix * Rinit;
for (i = 0; i < passNum; i++)
stride *= radixArr[i];
{
stride *= radixArr[i];
}
localString += string("indexOut = mad24(i, ") + num2str(stride) + string(", j + ") + string("(xNum << ") + num2str((int)log2(n * BS)) + string("));\n");
localString += string("bNum = groupId;\n");
}
@ -1041,7 +1174,9 @@ createGlobalFFTKernelString(cl_fft_plan *plan, int n, int BS, cl_fft_kernel_dir
localString += string("j = tid & ") + num2str(strideO - 1) + string(";\n");
int stride = radix * Rinit;
for (i = 0; i < passNum; i++)
stride *= radixArr[i];
{
stride *= radixArr[i];
}
localString += string("indexOut = mad24(i, ") + num2str(stride) + string(", j);\n");
localString += string("indexIn += (xNum << ") + num2str(m) + string(");\n");
localString += string("indexOut += (xNum << ") + num2str(m) + string(");\n");
@ -1059,15 +1194,21 @@ createGlobalFFTKernelString(cl_fft_plan *plan, int n, int BS, cl_fft_kernel_dir
localString += string("in_real += indexIn;\n");
localString += string("in_imag += indexIn;\n");
for (j = 0; j < R1; j++)
localString += string("a[") + num2str(j) + string("].x = in_real[") + num2str(j * gInInc * strideI) + string("];\n");
{
localString += string("a[") + num2str(j) + string("].x = in_real[") + num2str(j * gInInc * strideI) + string("];\n");
}
for (j = 0; j < R1; j++)
localString += string("a[") + num2str(j) + string("].y = in_imag[") + num2str(j * gInInc * strideI) + string("];\n");
{
localString += string("a[") + num2str(j) + string("].y = in_imag[") + num2str(j * gInInc * strideI) + string("];\n");
}
}
else
{
localString += string("in += indexIn;\n");
for (j = 0; j < R1; j++)
localString += string("a[") + num2str(j) + string("] = in[") + num2str(j * gInInc * strideI) + string("];\n");
{
localString += string("a[") + num2str(j) + string("] = in[") + num2str(j * gInInc * strideI) + string("];\n");
}
}
localString += string("fftKernel") + num2str(R1) + string("(a, dir);\n");
@ -1088,22 +1229,36 @@ createGlobalFFTKernelString(cl_fft_plan *plan, int n, int BS, cl_fft_kernel_dir
localString += string("lMemStore = sMem + tid;\n");
localString += string("lMemLoad = sMem + indexIn;\n");
for (k = 0; k < R1; k++)
localString += string("lMemStore[") + num2str(k * threadsPerBlock) + string("] = a[") + num2str(k) + string("].x;\n");
{
localString += string("lMemStore[") + num2str(k * threadsPerBlock) + string("] = a[") + num2str(k) + string("].x;\n");
}
localString += string("barrier(CLK_LOCAL_MEM_FENCE);\n");
for (k = 0; k < numIter; k++)
for (t = 0; t < R2; t++)
localString += string("a[") + num2str(k * R2 + t) + string("].x = lMemLoad[") + num2str(t * batchSize + k * threadsPerBlock) + string("];\n");
{
for (t = 0; t < R2; t++)
{
localString += string("a[") + num2str(k * R2 + t) + string("].x = lMemLoad[") + num2str(t * batchSize + k * threadsPerBlock) + string("];\n");
}
}
localString += string("barrier(CLK_LOCAL_MEM_FENCE);\n");
for (k = 0; k < R1; k++)
localString += string("lMemStore[") + num2str(k * threadsPerBlock) + string("] = a[") + num2str(k) + string("].y;\n");
{
localString += string("lMemStore[") + num2str(k * threadsPerBlock) + string("] = a[") + num2str(k) + string("].y;\n");
}
localString += string("barrier(CLK_LOCAL_MEM_FENCE);\n");
for (k = 0; k < numIter; k++)
for (t = 0; t < R2; t++)
localString += string("a[") + num2str(k * R2 + t) + string("].y = lMemLoad[") + num2str(t * batchSize + k * threadsPerBlock) + string("];\n");
{
for (t = 0; t < R2; t++)
{
localString += string("a[") + num2str(k * R2 + t) + string("].y = lMemLoad[") + num2str(t * batchSize + k * threadsPerBlock) + string("];\n");
}
}
localString += string("barrier(CLK_LOCAL_MEM_FENCE);\n");
for (j = 0; j < numIter; j++)
localString += string("fftKernel") + num2str(R2) + string("(a + ") + num2str(j * R2) + string(", dir);\n");
{
localString += string("fftKernel") + num2str(R2) + string("(a + ") + num2str(j * R2) + string(", dir);\n");
}
}
// twiddle
@ -1127,40 +1282,60 @@ createGlobalFFTKernelString(cl_fft_plan *plan, int n, int BS, cl_fft_kernel_dir
localString += string("lMemLoad = sMem + mad24(tid >> ") + num2str((int)log2(radix)) + string(", ") + num2str(radix + 1) + string(", tid & ") + num2str(radix - 1) + string(");\n");
for (i = 0; i < R1 / R2; i++)
for (j = 0; j < R2; j++)
localString += string("lMemStore[ ") + num2str(i + j * R1) + string("] = a[") + num2str(i * R2 + j) + string("].x;\n");
{
for (j = 0; j < R2; j++)
{
localString += string("lMemStore[ ") + num2str(i + j * R1) + string("] = a[") + num2str(i * R2 + j) + string("].x;\n");
}
}
localString += string("barrier(CLK_LOCAL_MEM_FENCE);\n");
if (threadsPerBlock >= radix)
{
for (i = 0; i < R1; i++)
localString += string("a[") + num2str(i) + string("].x = lMemLoad[") + num2str(i * (radix + 1) * (threadsPerBlock / radix)) + string("];\n");
{
localString += string("a[") + num2str(i) + string("].x = lMemLoad[") + num2str(i * (radix + 1) * (threadsPerBlock / radix)) + string("];\n");
}
}
else
{
int innerIter = radix / threadsPerBlock;
int outerIter = R1 / innerIter;
for (i = 0; i < outerIter; i++)
for (j = 0; j < innerIter; j++)
localString += string("a[") + num2str(i * innerIter + j) + string("].x = lMemLoad[") + num2str(j * threadsPerBlock + i * (radix + 1)) + string("];\n");
{
for (j = 0; j < innerIter; j++)
{
localString += string("a[") + num2str(i * innerIter + j) + string("].x = lMemLoad[") + num2str(j * threadsPerBlock + i * (radix + 1)) + string("];\n");
}
}
}
localString += string("barrier(CLK_LOCAL_MEM_FENCE);\n");
for (i = 0; i < R1 / R2; i++)
for (j = 0; j < R2; j++)
localString += string("lMemStore[ ") + num2str(i + j * R1) + string("] = a[") + num2str(i * R2 + j) + string("].y;\n");
{
for (j = 0; j < R2; j++)
{
localString += string("lMemStore[ ") + num2str(i + j * R1) + string("] = a[") + num2str(i * R2 + j) + string("].y;\n");
}
}
localString += string("barrier(CLK_LOCAL_MEM_FENCE);\n");
if (threadsPerBlock >= radix)
{
for (i = 0; i < R1; i++)
localString += string("a[") + num2str(i) + string("].y = lMemLoad[") + num2str(i * (radix + 1) * (threadsPerBlock / radix)) + string("];\n");
{
localString += string("a[") + num2str(i) + string("].y = lMemLoad[") + num2str(i * (radix + 1) * (threadsPerBlock / radix)) + string("];\n");
}
}
else
{
int innerIter = radix / threadsPerBlock;
int outerIter = R1 / innerIter;
for (i = 0; i < outerIter; i++)
for (j = 0; j < innerIter; j++)
localString += string("a[") + num2str(i * innerIter + j) + string("].y = lMemLoad[") + num2str(j * threadsPerBlock + i * (radix + 1)) + string("];\n");
{
for (j = 0; j < innerIter; j++)
{
localString += string("a[") + num2str(i * innerIter + j) + string("].y = lMemLoad[") + num2str(j * threadsPerBlock + i * (radix + 1)) + string("];\n");
}
}
}
localString += string("barrier(CLK_LOCAL_MEM_FENCE);\n");
@ -1170,15 +1345,21 @@ createGlobalFFTKernelString(cl_fft_plan *plan, int n, int BS, cl_fft_kernel_dir
localString += string("out_real += indexOut;\n");
localString += string("out_imag += indexOut;\n");
for (k = 0; k < R1; k++)
localString += string("out_real[") + num2str(k * threadsPerBlock) + string("] = a[") + num2str(k) + string("].x;\n");
{
localString += string("out_real[") + num2str(k * threadsPerBlock) + string("] = a[") + num2str(k) + string("].x;\n");
}
for (k = 0; k < R1; k++)
localString += string("out_imag[") + num2str(k * threadsPerBlock) + string("] = a[") + num2str(k) + string("].y;\n");
{
localString += string("out_imag[") + num2str(k * threadsPerBlock) + string("] = a[") + num2str(k) + string("].y;\n");
}
}
else
{
localString += string("out += indexOut;\n");
for (k = 0; k < R1; k++)
localString += string("out[") + num2str(k * threadsPerBlock) + string("] = a[") + num2str(k) + string("];\n");
{
localString += string("out[") + num2str(k * threadsPerBlock) + string("] = a[") + num2str(k) + string("];\n");
}
}
}
else
@ -1189,22 +1370,30 @@ createGlobalFFTKernelString(cl_fft_plan *plan, int n, int BS, cl_fft_kernel_dir
localString += string("out_real += indexOut;\n");
localString += string("out_imag += indexOut;\n");
for (k = 0; k < R1; k++)
localString += string("out_real[") + num2str(((k % R2) * R1 + (k / R2)) * strideO) + string("] = a[") + num2str(k) + string("].x;\n");
{
localString += string("out_real[") + num2str(((k % R2) * R1 + (k / R2)) * strideO) + string("] = a[") + num2str(k) + string("].x;\n");
}
for (k = 0; k < R1; k++)
localString += string("out_imag[") + num2str(((k % R2) * R1 + (k / R2)) * strideO) + string("] = a[") + num2str(k) + string("].y;\n");
{
localString += string("out_imag[") + num2str(((k % R2) * R1 + (k / R2)) * strideO) + string("] = a[") + num2str(k) + string("].y;\n");
}
}
else
{
localString += string("out += indexOut;\n");
for (k = 0; k < R1; k++)
localString += string("out[") + num2str(((k % R2) * R1 + (k / R2)) * strideO) + string("] = a[") + num2str(k) + string("];\n");
{
localString += string("out[") + num2str(((k % R2) * R1 + (k / R2)) * strideO) + string("] = a[") + num2str(k) + string("];\n");
}
}
}
insertHeader(*kernelString, kernelName, dataFormat);
*kernelString += string("{\n");
if ((*kInfo)->lmem_size)
*kernelString += string(" __local float sMem[") + num2str((*kInfo)->lmem_size) + string("];\n");
{
*kernelString += string(" __local float sMem[") + num2str((*kInfo)->lmem_size) + string("];\n");
}
*kernelString += localString;
*kernelString += string("}\n");
@ -1214,6 +1403,7 @@ createGlobalFFTKernelString(cl_fft_plan *plan, int n, int BS, cl_fft_kernel_dir
}
}
void FFT1D(cl_fft_plan *plan, cl_fft_kernel_dir dir)
{
unsigned int radixArray[10];
@ -1237,21 +1427,29 @@ void FFT1D(cl_fft_plan *plan, cl_fft_kernel_dir dir)
{
getRadixArray(plan->n.x, radixArray, &numRadix, plan->max_radix);
if (plan->n.x / radixArray[0] <= plan->max_work_item_per_workgroup)
createLocalMemfftKernelString(plan);
{
createLocalMemfftKernelString(plan);
}
else
createGlobalFFTKernelString(plan, plan->n.x, 1, cl_fft_kernel_x, 1);
{
createGlobalFFTKernelString(plan, plan->n.x, 1, cl_fft_kernel_x, 1);
}
}
}
break;
case cl_fft_kernel_y:
if (plan->n.y > 1)
createGlobalFFTKernelString(plan, plan->n.y, plan->n.x, cl_fft_kernel_y, 1);
{
createGlobalFFTKernelString(plan, plan->n.y, plan->n.x, cl_fft_kernel_y, 1);
}
break;
case cl_fft_kernel_z:
if (plan->n.z > 1)
createGlobalFFTKernelString(plan, plan->n.z, plan->n.x * plan->n.y, cl_fft_kernel_z, 1);
{
createGlobalFFTKernelString(plan, plan->n.z, plan->n.x * plan->n.y, cl_fft_kernel_z, 1);
}
default:
return;
}

63
src/algorithms/libs/opencl/fft_setup.cc
View File

@ -31,9 +31,13 @@ getBlockConfigAndKernelString(cl_fft_plan *plan)
*plan->kernel_string += baseKernels;
if (plan->format == clFFT_SplitComplexFormat)
*plan->kernel_string += twistKernelPlannar;
{
*plan->kernel_string += twistKernelPlannar;
}
else
*plan->kernel_string += twistKernelInterleaved;
{
*plan->kernel_string += twistKernelInterleaved;
}
switch (plan->dim)
{
@ -72,13 +76,18 @@ deleteKernelInfo(cl_fft_kernel_info *kInfo)
if (kInfo)
{
if (kInfo->kernel_name)
free(kInfo->kernel_name);
{
free(kInfo->kernel_name);
}
if (kInfo->kernel)
clReleaseKernel(kInfo->kernel);
{
clReleaseKernel(kInfo->kernel);
}
free(kInfo);
}
}
static void
destroy_plan(cl_fft_plan *Plan)
{
@ -125,6 +134,7 @@ destroy_plan(cl_fft_plan *Plan)
}
}
static int
createKernelList(cl_fft_plan *plan)
{
@ -136,21 +146,30 @@ createKernelList(cl_fft_plan *plan)
{
kernel_info->kernel = clCreateKernel(program, kernel_info->kernel_name, &err);
if (!kernel_info->kernel || err != CL_SUCCESS)
return err;
{
return err;
}
kernel_info = kernel_info->next;
}
if (plan->format == clFFT_SplitComplexFormat)
plan->twist_kernel = clCreateKernel(program, "clFFT_1DTwistSplit", &err);
{
plan->twist_kernel = clCreateKernel(program, "clFFT_1DTwistSplit", &err);
}
else
plan->twist_kernel = clCreateKernel(program, "clFFT_1DTwistInterleaved", &err);
{
plan->twist_kernel = clCreateKernel(program, "clFFT_1DTwistInterleaved", &err);
}
if (!plan->twist_kernel || err)
return err;
{
return err;
}
return CL_SUCCESS;
}
int getMaxKernelWorkGroupSize(cl_fft_plan *plan, unsigned int *max_wg_size, unsigned int num_devices, cl_device_id *devices)
{
int reg_needed = 0;
@ -166,13 +185,19 @@ int getMaxKernelWorkGroupSize(cl_fft_plan *plan, unsigned int *max_wg_size, unsi
{
err = clGetKernelWorkGroupInfo(kInfo->kernel, devices[i], CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &wg_size, nullptr);
if (err != CL_SUCCESS)
return -1;
{
return -1;
}
if (wg_size < kInfo->num_workitems_per_workgroup)
reg_needed |= 1;
{
reg_needed |= 1;
}
if (*max_wg_size > wg_size)
*max_wg_size = wg_size;
{
*max_wg_size = wg_size;
}
kInfo = kInfo->next;
}
@ -181,6 +206,7 @@ int getMaxKernelWorkGroupSize(cl_fft_plan *plan, unsigned int *max_wg_size, unsi
return reg_needed;
}
#define ERR_MACRO(err) \
{ \
if ((err) != CL_SUCCESS) \
@ -192,6 +218,7 @@ int getMaxKernelWorkGroupSize(cl_fft_plan *plan, unsigned int *max_wg_size, unsi
} \
}
clFFT_Plan
clFFT_CreatePlan(cl_context context, clFFT_Dim3 n, clFFT_Dimension dim, clFFT_DataFormat dataFormat, cl_int *error_code)
{
@ -326,11 +353,14 @@ patch_kernel_source:
}
if (error_code)
*error_code = CL_SUCCESS;
{
*error_code = CL_SUCCESS;
}
return (clFFT_Plan)plan;
}
void clFFT_DestroyPlan(clFFT_Plan plan)
{
auto *Plan = (cl_fft_plan *)plan;
@ -342,15 +372,20 @@ void clFFT_DestroyPlan(clFFT_Plan plan)
}
}
void clFFT_DumpPlan(clFFT_Plan Plan, FILE *file)
{
size_t gDim;
size_t lDim;
FILE *out;
if (!file)
out = stdout;
{
out = stdout;
}
else
out = file;
{
out = file;
}
auto *plan = (cl_fft_plan *)Plan;
cl_fft_kernel_info *kInfo = plan->kernel_info;

2
src/algorithms/libs/volk_gnsssdr_module/volk_gnsssdr/cmake/Toolchains/arm_cortex_a53_hardfp_native.cmake
View File

@ -5,7 +5,7 @@
# SPDX-License-Identifier: BSD-3-Clause
########################################################################
# Toolchain file for building native on a ARM Cortex A72 w/ NEON
# Toolchain file for building native on a ARM Cortex A53 w/ NEON
# Usage: cmake -DCMAKE_TOOLCHAIN_FILE=<this file> <source directory>
########################################################################
set(CMAKE_CXX_COMPILER g++)

22
src/algorithms/signal_source/adapters/ad9361_fpga_signal_source.cc
View File

@ -283,7 +283,7 @@ Ad9361FpgaSignalSource::Ad9361FpgaSignalSource(const ConfigurationInterface *con
enable_ovf_check_buffer_monitor_active_ = false; // check buffer overflow and buffer monitor disabled by default
// some basic checks
if ((rf_port_select_ != "A_BALANCED") and (rf_port_select_ != "B_BALANCED") and (rf_port_select_ != "A_N") and (rf_port_select_ != "B_N") and (rf_port_select_ != "B_P") and (rf_port_select_ != "C_N") and (rf_port_select_ != "C_P") and (rf_port_select_ != "TX_MONITOR1") and (rf_port_select_ != "TX_MONITOR2") and (rf_port_select_ != "TX_MONITOR1_2"))
if ((rf_port_select_ != "A_BALANCED") && (rf_port_select_ != "B_BALANCED") && (rf_port_select_ != "A_N") && (rf_port_select_ != "B_N") && (rf_port_select_ != "B_P") && (rf_port_select_ != "C_N") && (rf_port_select_ != "C_P") && (rf_port_select_ != "TX_MONITOR1") && (rf_port_select_ != "TX_MONITOR2") && (rf_port_select_ != "TX_MONITOR1_2"))
{
std::cout << "Configuration parameter rf_port_select should take one of these values:\n";
std::cout << " A_BALANCED, B_BALANCED, A_N, B_N, B_P, C_N, C_P, TX_MONITOR1, TX_MONITOR2, TX_MONITOR1_2\n";
@ -293,7 +293,7 @@ Ad9361FpgaSignalSource::Ad9361FpgaSignalSource(const ConfigurationInterface *con
LOG(WARNING) << "Invalid configuration value for rf_port_select parameter. Set to rf_port_select=" << default_rf_port_select;
}
if ((gain_mode_rx1_ != "manual") and (gain_mode_rx1_ != "slow_attack") and (gain_mode_rx1_ != "fast_attack") and (gain_mode_rx1_ != "hybrid"))
if ((gain_mode_rx1_ != "manual") && (gain_mode_rx1_ != "slow_attack") && (gain_mode_rx1_ != "fast_attack") && (gain_mode_rx1_ != "hybrid"))
{
std::cout << "Configuration parameter gain_mode_rx1 should take one of these values:\n";
std::cout << " manual, slow_attack, fast_attack, hybrid\n";
@ -303,7 +303,7 @@ Ad9361FpgaSignalSource::Ad9361FpgaSignalSource(const ConfigurationInterface *con
LOG(WARNING) << "Invalid configuration value for gain_mode_rx1 parameter. Set to gain_mode_rx1=" << default_gain_mode;
}
if ((gain_mode_rx2_ != "manual") and (gain_mode_rx2_ != "slow_attack") and (gain_mode_rx2_ != "fast_attack") and (gain_mode_rx2_ != "hybrid"))
if ((gain_mode_rx2_ != "manual") && (gain_mode_rx2_ != "slow_attack") && (gain_mode_rx2_ != "fast_attack") && (gain_mode_rx2_ != "hybrid"))
{
std::cout << "Configuration parameter gain_mode_rx2 should take one of these values:\n";
std::cout << " manual, slow_attack, fast_attack, hybrid\n";
@ -315,7 +315,7 @@ Ad9361FpgaSignalSource::Ad9361FpgaSignalSource(const ConfigurationInterface *con
if (gain_mode_rx1_ == "manual")
{
if (rf_gain_rx1_ > 73.0 or rf_gain_rx1_ < -1.0)
if (rf_gain_rx1_ > 73.0 || rf_gain_rx1_ < -1.0)
{
std::cout << "Configuration parameter rf_gain_rx1 should take values between -1.0 and 73 dB\n";
std::cout << "Error: provided value rf_gain_rx1=" << rf_gain_rx1_ << " is not among valid values\n";
@ -327,7 +327,7 @@ Ad9361FpgaSignalSource::Ad9361FpgaSignalSource(const ConfigurationInterface *con
if (gain_mode_rx2_ == "manual")
{
if (rf_gain_rx2_ > 73.0 or rf_gain_rx2_ < -1.0)
if (rf_gain_rx2_ > 73.0 || rf_gain_rx2_ < -1.0)
{
std::cout << "Configuration parameter rf_gain_rx2 should take values between -1.0 and 73 dB\n";
std::cout << "Error: provided value rf_gain_rx2=" << rf_gain_rx2_ << " is not among valid values\n";
@ -337,7 +337,7 @@ Ad9361FpgaSignalSource::Ad9361FpgaSignalSource(const ConfigurationInterface *con
}
}
if ((filter_source_ != "Off") and (filter_source_ != "Auto") and (filter_source_ != "File") and (filter_source_ != "Design"))
if ((filter_source_ != "Off") && (filter_source_ != "Auto") && (filter_source_ != "File") && (filter_source_ != "Design"))
{
std::cout << "Configuration parameter filter_source should take one of these values:\n";
std::cout << " Off: Disable filter\n";
@ -350,7 +350,7 @@ Ad9361FpgaSignalSource::Ad9361FpgaSignalSource(const ConfigurationInterface *con
LOG(WARNING) << "Invalid configuration value for filter_source parameter. Set to filter_source=Off";
}
if (bandwidth_ < 200000 or bandwidth_ > 56000000)
if (bandwidth_ < 200000 || bandwidth_ > 56000000)
{
std::cout << "Configuration parameter bandwidth should take values between 200000 and 56000000 Hz\n";
std::cout << "Error: provided value bandwidth=" << bandwidth_ << " is not among valid values\n";
@ -389,7 +389,7 @@ Ad9361FpgaSignalSource::Ad9361FpgaSignalSource(const ConfigurationInterface *con
// LOCAL OSCILLATOR DDS GENERATOR FOR DUAL FREQUENCY OPERATION
if (enable_dds_lo_ == true)
{
if (tx_bandwidth_ < static_cast<uint64_t>(std::floor(static_cast<float>(freq_dds_tx_hz_) * 1.1)) or (tx_bandwidth_ < 200000) or (tx_bandwidth_ > 1000000))
if (tx_bandwidth_ < static_cast<uint64_t>(std::floor(static_cast<float>(freq_dds_tx_hz_) * 1.1)) || (tx_bandwidth_ < 200000) || (tx_bandwidth_ > 1000000))
{
std::cout << "Configuration parameter tx_bandwidth value should be between " << std::max(static_cast<float>(freq_dds_tx_hz_) * 1.1, 200000.0) << " and 1000000 Hz\n";
std::cout << "Error: provided value tx_bandwidth=" << tx_bandwidth_ << " is not among valid values\n";
@ -397,7 +397,7 @@ Ad9361FpgaSignalSource::Ad9361FpgaSignalSource(const ConfigurationInterface *con
tx_bandwidth_ = 500000;
LOG(WARNING) << "Invalid configuration value for tx_bandwidth parameter. Set to tx_bandwidth=500000";
}
if (tx_attenuation_db_ > 0.0 or tx_attenuation_db_ < -89.75)
if (tx_attenuation_db_ > 0.0 || tx_attenuation_db_ < -89.75)
{
std::cout << "Configuration parameter tx_attenuation_db should take values between 0.0 and -89.95 in 0.25 dB steps\n";
std::cout << "Error: provided value tx_attenuation_db=" << tx_attenuation_db_ << " is not among valid values\n";
@ -908,11 +908,11 @@ void Ad9361FpgaSignalSource::disconnect(gr::top_block_sptr top_block)
gr::basic_block_sptr Ad9361FpgaSignalSource::get_left_block()
{
LOG(WARNING) << "Trying to get signal source left block.";
return gr::basic_block_sptr();
return {};
}
gr::basic_block_sptr Ad9361FpgaSignalSource::get_right_block()
{
return gr::basic_block_sptr();
return {};
}

8
src/algorithms/signal_source/adapters/file_source_base.cc
View File

@ -65,6 +65,14 @@ FileSourceBase::FileSourceBase(ConfigurationInterface const* configuration, std:
{
filename_ = FLAGS_s;
}
if (sampling_frequency_ == 0)
{
std::cerr << "Warning: parameter " << role << ".sampling_frequency is not set, this could lead to wrong results.\n"
<< " Please set the " << role << ".sampling_frequency parameter in your configuration file.\n"
<< " If not set, " << role << ".sampling_frequency=" << configuration->property("GNSS-SDR.internal_fs_sps"s, int64_t(0))
<< " will be assumed.\n";
sampling_frequency_ = configuration->property("GNSS-SDR.internal_fs_sps"s, int64_t(0));
}
}

28
src/algorithms/signal_source/adapters/fmcomms2_signal_source.cc
View File

@ -91,7 +91,7 @@ Fmcomms2SignalSource::Fmcomms2SignalSource(const ConfigurationInterface *configu
item_size_ = sizeof(gr_complex);
// some basic checks
if ((rf_port_select_ != "A_BALANCED") and (rf_port_select_ != "B_BALANCED") and (rf_port_select_ != "A_N") and (rf_port_select_ != "B_N") and (rf_port_select_ != "B_P") and (rf_port_select_ != "C_N") and (rf_port_select_ != "C_P") and (rf_port_select_ != "TX_MONITOR1") and (rf_port_select_ != "TX_MONITOR2") and (rf_port_select_ != "TX_MONITOR1_2"))
if ((rf_port_select_ != "A_BALANCED") && (rf_port_select_ != "B_BALANCED") && (rf_port_select_ != "A_N") && (rf_port_select_ != "B_N") && (rf_port_select_ != "B_P") && (rf_port_select_ != "C_N") && (rf_port_select_ != "C_P") && (rf_port_select_ != "TX_MONITOR1") && (rf_port_select_ != "TX_MONITOR2") && (rf_port_select_ != "TX_MONITOR1_2"))
{
std::cout << "Configuration parameter rf_port_select should take one of these values:\n";
std::cout << " A_BALANCED, B_BALANCED, A_N, B_N, B_P, C_N, C_P, TX_MONITOR1, TX_MONITOR2, TX_MONITOR1_2\n";
@ -101,7 +101,7 @@ Fmcomms2SignalSource::Fmcomms2SignalSource(const ConfigurationInterface *configu
LOG(WARNING) << "Invalid configuration value for rf_port_select parameter. Set to rf_port_select=A_BALANCED";
}
if ((gain_mode_rx1_ != "manual") and (gain_mode_rx1_ != "slow_attack") and (gain_mode_rx1_ != "fast_attack") and (gain_mode_rx1_ != "hybrid"))
if ((gain_mode_rx1_ != "manual") && (gain_mode_rx1_ != "slow_attack") && (gain_mode_rx1_ != "fast_attack") && (gain_mode_rx1_ != "hybrid"))
{
std::cout << "Configuration parameter gain_mode_rx1 should take one of these values:\n";
std::cout << " manual, slow_attack, fast_attack, hybrid\n";
@ -111,7 +111,7 @@ Fmcomms2SignalSource::Fmcomms2SignalSource(const ConfigurationInterface *configu
LOG(WARNING) << "Invalid configuration value for gain_mode_rx1 parameter. Set to gain_mode_rx1=" << default_gain_mode;
}
if ((gain_mode_rx2_ != "manual") and (gain_mode_rx2_ != "slow_attack") and (gain_mode_rx2_ != "fast_attack") and (gain_mode_rx2_ != "hybrid"))
if ((gain_mode_rx2_ != "manual") && (gain_mode_rx2_ != "slow_attack") && (gain_mode_rx2_ != "fast_attack") && (gain_mode_rx2_ != "hybrid"))
{
std::cout << "Configuration parameter gain_mode_rx2 should take one of these values:\n";
std::cout << " manual, slow_attack, fast_attack, hybrid\n";
@ -121,7 +121,7 @@ Fmcomms2SignalSource::Fmcomms2SignalSource(const ConfigurationInterface *configu
LOG(WARNING) << "Invalid configuration value for gain_mode_rx1 parameter. Set to gain_mode_rx2=" << default_gain_mode;
}
if ((filter_source_ != "Off") and (filter_source_ != "Auto") and (filter_source_ != "File") and (filter_source_ != "Design"))
if ((filter_source_ != "Off") && (filter_source_ != "Auto") && (filter_source_ != "File") && (filter_source_ != "Design"))
{
std::cout << "Configuration parameter filter_source should take one of these values:\n";
std::cout << " Off: Disable filter\n";
@ -138,7 +138,7 @@ Fmcomms2SignalSource::Fmcomms2SignalSource(const ConfigurationInterface *configu
if (gain_mode_rx1_ == "manual")
{
if (rf_gain_rx1_ > 73.0 or rf_gain_rx1_ < -1.0)
if (rf_gain_rx1_ > 73.0 || rf_gain_rx1_ < -1.0)
{
std::cout << "Configuration parameter rf_gain_rx1 should take values between -1.0 and 73 dB\n";
std::cout << "Error: provided value rf_gain_rx1=" << rf_gain_rx1_ << " is not among valid values\n";
@ -150,7 +150,7 @@ Fmcomms2SignalSource::Fmcomms2SignalSource(const ConfigurationInterface *configu
if (gain_mode_rx2_ == "manual")
{
if (rf_gain_rx2_ > 73.0 or rf_gain_rx2_ < -1.0)
if (rf_gain_rx2_ > 73.0 || rf_gain_rx2_ < -1.0)
{
std::cout << "Configuration parameter rf_gain_rx2 should take values between -1.0 and 73 dB\n";
std::cout << "Error: provided value rf_gain_rx2=" << rf_gain_rx2_ << " is not among valid values\n";
@ -160,7 +160,7 @@ Fmcomms2SignalSource::Fmcomms2SignalSource(const ConfigurationInterface *configu
}
}
if (bandwidth_ < 200000 or bandwidth_ > 56000000)
if (bandwidth_ < 200000 || bandwidth_ > 56000000)
{
std::cout << "Configuration parameter bandwidth should take values between 200000 and 56000000 Hz\n";
std::cout << "Error: provided value bandwidth=" << bandwidth_ << " is not among valid values\n";
@ -177,7 +177,7 @@ Fmcomms2SignalSource::Fmcomms2SignalSource(const ConfigurationInterface *configu
{
if (RF_channels_ == 1)
{
if (rx1_en_ and rx2_en_)
if (rx1_en_ && rx2_en_)
{
LOG(FATAL) << "Configuration error: both rx1 and rx2 are enabled but RF_channels=1 !";
}
@ -220,7 +220,7 @@ Fmcomms2SignalSource::Fmcomms2SignalSource(const ConfigurationInterface *configu
// configure LO
if (enable_dds_lo_ == true)
{
if (tx_bandwidth_ < static_cast<uint64_t>(std::floor(static_cast<float>(freq_dds_tx_hz_) * 1.1)) or (tx_bandwidth_ < 200000) or (tx_bandwidth_ > 1000000))
if (tx_bandwidth_ < static_cast<uint64_t>(std::floor(static_cast<float>(freq_dds_tx_hz_) * 1.1)) || (tx_bandwidth_ < 200000) || (tx_bandwidth_ > 1000000))
{
std::cout << "Configuration parameter tx_bandwidth value should be between " << std::max(static_cast<float>(freq_dds_tx_hz_) * 1.1, 200000.0) << " and 1000000 Hz\n";
std::cout << "Error: provided value tx_bandwidth=" << tx_bandwidth_ << " is not among valid values\n";
@ -228,7 +228,7 @@ Fmcomms2SignalSource::Fmcomms2SignalSource(const ConfigurationInterface *configu
tx_bandwidth_ = 500000;
LOG(WARNING) << "Invalid configuration value for tx_bandwidth parameter. Set to tx_bandwidth=500000";
}
if (tx_attenuation_db_ > 0.0 or tx_attenuation_db_ < -89.75)
if (tx_attenuation_db_ > 0.0 || tx_attenuation_db_ < -89.75)
{
std::cout << "Configuration parameter tx_attenuation_db should take values between 0.0 and -89.95 in 0.25 dB steps\n";
std::cout << "Error: provided value tx_attenuation_db=" << tx_attenuation_db_ << " is not among valid values\n";
@ -257,7 +257,7 @@ Fmcomms2SignalSource::Fmcomms2SignalSource(const ConfigurationInterface *configu
}
else if (RF_channels_ == 2)
{
if (!(rx1_en_ and rx2_en_))
if (!(rx1_en_ && rx2_en_))
{
LOG(FATAL) << "Configuration error: RF_channels=2 but are not enabled both receivers in FMCOMMS2 !";
}
@ -294,7 +294,7 @@ Fmcomms2SignalSource::Fmcomms2SignalSource(const ConfigurationInterface *configu
// configure LO
if (enable_dds_lo_ == true)
{
if (tx_bandwidth_ < static_cast<uint64_t>(std::floor(static_cast<float>(freq_dds_tx_hz_) * 1.1)) or (tx_bandwidth_ < 200000) or (tx_bandwidth_ > 1000000))
if (tx_bandwidth_ < static_cast<uint64_t>(std::floor(static_cast<float>(freq_dds_tx_hz_) * 1.1)) || (tx_bandwidth_ < 200000) || (tx_bandwidth_ > 1000000))
{
std::cout << "Configuration parameter tx_bandwidth value should be between " << std::max(static_cast<float>(freq_dds_tx_hz_) * 1.1, 200000.0) << " and 1000000 Hz\n";
std::cout << "Error: provided value tx_bandwidth=" << tx_bandwidth_ << " is not among valid values\n";
@ -302,7 +302,7 @@ Fmcomms2SignalSource::Fmcomms2SignalSource(const ConfigurationInterface *configu
tx_bandwidth_ = 500000;
LOG(WARNING) << "Invalid configuration value for tx_bandwidth parameter. Set to tx_bandwidth=500000";
}
if (tx_attenuation_db_ > 0.0 or tx_attenuation_db_ < -89.75)
if (tx_attenuation_db_ > 0.0 || tx_attenuation_db_ < -89.75)
{
std::cout << "Configuration parameter tx_attenuation_db should take values between 0.0 and -89.95 in 0.25 dB steps\n";
std::cout << "Error: provided value tx_attenuation_db=" << tx_attenuation_db_ << " is not among valid values\n";
@ -433,7 +433,7 @@ void Fmcomms2SignalSource::disconnect(gr::top_block_sptr top_block)
gr::basic_block_sptr Fmcomms2SignalSource::get_left_block()
{
LOG(WARNING) << "Trying to get signal source left block.";
return gr::basic_block_sptr();
return {};
}

2
src/algorithms/signal_source/adapters/limesdr_signal_source.cc
View File

@ -231,7 +231,7 @@ void LimesdrSignalSource::disconnect(gr::top_block_sptr top_block)
gr::basic_block_sptr LimesdrSignalSource::get_left_block()
{
LOG(WARNING) << "Trying to get signal source left block.";
return gr::basic_block_sptr();
return {};
}

2
src/algorithms/signal_source/adapters/osmosdr_signal_source.cc
View File

@ -210,7 +210,7 @@ void OsmosdrSignalSource::disconnect(gr::top_block_sptr top_block)
gr::basic_block_sptr OsmosdrSignalSource::get_left_block()
{
LOG(WARNING) << "Trying to get signal source left block.";
return gr::basic_block_sptr();
return {};
}

10
src/algorithms/signal_source/adapters/plutosdr_signal_source.cc
View File

@ -71,7 +71,7 @@ PlutosdrSignalSource::PlutosdrSignalSource(const ConfigurationInterface* configu
}
// basic check
if ((gain_mode_ != "manual") and (gain_mode_ != "slow_attack") and (gain_mode_ != "fast_attack") and (gain_mode_ != "hybrid"))
if ((gain_mode_ != "manual") && (gain_mode_ != "slow_attack") && (gain_mode_ != "fast_attack") && (gain_mode_ != "hybrid"))
{
std::cout << "Configuration parameter gain_mode should take one of these values:\n";
std::cout << " manual, slow_attack, fast_attack, hybrid\n";
@ -83,7 +83,7 @@ PlutosdrSignalSource::PlutosdrSignalSource(const ConfigurationInterface* configu
if (gain_mode_ == "manual")
{
if (rf_gain_ > 73.0 or rf_gain_ < -1.0)
if (rf_gain_ > 73.0 || rf_gain_ < -1.0)
{
std::cout << "Configuration parameter rf_gain should take values between -1.0 and 73 dB\n";
std::cout << "Error: provided value rf_gain=" << rf_gain_ << " is not among valid values\n";
@ -93,7 +93,7 @@ PlutosdrSignalSource::PlutosdrSignalSource(const ConfigurationInterface* configu
}
}
if ((filter_source_ != "Off") and (filter_source_ != "Auto") and (filter_source_ != "File") and (filter_source_ != "Design"))
if ((filter_source_ != "Off") && (filter_source_ != "Auto") && (filter_source_ != "File") && (filter_source_ != "Design"))
{
std::cout << "Configuration parameter filter_source should take one of these values:\n";
std::cout << " Off: Disable filter\n";
@ -108,7 +108,7 @@ PlutosdrSignalSource::PlutosdrSignalSource(const ConfigurationInterface* configu
LOG(WARNING) << "Invalid configuration value for filter_source parameter. Set to filter_source=Off";
}
if (bandwidth_ < 200000 or bandwidth_ > 56000000)
if (bandwidth_ < 200000 || bandwidth_ > 56000000)
{
std::cout << "Configuration parameter bandwidth should take values between 200000 and 56000000 Hz\n";
std::cout << "Error: provided value bandwidth=" << bandwidth_ << " is not among valid values\n";
@ -216,7 +216,7 @@ void PlutosdrSignalSource::disconnect(gr::top_block_sptr top_block)
gr::basic_block_sptr PlutosdrSignalSource::get_left_block()
{
LOG(WARNING) << "Trying to get signal source left block.";
return gr::basic_block_sptr();
return {};
}

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

@ -84,6 +84,7 @@ kf_vtl_tracking::kf_vtl_tracking(const Kf_Conf &conf_)
d_trk_parameters(conf_),
d_acquisition_gnss_synchro(nullptr),
d_signal_type(d_trk_parameters.signal),
d_code_chip_rate(0.0),
d_acq_code_phase_samples(0.0),
d_acq_carrier_doppler_hz(0.0),
d_current_correlation_time_s(0.0),

1
src/core/receiver/control_thread.cc
View File

@ -1207,6 +1207,7 @@ void ControlThread::keyboard_listener()
{
std::cout << "Quit keystroke order received, stopping GNSS-SDR !!\n";
control_queue_->push(pmt::make_any(command_event_make(200, 0)));
stop_ = true;
read_keys = false;
}
else

727
src/core/receiver/gnss_flowgraph.cc
View File

File diff suppressed because it is too large Load Diff

20
src/core/receiver/gnss_flowgraph.h
View File

@ -183,36 +183,18 @@ private:
int connect_tracking_monitor();
int connect_navdata_monitor();
int disconnect_desktop_flowgraph();
int disconnect_signal_sources();
int disconnect_signal_conditioners();
int disconnect_channels();
int disconnect_observables();
int disconnect_pvt();
int disconnect_sample_counter();
int disconnect_signal_sources_from_signal_conditioners();
int disconnect_signal_conditioners_from_channels();
int disconnect_channels_from_observables();
int disconnect_observables_from_pvt();
int disconnect_monitors();
#if ENABLE_FPGA
int connect_fpga_flowgraph();
int disconnect_fpga_flowgraph();
int connect_fpga_sample_counter();
int disconnect_fpga_sample_counter();
#endif
void assign_channels();
int assign_channels();
void check_signal_conditioners();
void set_signals_list();
void set_channels_state(); // Initializes the channels state (start acquisition or keep standby)
// using the configuration parameters (number of channels and max channels in acquisition)
Gnss_Signal search_next_signal(const std::string& searched_signal,
const bool pop,
bool& is_primary_frequency,
bool& assistance_available,
float& estimated_doppler,

2
src/tests/CMakeLists.txt
View File

@ -30,7 +30,7 @@ if(NOT GOOGLETEST_FOUND)
"--build" "${CMAKE_BINARY_DIR}/gtest-${GNSSSDR_GTEST_LOCAL_VERSION}"
"--config" $<$<CONFIG:Debug>:Debug>$<$<CONFIG:Release>:Release>$<$<CONFIG:RelWithDebInfo>:RelWithDebInfo>$<$<CONFIG:MinSizeRel>:MinSizeRel>$<$<CONFIG:NoOptWithASM>:Debug>$<$<CONFIG:Coverage>:Debug>$<$<CONFIG:O2WithASM>:RelWithDebInfo>$<$<CONFIG:O3WithASM>:RelWithDebInfo>$<$<CONFIG:ASAN>:Debug>
)
if(CMAKE_VERSION VERSION_GREATER 3.12)
if(CMAKE_VERSION VERSION_GREATER 3.12 AND NOT CMAKE_CROSSCOMPILING)
# Parallel building of gtest causes problems in some environments (e.g. Macports buildings)
set(GTEST_BUILD_COMMAND ${GTEST_BUILD_COMMAND} "--parallel 1")
endif()

78
src/tests/system-tests/libs/spirent_motion_csv_dump_reader.cc
View File

@ -22,46 +22,46 @@
Spirent_Motion_Csv_Dump_Reader::Spirent_Motion_Csv_Dump_Reader()
: header_lines(2),
TOW_ms(0.0),
Pos_X(0.0),
Pos_Y(0.0),
Pos_Z(0.0),
Vel_X(0.0),
Vel_Y(0.0),
Vel_Z(0.0),
Acc_X(0.0),
Acc_Y(0.0),
Acc_Z(0.0),
Jerk_X(0.0),
Jerk_Y(0.0),
Jerk_Z(0.0),
Lat(0.0),
Long(0.0),
Height(0.0),
Heading(0.0),
Elevation(0.0),
Bank(0.0),
Ang_vel_X(0.0),
Ang_vel_Y(0.0),
Ang_vel_Z(0.0),
Ang_acc_X(0.0),
Ang_acc_Y(0.0),
Ang_acc_Z(0.0),
Ant1_Pos_X(0.0),
Ant1_Pos_Y(0.0),
Ant1_Pos_Z(0.0),
Ant1_Vel_X(0.0),
Ant1_Vel_Y(0.0),
Ant1_Vel_Z(0.0),
Ant1_Acc_X(0.0),
Ant1_Acc_Y(0.0),
Ant1_Acc_Z(0.0),
Ant1_Lat(0.0),
Ant1_Long(0.0),
Ant1_Height(0.0),
Ant1_DOP(0.0)
{
header_lines = 2;
TOW_ms = 0.0;
Pos_X = 0.0;
Pos_Y = 0.0;
Pos_Z = 0.0;
Vel_X = 0.0;
Vel_Y = 0.0;
Vel_Z = 0.0;
Acc_X = 0.0;
Acc_Y = 0.0;
Acc_Z = 0.0;
Jerk_X = 0.0;
Jerk_Y = 0.0;
Jerk_Z = 0.0;
Lat = 0.0;
Long = 0.0;
Height = 0.0;
Heading = 0.0;
Elevation = 0.0;
Bank = 0.0;
Ang_vel_X = 0.0;
Ang_vel_Y = 0.0;
Ang_vel_Z = 0.0;
Ang_acc_X = 0.0;
Ang_acc_Y = 0.0;
Ang_acc_Z = 0.0;
Ant1_Pos_X = 0.0;
Ant1_Pos_Y = 0.0;
Ant1_Pos_Z = 0.0;
Ant1_Vel_X = 0.0;
Ant1_Vel_Y = 0.0;
Ant1_Vel_Z = 0.0;
Ant1_Acc_X = 0.0;
Ant1_Acc_Y = 0.0;
Ant1_Acc_Z = 0.0;
Ant1_Lat = 0.0;
Ant1_Long = 0.0;
Ant1_Height = 0.0;
Ant1_DOP = 0.0;
}

16
src/tests/unit-tests/signal-processing-blocks/acquisition/gps_l1_ca_pcps_acquisition_test.cc
View File

@ -71,7 +71,7 @@ private:
GpsL1CaPcpsAcquisitionTest_msg_rx();
public:
int rx_message;
int rx_message{0};
~GpsL1CaPcpsAcquisitionTest_msg_rx() override; //!< Default destructor
};
@ -97,7 +97,9 @@ void GpsL1CaPcpsAcquisitionTest_msg_rx::msg_handler_channel_events(const pmt::pm
}
GpsL1CaPcpsAcquisitionTest_msg_rx::GpsL1CaPcpsAcquisitionTest_msg_rx() : gr::block("GpsL1CaPcpsAcquisitionTest_msg_rx", gr::io_signature::make(0, 0, 0), gr::io_signature::make(0, 0, 0))
GpsL1CaPcpsAcquisitionTest_msg_rx::GpsL1CaPcpsAcquisitionTest_msg_rx()
: gr::block("GpsL1CaPcpsAcquisitionTest_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"),
@ -110,7 +112,6 @@ GpsL1CaPcpsAcquisitionTest_msg_rx::GpsL1CaPcpsAcquisitionTest_msg_rx() : gr::blo
boost::bind(&GpsL1CaPcpsAcquisitionTest_msg_rx::msg_handler_channel_events, this, _1));
#endif
#endif
rx_message = 0;
}
@ -123,12 +124,11 @@ class GpsL1CaPcpsAcquisitionTest : public ::testing::Test
{
protected:
GpsL1CaPcpsAcquisitionTest()
: item_size(sizeof(gr_complex))
{
config = std::make_shared<InMemoryConfiguration>();
item_size = sizeof(gr_complex);
gnss_synchro = Gnss_Synchro();
doppler_max = 5000;
doppler_step = 100;
}
~GpsL1CaPcpsAcquisitionTest() override = default;
@ -140,8 +140,8 @@ protected:
std::shared_ptr<InMemoryConfiguration> config;
Gnss_Synchro gnss_synchro{};
size_t item_size;
unsigned int doppler_max;
unsigned int doppler_step;
unsigned int doppler_max{5000};
unsigned int doppler_step{100};
};

4
src/tests/unit-tests/signal-processing-blocks/adapter/adapter_test.cc
View File

@ -53,9 +53,9 @@ public:
DataTypeAdapter::DataTypeAdapter()
: file_name_input("adapter_test_input.dat"),
file_name_output("adapter_test_output.dat")
{
file_name_input = "adapter_test_input.dat";
file_name_output = "adapter_test_output.dat";
std::array<int8_t, 6> input_bytes{2, 23, -1, 127, -127, 0};
std::array<int16_t, 8> input_shorts{2, 23, -1, 127, -127, 0, 255, 255};

3
src/tests/unit-tests/signal-processing-blocks/filter/fir_filter_test.cc
View File

@ -44,10 +44,9 @@ DEFINE_int32(filter_test_nsamples, 1000000, "Number of samples to filter in the
class FirFilterTest : public ::testing::Test
{
protected:
FirFilterTest()
FirFilterTest() : item_size(sizeof(gr_complex))
{
queue = std::make_shared<Concurrent_Queue<pmt::pmt_t>>();
item_size = sizeof(gr_complex);
config = std::make_shared<InMemoryConfiguration>();
}
~FirFilterTest() override = default;

4
src/tests/unit-tests/signal-processing-blocks/filter/notch_filter_lite_test.cc
View File

@ -43,12 +43,10 @@ DEFINE_int32(notch_filter_lite_test_nsamples, 1000000, "Number of samples to fil
class NotchFilterLiteTest : public ::testing::Test
{
protected:
NotchFilterLiteTest()
NotchFilterLiteTest() : item_size(sizeof(gr_complex)), nsamples(FLAGS_notch_filter_lite_test_nsamples)
{
queue = std::make_shared<Concurrent_Queue<pmt::pmt_t>>();
item_size = sizeof(gr_complex);
config = std::make_shared<InMemoryConfiguration>();
nsamples = FLAGS_notch_filter_lite_test_nsamples;
}
~NotchFilterLiteTest() override = default;

4
src/tests/unit-tests/signal-processing-blocks/filter/notch_filter_test.cc
View File

@ -43,12 +43,10 @@ DEFINE_int32(notch_filter_test_nsamples, 1000000, "Number of samples to filter i
class NotchFilterTest : public ::testing::Test
{
protected:
NotchFilterTest()
NotchFilterTest() : item_size(sizeof(gr_complex)), nsamples(FLAGS_notch_filter_test_nsamples)
{
queue = std::make_shared<Concurrent_Queue<pmt::pmt_t>>();
item_size = sizeof(gr_complex);
config = std::make_shared<InMemoryConfiguration>();
nsamples = FLAGS_notch_filter_test_nsamples;
}
~NotchFilterTest() override = default;

4
src/tests/unit-tests/signal-processing-blocks/filter/pulse_blanking_filter_test.cc
View File

@ -43,12 +43,10 @@ DEFINE_int32(pb_filter_test_nsamples, 1000000, "Number of samples to filter in t
class PulseBlankingFilterTest : public ::testing::Test
{
protected:
PulseBlankingFilterTest()
PulseBlankingFilterTest() : item_size(sizeof(gr_complex)), nsamples(FLAGS_pb_filter_test_nsamples)
{
queue = std::make_shared<Concurrent_Queue<pmt::pmt_t>>();
item_size = sizeof(gr_complex);
config = std::make_shared<InMemoryConfiguration>();
nsamples = FLAGS_pb_filter_test_nsamples;
}
~PulseBlankingFilterTest() override = default;
bool stop = false;

19
src/tests/unit-tests/signal-processing-blocks/libs/acquisition_dump_reader.cc
View File

@ -196,21 +196,12 @@ Acquisition_Dump_Reader::Acquisition_Dump_Reader(const std::string& basename,
unsigned int samples_per_code,
int channel,
int execution)
: d_basename(basename),
d_sat(sat),
d_doppler_max(doppler_max),
d_doppler_step(doppler_step),
d_samples_per_code(samples_per_code)
{
d_basename = basename;
d_sat = sat;
d_doppler_max = doppler_max;
d_doppler_step = doppler_step;
d_samples_per_code = samples_per_code;
acq_doppler_hz = 0.0;
acq_delay_samples = 0.0;
test_statistic = 0.0;
input_power = 0.0;
threshold = 0.0;
positive_acq = 0;
sample_counter = 0;
num_dwells = 0;
PRN = 0;
if (d_doppler_step == 0)
{
d_doppler_step = 1;

5
src/tests/unit-tests/signal-processing-blocks/libs/acquisition_msg_rx.cc
View File

@ -56,7 +56,9 @@ void Acquisition_msg_rx::msg_handler_channel_events(const pmt::pmt_t& msg)
}
Acquisition_msg_rx::Acquisition_msg_rx() : gr::block("Acquisition_msg_rx", gr::io_signature::make(0, 0, 0), gr::io_signature::make(0, 0, 0))
Acquisition_msg_rx::Acquisition_msg_rx()
: gr::block("Acquisition_msg_rx", gr::io_signature::make(0, 0, 0), gr::io_signature::make(0, 0, 0)),
rx_message(0)
{
this->message_port_register_in(pmt::mp("events"));
this->set_msg_handler(pmt::mp("events"),
@ -69,7 +71,6 @@ Acquisition_msg_rx::Acquisition_msg_rx() : gr::block("Acquisition_msg_rx", gr::i
boost::bind(&Acquisition_msg_rx::msg_handler_channel_events, this, _1));
#endif
#endif
rx_message = 0;
}

2
src/tests/unit-tests/signal-processing-blocks/libs/observables_dump_reader.cc
View File

@ -104,8 +104,8 @@ void Observables_Dump_Reader::close_obs_file()
Observables_Dump_Reader::Observables_Dump_Reader(int n_channels_)
: n_channels(n_channels_)
{
n_channels = n_channels_;
RX_time = std::vector<double>(n_channels);
TOW_at_current_symbol_s = std::vector<double>(n_channels);
Carrier_Doppler_hz = std::vector<double>(n_channels);

4
src/tests/unit-tests/signal-processing-blocks/tracking/cubature_filter_test.cc
View File

@ -25,7 +25,7 @@
class TransitionModel : public ModelFunction
{
public:
explicit TransitionModel(const arma::mat& kf_F) { coeff_mat = kf_F; };
explicit TransitionModel(const arma::mat& kf_F) : coeff_mat(kf_F){};
arma::vec operator()(const arma::vec& input) override { return coeff_mat * input; };
private:
@ -35,7 +35,7 @@ private:
class MeasurementModel : public ModelFunction
{
public:
explicit MeasurementModel(const arma::mat& kf_H) { coeff_mat = kf_H; };
explicit MeasurementModel(const arma::mat& kf_H) : coeff_mat(kf_H){};
arma::vec operator()(const arma::vec& input) override { return coeff_mat * input; };
private:

8
src/tests/unit-tests/signal-processing-blocks/tracking/galileo_e1_dll_pll_veml_tracking_test.cc
View File

@ -41,12 +41,10 @@ class GalileoE1DllPllVemlTrackingInternalTest : public ::testing::Test
{
protected:
GalileoE1DllPllVemlTrackingInternalTest()
: item_size(sizeof(gr_complex))
{
factory = std::make_shared<GNSSBlockFactory>();
config = std::make_shared<InMemoryConfiguration>();
item_size = sizeof(gr_complex);
stop = false;
message = 0;
gnss_synchro = Gnss_Synchro();
}
@ -60,8 +58,8 @@ protected:
std::shared_ptr<InMemoryConfiguration> config;
Gnss_Synchro gnss_synchro{};
size_t item_size;
bool stop;
int message;
bool stop{false};
int message{0};
};

8
src/tests/unit-tests/signal-processing-blocks/tracking/gps_l1_ca_dll_pll_tracking_test.cc
View File

@ -102,7 +102,9 @@ void GpsL1CADllPllTrackingTest_msg_rx::msg_handler_channel_events(const pmt::pmt
}
GpsL1CADllPllTrackingTest_msg_rx::GpsL1CADllPllTrackingTest_msg_rx() : gr::block("GpsL1CADllPllTrackingTest_msg_rx", gr::io_signature::make(0, 0, 0), gr::io_signature::make(0, 0, 0))
GpsL1CADllPllTrackingTest_msg_rx::GpsL1CADllPllTrackingTest_msg_rx()
: gr::block("GpsL1CADllPllTrackingTest_msg_rx", gr::io_signature::make(0, 0, 0), gr::io_signature::make(0, 0, 0)),
rx_message(0)
{
this->message_port_register_in(pmt::mp("events"));
this->set_msg_handler(pmt::mp("events"),
@ -115,7 +117,6 @@ GpsL1CADllPllTrackingTest_msg_rx::GpsL1CADllPllTrackingTest_msg_rx() : gr::block
boost::bind(&GpsL1CADllPllTrackingTest_msg_rx::msg_handler_channel_events, this, _1));
#endif
#endif
rx_message = 0;
}
@ -166,11 +167,10 @@ public:
double& rmse);
bool save_mat_xy(std::vector<double>& x, std::vector<double>& y, std::string filename);
GpsL1CADllPllTrackingTest()
GpsL1CADllPllTrackingTest() : item_size(sizeof(gr_complex))
{
factory = std::make_shared<GNSSBlockFactory>();
config = std::make_shared<InMemoryConfiguration>();
item_size = sizeof(gr_complex);
gnss_synchro = Gnss_Synchro();
}

13
src/tests/unit-tests/signal-processing-blocks/tracking/gps_l1_ca_dll_pll_tracking_test_fpga.cc
View File

@ -197,9 +197,11 @@ void GpsL1CADllPllTrackingTestFpga_msg_rx::msg_handler_channel_events(const pmt:
}
GpsL1CADllPllTrackingTestFpga_msg_rx::GpsL1CADllPllTrackingTestFpga_msg_rx() : gr::block("GpsL1CADllPllTrackingTestFpga_msg_rx",
gr::io_signature::make(0, 0, 0),
gr::io_signature::make(0, 0, 0))
GpsL1CADllPllTrackingTestFpga_msg_rx::GpsL1CADllPllTrackingTestFpga_msg_rx()
: gr::block("GpsL1CADllPllTrackingTestFpga_msg_rx",
gr::io_signature::make(0, 0, 0),
gr::io_signature::make(0, 0, 0)),
rx_message(0)
{
this->message_port_register_in(pmt::mp("events"));
this->set_msg_handler(pmt::mp("events"),
@ -212,7 +214,6 @@ GpsL1CADllPllTrackingTestFpga_msg_rx::GpsL1CADllPllTrackingTestFpga_msg_rx() : g
boost::bind(&GpsL1CADllPllTrackingTestFpga_msg_rx::msg_handler_channel_events, this, _1));
#endif
#endif
rx_message = 0;
}
@ -245,10 +246,10 @@ public:
void check_results_codephase(arma::vec &true_time_s, arma::vec &true_value,
arma::vec &meas_time_s, arma::vec &meas_value);
GpsL1CADllPllTrackingTestFpga()
GpsL1CADllPllTrackingTestFpga() : item_size(sizeof(gr_complex))
{
config = std::make_shared<InMemoryConfiguration>();
item_size = sizeof(gr_complex);
gnss_synchro = Gnss_Synchro();
}

6
src/utils/front-end-cal/main.cc
View File

@ -117,7 +117,7 @@ private:
FrontEndCal_msg_rx();
public:
int rx_message;
int rx_message{0};
};
@ -143,7 +143,8 @@ void FrontEndCal_msg_rx::msg_handler_channel_events(const pmt::pmt_t& msg)
}
FrontEndCal_msg_rx::FrontEndCal_msg_rx() : gr::block("FrontEndCal_msg_rx", gr::io_signature::make(0, 0, 0), gr::io_signature::make(0, 0, 0))
FrontEndCal_msg_rx::FrontEndCal_msg_rx()
: gr::block("FrontEndCal_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"),
@ -156,7 +157,6 @@ FrontEndCal_msg_rx::FrontEndCal_msg_rx() : gr::block("FrontEndCal_msg_rx", gr::i
boost::bind(&FrontEndCal_msg_rx::msg_handler_channel_events, this, _1));
#endif
#endif
rx_message = 0;
}