mirror of
https://github.com/gnss-sdr/gnss-sdr
synced 2025-01-31 11:19:18 +00:00
Merge branch 'next' of https://github.com/gnss-sdr/gnss-sdr into ssr
This commit is contained in:
commit
ebf9aaf9cc
409
.clang-tidy
409
.clang-tidy
@ -56,6 +56,7 @@ Checks: '-*,
|
|||||||
clang-analyzer-optin.performance.*,
|
clang-analyzer-optin.performance.*,
|
||||||
clang-analyzer-optin.portability.UnixAPI,
|
clang-analyzer-optin.portability.UnixAPI,
|
||||||
clang-analyzer-security.*,
|
clang-analyzer-security.*,
|
||||||
|
cppcoreguidelines-prefer-member-initializer,
|
||||||
cppcoreguidelines-pro-type-cstyle-cast,
|
cppcoreguidelines-pro-type-cstyle-cast,
|
||||||
cppcoreguidelines-pro-type-static-cast-downcast,
|
cppcoreguidelines-pro-type-static-cast-downcast,
|
||||||
cppcoreguidelines-slicing,
|
cppcoreguidelines-slicing,
|
||||||
@ -120,70 +121,16 @@ HeaderFilterRegex: ''
|
|||||||
AnalyzeTemporaryDtors: false
|
AnalyzeTemporaryDtors: false
|
||||||
FormatStyle: 'file'
|
FormatStyle: 'file'
|
||||||
CheckOptions:
|
CheckOptions:
|
||||||
- key: bugprone-argument-comment.CommentBoolLiterals
|
- key: performance-unnecessary-copy-initialization.ExcludedContainerTypes
|
||||||
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
|
|
||||||
value: ''
|
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
|
- key: bugprone-reserved-identifier.Invert
|
||||||
value: 'false'
|
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
|
- key: bugprone-unused-return-value.CheckedFunctions
|
||||||
value: 'std::async;
|
value: 'std::async;
|
||||||
std::launder;
|
std::launder;
|
||||||
@ -276,148 +223,218 @@ std::vector::at;
|
|||||||
::ttyname'
|
::ttyname'
|
||||||
- key: cert-dcl16-c.NewSuffixes
|
- key: cert-dcl16-c.NewSuffixes
|
||||||
value: 'L;LL;LU;LLU'
|
value: 'L;LL;LU;LLU'
|
||||||
- key: cert-msc51-cpp.DisallowedSeedTypes
|
- key: readability-identifier-naming.GetConfigPerFile
|
||||||
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
|
|
||||||
value: 'true'
|
value: 'true'
|
||||||
- key: readability-inconsistent-declaration-parameter-name.Strict
|
- key: readability-inconsistent-declaration-parameter-name.Strict
|
||||||
value: 'false'
|
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
|
- key: readability-qualified-auto.AddConstToQualified
|
||||||
value: 'true'
|
value: 'true'
|
||||||
- key: readability-redundant-member-init.IgnoreBaseInCopyConstructors
|
- key: bugprone-sizeof-expression.WarnOnSizeOfThis
|
||||||
value: 'false'
|
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
|
- key: readability-uppercase-literal-suffix.IgnoreMacros
|
||||||
value: 'true'
|
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
|
- key: readability-uppercase-literal-suffix.NewSuffixes
|
||||||
value: ''
|
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'
|
||||||
...
|
...
|
||||||
|
@ -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(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)
|
if(ENABLE_PACKAGING)
|
||||||
set(ENABLE_GENERIC_ARCH ON)
|
set(ENABLE_GENERIC_ARCH ON)
|
||||||
set(ENABLE_ARMA_NO_DEBUG 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()
|
endif()
|
||||||
if(FILESYSTEM_FOUND)
|
if(FILESYSTEM_FOUND)
|
||||||
set(CMAKE_CXX_STANDARD 17)
|
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
|
# UHD 3.15.0.0-5 does not support C++20
|
||||||
# GNU Radio 3.10.0.git 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))
|
# 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)
|
# set(CMAKE_CXX_STANDARD 20)
|
||||||
if(CMAKE_VERSION VERSION_GREATER 3.20.99)
|
# 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
|
# 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")))
|
# ((CMAKE_CXX_COMPILER_ID STREQUAL "Clang") AND NOT (CMAKE_CXX_COMPILER_VERSION VERSION_LESS "12.0")))
|
||||||
# set(CMAKE_CXX_STANDARD 23)
|
# set(CMAKE_CXX_STANDARD 23)
|
||||||
endif()
|
# endif()
|
||||||
endif()
|
# endif()
|
||||||
endif()
|
# endif()
|
||||||
endif()
|
# endif()
|
||||||
set(CMAKE_CXX_STANDARD_REQUIRED ON)
|
set(CMAKE_CXX_STANDARD_REQUIRED ON)
|
||||||
endif()
|
endif()
|
||||||
endif()
|
endif()
|
||||||
@ -711,9 +713,6 @@ set(Boost_ADDITIONAL_VERSIONS
|
|||||||
"1.70.0" "1.70" "1.71.0" "1.71"
|
"1.70.0" "1.70" "1.71.0" "1.71"
|
||||||
)
|
)
|
||||||
set(Boost_USE_MULTITHREAD ON)
|
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)
|
set(BOOST_COMPONENTS atomic chrono date_time serialization system thread)
|
||||||
if(NOT ${FILESYSTEM_FOUND})
|
if(NOT ${FILESYSTEM_FOUND})
|
||||||
set(BOOST_COMPONENTS ${BOOST_COMPONENTS} filesystem)
|
set(BOOST_COMPONENTS ${BOOST_COMPONENTS} filesystem)
|
||||||
@ -1807,6 +1806,8 @@ if(NOT TARGET BLAS::BLAS)
|
|||||||
)
|
)
|
||||||
endif()
|
endif()
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
################################################################################
|
################################################################################
|
||||||
# Check that LAPACK (Linear Algebra PACKage) is found in the system
|
# Check that LAPACK (Linear Algebra PACKage) is found in the system
|
||||||
# See https://www.netlib.org/lapack/
|
# See https://www.netlib.org/lapack/
|
||||||
@ -1841,6 +1842,7 @@ if(NOT TARGET LAPACK::LAPACK)
|
|||||||
endif()
|
endif()
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
################################################################################
|
################################################################################
|
||||||
# Armadillo - http://arma.sourceforge.net/
|
# Armadillo - http://arma.sourceforge.net/
|
||||||
################################################################################
|
################################################################################
|
||||||
@ -2981,6 +2983,7 @@ else()
|
|||||||
endif()
|
endif()
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
##########################################
|
##########################################
|
||||||
# gr-limesdr - OPTIONAL
|
# gr-limesdr - OPTIONAL
|
||||||
# https://github.com/myriadrf/gr-limesdr
|
# https://github.com/myriadrf/gr-limesdr
|
||||||
@ -3008,6 +3011,7 @@ else()
|
|||||||
endif()
|
endif()
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
##############################################
|
##############################################
|
||||||
# gr-iio - OPTIONAL
|
# gr-iio - OPTIONAL
|
||||||
# IIO blocks for GNU Radio
|
# IIO blocks for GNU Radio
|
||||||
@ -3054,6 +3058,7 @@ if(ENABLE_PLUTOSDR OR ENABLE_FMCOMMS2)
|
|||||||
endif()
|
endif()
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
#####################################################################
|
#####################################################################
|
||||||
# libiio - OPTIONAL
|
# libiio - OPTIONAL
|
||||||
# A library for interfacing with local and remote Linux IIO devices
|
# A library for interfacing with local and remote Linux IIO devices
|
||||||
@ -3081,6 +3086,7 @@ if(ENABLE_AD9361 OR ENABLE_FMCOMMS2)
|
|||||||
endif()
|
endif()
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
##############################################
|
##############################################
|
||||||
# TELEORBIT FLEXIBAND FRONTEND - OPTIONAL
|
# TELEORBIT FLEXIBAND FRONTEND - OPTIONAL
|
||||||
##############################################
|
##############################################
|
||||||
@ -3112,6 +3118,7 @@ if(ENABLE_FLEXIBAND)
|
|||||||
endif()
|
endif()
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
#######################################################
|
#######################################################
|
||||||
# CTTC's digital array beamformer prototype - OPTIONAL
|
# 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_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_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(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 "")
|
||||||
message(STATUS "***************************************")
|
message(STATUS "***************************************")
|
||||||
|
@ -5,7 +5,7 @@
|
|||||||
# SPDX-License-Identifier: BSD-3-Clause
|
# 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>
|
# Usage: cmake -DCMAKE_TOOLCHAIN_FILE=<this file> <source directory>
|
||||||
########################################################################
|
########################################################################
|
||||||
set(CMAKE_CXX_COMPILER g++)
|
set(CMAKE_CXX_COMPILER g++)
|
||||||
|
@ -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
|
readability of the code, could potentially increase performance, and allows
|
||||||
for easier detection of unused data members (see
|
for easier detection of unused data members (see
|
||||||
https://github.com/isocpp/CppCoreGuidelines/blob/master/CppCoreGuidelines.md/#Rc-initialize).
|
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: Fixed formatting defects detected by clang-format 13.0.
|
||||||
|
- Non-functional change: Simplified flow graph disconnection.
|
||||||
- Updated GSL implementation to v0.40.0. See
|
- Updated GSL implementation to v0.40.0. See
|
||||||
https://github.com/gsl-lite/gsl-lite/releases/tag/v0.40.0
|
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
|
- 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
|
- 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
|
output that indicates if the PLL got locked at 180 degrees, so the symbol sign
|
||||||
is reversed.
|
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
|
See the definitions of concepts and metrics at
|
||||||
https://gnss-sdr.org/design-forces/
|
https://gnss-sdr.org/design-forces/
|
||||||
|
@ -224,7 +224,7 @@ void GpsL1CaPcpsOpenClAcquisition::reset()
|
|||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
float GpsL1CaPcpsOpenClAcquisition::calculate_threshold(float pfa)
|
float GpsL1CaPcpsOpenClAcquisition::calculate_threshold(float pfa) const
|
||||||
{
|
{
|
||||||
// Calculate the threshold
|
// Calculate the threshold
|
||||||
unsigned int frequency_bins = 0;
|
unsigned int frequency_bins = 0;
|
||||||
|
@ -146,7 +146,7 @@ public:
|
|||||||
}
|
}
|
||||||
|
|
||||||
private:
|
private:
|
||||||
float calculate_threshold(float pfa);
|
float calculate_threshold(float pfa) const;
|
||||||
const ConfigurationInterface* configuration_;
|
const ConfigurationInterface* configuration_;
|
||||||
pcps_opencl_acquisition_cc_sptr acquisition_cc_;
|
pcps_opencl_acquisition_cc_sptr acquisition_cc_;
|
||||||
gr::blocks::stream_to_vector::sptr stream_to_vector_;
|
gr::blocks::stream_to_vector::sptr stream_to_vector_;
|
||||||
|
@ -89,8 +89,8 @@ Channel::Channel(const ConfigurationInterface* configuration,
|
|||||||
acq_->set_threshold(threshold);
|
acq_->set_threshold(threshold);
|
||||||
|
|
||||||
acq_->init();
|
acq_->init();
|
||||||
|
repeat_ = configuration->property("Acquisition_" + signal_str + ".repeat_satellite", false);
|
||||||
repeat_ = configuration->property("Acquisition_" + signal_str + std::to_string(channel_) + ".repeat_satellite", false);
|
repeat_ = configuration->property("Acquisition_" + signal_str + std::to_string(channel_) + ".repeat_satellite", repeat_);
|
||||||
DLOG(INFO) << "Channel " << channel_ << " satellite repeat = " << repeat_;
|
DLOG(INFO) << "Channel " << channel_ << " satellite repeat = " << repeat_;
|
||||||
|
|
||||||
channel_fsm_->set_acquisition(acq_);
|
channel_fsm_->set_acquisition(acq_);
|
||||||
|
@ -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);
|
size_t tmpLength = plan->n.x * plan->n.y * plan->n.z * batchSize * 2 * sizeof(cl_float);
|
||||||
|
|
||||||
if (plan->tempmemobj)
|
if (plan->tempmemobj)
|
||||||
clReleaseMemObject(plan->tempmemobj);
|
{
|
||||||
|
clReleaseMemObject(plan->tempmemobj);
|
||||||
|
}
|
||||||
|
|
||||||
plan->tempmemobj = clCreateBuffer(plan->context, CL_MEM_READ_WRITE, tmpLength, nullptr, &err);
|
plan->tempmemobj = clCreateBuffer(plan->context, CL_MEM_READ_WRITE, tmpLength, nullptr, &err);
|
||||||
}
|
}
|
||||||
return err;
|
return err;
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
static cl_int
|
static cl_int
|
||||||
allocateTemporaryBufferPlannar(cl_fft_plan *plan, cl_uint batchSize)
|
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);
|
size_t tmpLength = plan->n.x * plan->n.y * plan->n.z * batchSize * sizeof(cl_float);
|
||||||
|
|
||||||
if (plan->tempmemobj_real)
|
if (plan->tempmemobj_real)
|
||||||
clReleaseMemObject(plan->tempmemobj_real);
|
{
|
||||||
|
clReleaseMemObject(plan->tempmemobj_real);
|
||||||
|
}
|
||||||
|
|
||||||
if (plan->tempmemobj_imag)
|
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_real = clCreateBuffer(plan->context, CL_MEM_READ_WRITE, tmpLength, nullptr, &err);
|
||||||
plan->tempmemobj_imag = clCreateBuffer(plan->context, CL_MEM_READ_WRITE, tmpLength, nullptr, &terr);
|
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;
|
return err;
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
void getKernelWorkDimensions(cl_fft_plan *plan, cl_fft_kernel_info *kernelInfo, cl_int *batchSize, size_t *gWorkItems, size_t *lWorkItems)
|
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;
|
*lWorkItems = kernelInfo->num_workitems_per_workgroup;
|
||||||
@ -83,6 +91,7 @@ void getKernelWorkDimensions(cl_fft_plan *plan, cl_fft_kernel_info *kernelInfo,
|
|||||||
*gWorkItems = numWorkGroups * *lWorkItems;
|
*gWorkItems = numWorkGroups * *lWorkItems;
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
cl_int
|
cl_int
|
||||||
clFFT_ExecuteInterleaved(cl_command_queue queue, clFFT_Plan Plan, cl_int batchSize, clFFT_Direction dir,
|
clFFT_ExecuteInterleaved(cl_command_queue queue, clFFT_Plan Plan, cl_int batchSize, clFFT_Direction dir,
|
||||||
cl_mem data_in, cl_mem data_out,
|
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;
|
int s;
|
||||||
auto *plan = (cl_fft_plan *)Plan;
|
auto *plan = (cl_fft_plan *)Plan;
|
||||||
if (plan->format != clFFT_InterleavedComplexFormat)
|
if (plan->format != clFFT_InterleavedComplexFormat)
|
||||||
return CL_INVALID_VALUE;
|
{
|
||||||
|
return CL_INVALID_VALUE;
|
||||||
|
}
|
||||||
|
|
||||||
cl_int err;
|
cl_int err;
|
||||||
size_t gWorkItems;
|
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;
|
cl_int isInPlace = data_in == data_out ? 1 : 0;
|
||||||
|
|
||||||
if ((err = allocateTemporaryBufferInterleaved(plan, batchSize)) != CL_SUCCESS)
|
if ((err = allocateTemporaryBufferInterleaved(plan, batchSize)) != CL_SUCCESS)
|
||||||
return err;
|
{
|
||||||
|
return err;
|
||||||
|
}
|
||||||
|
|
||||||
cl_mem memObj[3];
|
cl_mem memObj[3];
|
||||||
memObj[0] = data_in;
|
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);
|
err |= clEnqueueNDRangeKernel(queue, kernelInfo->kernel, 1, nullptr, &gWorkItems, &lWorkItems, 0, nullptr, nullptr);
|
||||||
if (err)
|
if (err)
|
||||||
return err;
|
{
|
||||||
|
return err;
|
||||||
|
}
|
||||||
|
|
||||||
currRead = (currWrite == 1) ? 1 : 2;
|
currRead = (currWrite == 1) ? 1 : 2;
|
||||||
currWrite = (currWrite == 1) ? 2 : 1;
|
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);
|
err |= clEnqueueNDRangeKernel(queue, kernelInfo->kernel, 1, nullptr, &gWorkItems, &lWorkItems, 0, nullptr, nullptr);
|
||||||
if (err)
|
if (err)
|
||||||
return err;
|
{
|
||||||
|
return err;
|
||||||
|
}
|
||||||
|
|
||||||
currRead = 1;
|
currRead = 1;
|
||||||
currWrite = 1;
|
currWrite = 1;
|
||||||
@ -181,6 +198,7 @@ clFFT_ExecuteInterleaved(cl_command_queue queue, clFFT_Plan Plan, cl_int batchSi
|
|||||||
return err;
|
return err;
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
cl_int
|
cl_int
|
||||||
clFFT_ExecutePlannar(cl_command_queue queue, clFFT_Plan Plan, cl_int batchSize, clFFT_Direction dir,
|
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,
|
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;
|
auto *plan = (cl_fft_plan *)Plan;
|
||||||
|
|
||||||
if (plan->format != clFFT_SplitComplexFormat)
|
if (plan->format != clFFT_SplitComplexFormat)
|
||||||
return CL_INVALID_VALUE;
|
{
|
||||||
|
return CL_INVALID_VALUE;
|
||||||
|
}
|
||||||
|
|
||||||
cl_int err;
|
cl_int err;
|
||||||
size_t gWorkItems;
|
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;
|
cl_int isInPlace = ((data_in_real == data_out_real) && (data_in_imag == data_out_imag)) ? 1 : 0;
|
||||||
|
|
||||||
if ((err = allocateTemporaryBufferPlannar(plan, batchSize)) != CL_SUCCESS)
|
if ((err = allocateTemporaryBufferPlannar(plan, batchSize)) != CL_SUCCESS)
|
||||||
return err;
|
{
|
||||||
|
return err;
|
||||||
|
}
|
||||||
|
|
||||||
cl_mem memObj_real[3];
|
cl_mem memObj_real[3];
|
||||||
cl_mem memObj_imag[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);
|
err |= clEnqueueNDRangeKernel(queue, kernelInfo->kernel, 1, nullptr, &gWorkItems, &lWorkItems, 0, nullptr, nullptr);
|
||||||
if (err)
|
if (err)
|
||||||
return err;
|
{
|
||||||
|
return err;
|
||||||
|
}
|
||||||
|
|
||||||
currRead = (currWrite == 1) ? 1 : 2;
|
currRead = (currWrite == 1) ? 1 : 2;
|
||||||
currWrite = (currWrite == 1) ? 2 : 1;
|
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);
|
err |= clEnqueueNDRangeKernel(queue, kernelInfo->kernel, 1, nullptr, &gWorkItems, &lWorkItems, 0, nullptr, nullptr);
|
||||||
if (err)
|
if (err)
|
||||||
return err;
|
{
|
||||||
|
return err;
|
||||||
|
}
|
||||||
|
|
||||||
currRead = 1;
|
currRead = 1;
|
||||||
currWrite = 1;
|
currWrite = 1;
|
||||||
@ -288,6 +314,7 @@ clFFT_ExecutePlannar(cl_command_queue queue, clFFT_Plan Plan, cl_int batchSize,
|
|||||||
return err;
|
return err;
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
cl_int
|
cl_int
|
||||||
clFFT_1DTwistInterleaved(clFFT_Plan Plan, cl_command_queue queue, cl_mem array,
|
clFFT_1DTwistInterleaved(clFFT_Plan Plan, cl_command_queue queue, cl_mem array,
|
||||||
unsigned numRows, unsigned numCols, unsigned startRow, unsigned rowsToProcess, clFFT_Direction dir)
|
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;
|
cl_device_id device_id;
|
||||||
err = clGetCommandQueueInfo(queue, CL_QUEUE_DEVICE, sizeof(cl_device_id), &device_id, nullptr);
|
err = clGetCommandQueueInfo(queue, CL_QUEUE_DEVICE, sizeof(cl_device_id), &device_id, nullptr);
|
||||||
if (err)
|
if (err)
|
||||||
return err;
|
{
|
||||||
|
return err;
|
||||||
|
}
|
||||||
|
|
||||||
size_t gSize;
|
size_t gSize;
|
||||||
err = clGetKernelWorkGroupInfo(plan->twist_kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &gSize, nullptr);
|
err = clGetKernelWorkGroupInfo(plan->twist_kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &gSize, nullptr);
|
||||||
if (err)
|
if (err)
|
||||||
return err;
|
{
|
||||||
|
return err;
|
||||||
|
}
|
||||||
|
|
||||||
gSize = min(128, gSize);
|
gSize = min(128, gSize);
|
||||||
size_t numGlobalThreads[1] = {max(numCols / gSize, 1) * 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;
|
return err;
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
cl_int
|
cl_int
|
||||||
clFFT_1DTwistPlannar(clFFT_Plan Plan, cl_command_queue queue, cl_mem array_real, cl_mem array_imag,
|
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)
|
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;
|
cl_device_id device_id;
|
||||||
err = clGetCommandQueueInfo(queue, CL_QUEUE_DEVICE, sizeof(cl_device_id), &device_id, nullptr);
|
err = clGetCommandQueueInfo(queue, CL_QUEUE_DEVICE, sizeof(cl_device_id), &device_id, nullptr);
|
||||||
if (err)
|
if (err)
|
||||||
return err;
|
{
|
||||||
|
return err;
|
||||||
|
}
|
||||||
|
|
||||||
size_t gSize;
|
size_t gSize;
|
||||||
err = clGetKernelWorkGroupInfo(plan->twist_kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &gSize, nullptr);
|
err = clGetKernelWorkGroupInfo(plan->twist_kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &gSize, nullptr);
|
||||||
if (err)
|
if (err)
|
||||||
return err;
|
{
|
||||||
|
return err;
|
||||||
|
}
|
||||||
|
|
||||||
gSize = min(128, gSize);
|
gSize = min(128, gSize);
|
||||||
size_t numGlobalThreads[1] = {max(numCols / gSize, 1) * gSize};
|
size_t numGlobalThreads[1] = {max(numCols / gSize, 1) * gSize};
|
||||||
|
@ -30,7 +30,7 @@ num2str(int num)
|
|||||||
{
|
{
|
||||||
char temp[200];
|
char temp[200];
|
||||||
snprintf(temp, sizeof(temp), "%d", num);
|
snprintf(temp, sizeof(temp), "%d", num);
|
||||||
return string(temp);
|
return {temp};
|
||||||
}
|
}
|
||||||
|
|
||||||
// For any n, this function decomposes n into factors for loacal memory tranpose
|
// 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
|
static void
|
||||||
insertHeader(string &kernelString, string &kernelName, clFFT_DataFormat dataFormat)
|
insertHeader(string &kernelString, string &kernelName, clFFT_DataFormat dataFormat)
|
||||||
{
|
{
|
||||||
if (dataFormat == clFFT_SplitComplexFormat)
|
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
|
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
|
static void
|
||||||
insertVariables(string &kStream, int maxRadix)
|
insertVariables(string &kStream, int maxRadix)
|
||||||
{
|
{
|
||||||
@ -177,11 +183,14 @@ insertVariables(string &kStream, int maxRadix)
|
|||||||
kStream += string(" int groupId = get_group_id( 0 );\n");
|
kStream += string(" int groupId = get_group_id( 0 );\n");
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
static void
|
static void
|
||||||
formattedLoad(string &kernelString, int aIndex, int gIndex, clFFT_DataFormat dataFormat)
|
formattedLoad(string &kernelString, int aIndex, int gIndex, clFFT_DataFormat dataFormat)
|
||||||
{
|
{
|
||||||
if (dataFormat == clFFT_InterleavedComplexFormat)
|
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
|
else
|
||||||
{
|
{
|
||||||
kernelString += string(" a[") + num2str(aIndex) + string("].x = in_real[") + num2str(gIndex) + string("];\n");
|
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
|
static void
|
||||||
formattedStore(string &kernelString, int aIndex, int gIndex, clFFT_DataFormat dataFormat)
|
formattedStore(string &kernelString, int aIndex, int gIndex, clFFT_DataFormat dataFormat)
|
||||||
{
|
{
|
||||||
if (dataFormat == clFFT_InterleavedComplexFormat)
|
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
|
else
|
||||||
{
|
{
|
||||||
kernelString += string(" out_real[") + num2str(gIndex) + string("] = a[") + num2str(aIndex) + string("].x;\n");
|
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
|
static int
|
||||||
insertGlobalLoadsAndTranspose(string &kernelString, int N, int numWorkItemsPerXForm, int numXFormsPerWG, int R0, int mem_coalesce_width, clFFT_DataFormat dataFormat)
|
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;
|
int lMemSize = 0;
|
||||||
|
|
||||||
if (numXFormsPerWG > 1)
|
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)
|
if (numWorkItemsPerXForm >= mem_coalesce_width)
|
||||||
{
|
{
|
||||||
@ -234,7 +249,9 @@ insertGlobalLoadsAndTranspose(string &kernelString, int N, int numWorkItemsPerXF
|
|||||||
kernelString += string(" out_imag += offset;\n");
|
kernelString += string(" out_imag += offset;\n");
|
||||||
}
|
}
|
||||||
for (i = 0; i < R0; i++)
|
for (i = 0; i < R0; i++)
|
||||||
formattedLoad(kernelString, i, i * numWorkItemsPerXForm, dataFormat);
|
{
|
||||||
|
formattedLoad(kernelString, i, i * numWorkItemsPerXForm, dataFormat);
|
||||||
|
}
|
||||||
kernelString += string(" }\n");
|
kernelString += string(" }\n");
|
||||||
}
|
}
|
||||||
else
|
else
|
||||||
@ -255,7 +272,9 @@ insertGlobalLoadsAndTranspose(string &kernelString, int N, int numWorkItemsPerXF
|
|||||||
kernelString += string(" out_imag += offset;\n");
|
kernelString += string(" out_imag += offset;\n");
|
||||||
}
|
}
|
||||||
for (i = 0; i < R0; i++)
|
for (i = 0; i < R0; i++)
|
||||||
formattedLoad(kernelString, i, i * numWorkItemsPerXForm, dataFormat);
|
{
|
||||||
|
formattedLoad(kernelString, i, i * numWorkItemsPerXForm, dataFormat);
|
||||||
|
}
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
else if (N >= mem_coalesce_width)
|
else if (N >= mem_coalesce_width)
|
||||||
@ -286,17 +305,23 @@ insertGlobalLoadsAndTranspose(string &kernelString, int N, int numWorkItemsPerXF
|
|||||||
{
|
{
|
||||||
kernelString += string(" if( jj < s ) {\n");
|
kernelString += string(" if( jj < s ) {\n");
|
||||||
for (j = 0; j < numInnerIter; j++)
|
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");
|
kernelString += string(" }\n");
|
||||||
if (i != numOuterIter - 1)
|
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("}\n ");
|
||||||
kernelString += string("else {\n");
|
kernelString += string("else {\n");
|
||||||
for (i = 0; i < numOuterIter; i++)
|
for (i = 0; i < numOuterIter; i++)
|
||||||
{
|
{
|
||||||
for (j = 0; j < numInnerIter; j++)
|
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");
|
kernelString += string("}\n");
|
||||||
|
|
||||||
@ -315,7 +340,9 @@ insertGlobalLoadsAndTranspose(string &kernelString, int N, int numWorkItemsPerXF
|
|||||||
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
|
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
|
||||||
|
|
||||||
for (i = 0; i < R0; i++)
|
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");
|
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
|
||||||
|
|
||||||
for (i = 0; i < numOuterIter; i++)
|
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");
|
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
|
||||||
|
|
||||||
for (i = 0; i < R0; i++)
|
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");
|
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
|
||||||
|
|
||||||
lMemSize = (N + numWorkItemsPerXForm) * numXFormsPerWG;
|
lMemSize = (N + numWorkItemsPerXForm) * numXFormsPerWG;
|
||||||
@ -360,7 +389,9 @@ insertGlobalLoadsAndTranspose(string &kernelString, int N, int numWorkItemsPerXF
|
|||||||
kernelString += string(" if(jj < s )\n");
|
kernelString += string(" if(jj < s )\n");
|
||||||
formattedLoad(kernelString, i, i * groupSize, dataFormat);
|
formattedLoad(kernelString, i, i * groupSize, dataFormat);
|
||||||
if (i != R0 - 1)
|
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("}\n");
|
||||||
kernelString += string("else {\n");
|
kernelString += string("else {\n");
|
||||||
@ -385,19 +416,27 @@ insertGlobalLoadsAndTranspose(string &kernelString, int N, int numWorkItemsPerXF
|
|||||||
|
|
||||||
|
|
||||||
for (i = 0; i < R0; i++)
|
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");
|
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
|
||||||
|
|
||||||
for (i = 0; i < R0; i++)
|
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");
|
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
|
||||||
|
|
||||||
for (i = 0; i < R0; i++)
|
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");
|
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
|
||||||
|
|
||||||
for (i = 0; i < R0; i++)
|
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");
|
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
|
||||||
|
|
||||||
lMemSize = (N + numWorkItemsPerXForm) * numXFormsPerWG;
|
lMemSize = (N + numWorkItemsPerXForm) * numXFormsPerWG;
|
||||||
@ -406,6 +445,7 @@ insertGlobalLoadsAndTranspose(string &kernelString, int N, int numWorkItemsPerXF
|
|||||||
return lMemSize;
|
return lMemSize;
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
static int
|
static int
|
||||||
insertGlobalStoresAndTranspose(string &kernelString, int N, int maxRadix, int Nr, int numWorkItemsPerXForm, int numXFormsPerWG, int mem_coalesce_width, clFFT_DataFormat dataFormat)
|
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);
|
formattedStore(kernelString, ind, i * numWorkItemsPerXForm, dataFormat);
|
||||||
}
|
}
|
||||||
if (numXFormsPerWG > 1)
|
if (numXFormsPerWG > 1)
|
||||||
kernelString += string(" }\n");
|
{
|
||||||
|
kernelString += string(" }\n");
|
||||||
|
}
|
||||||
}
|
}
|
||||||
else if (N >= mem_coalesce_width)
|
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");
|
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
|
||||||
|
|
||||||
for (i = 0; i < numOuterIter; i++)
|
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");
|
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
|
||||||
|
|
||||||
for (i = 0; i < maxRadix; i++)
|
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");
|
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
|
||||||
|
|
||||||
for (i = 0; i < numOuterIter; i++)
|
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(" barrier( CLK_LOCAL_MEM_FENCE );\n");
|
||||||
|
|
||||||
kernelString += string("if((groupId == get_num_groups(0)-1) && s) {\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");
|
kernelString += string(" if( jj < s ) {\n");
|
||||||
for (j = 0; j < numInnerIter; j++)
|
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");
|
kernelString += string(" }\n");
|
||||||
if (i != numOuterIter - 1)
|
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("}\n");
|
||||||
kernelString += string("else {\n");
|
kernelString += string("else {\n");
|
||||||
for (i = 0; i < numOuterIter; i++)
|
for (i = 0; i < numOuterIter; i++)
|
||||||
{
|
{
|
||||||
for (j = 0; j < numInnerIter; j++)
|
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");
|
kernelString += string("}\n");
|
||||||
|
|
||||||
@ -512,7 +568,9 @@ insertGlobalStoresAndTranspose(string &kernelString, int N, int maxRadix, int Nr
|
|||||||
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
|
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
|
||||||
|
|
||||||
for (i = 0; i < maxRadix; i++)
|
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");
|
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
|
||||||
|
|
||||||
for (i = 0; i < maxRadix; i++)
|
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");
|
kernelString += string(" barrier( CLK_LOCAL_MEM_FENCE );\n");
|
||||||
|
|
||||||
for (i = 0; i < maxRadix; i++)
|
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(" barrier( CLK_LOCAL_MEM_FENCE );\n");
|
||||||
|
|
||||||
kernelString += string("if((groupId == get_num_groups(0)-1) && s) {\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);
|
formattedStore(kernelString, i, i * groupSize, dataFormat);
|
||||||
kernelString += string(" }\n");
|
kernelString += string(" }\n");
|
||||||
if (i != maxRadix - 1)
|
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("}\n");
|
||||||
kernelString += string("else {\n");
|
kernelString += string("else {\n");
|
||||||
@ -551,6 +613,7 @@ insertGlobalStoresAndTranspose(string &kernelString, int N, int maxRadix, int Nr
|
|||||||
return lMemSize;
|
return lMemSize;
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
static void
|
static void
|
||||||
insertfftKernel(string &kernelString, int Nr, int numIter)
|
insertfftKernel(string &kernelString, int Nr, int numIter)
|
||||||
{
|
{
|
||||||
@ -561,6 +624,7 @@ insertfftKernel(string &kernelString, int Nr, int numIter)
|
|||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
static void
|
static void
|
||||||
insertTwiddleKernel(string &kernelString, int Nr, int numIter, int Nprev, int len, int numWorkItemsPerXForm)
|
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 (z == 0)
|
||||||
{
|
{
|
||||||
if (Nprev > 1)
|
if (Nprev > 1)
|
||||||
kernelString += string(" angf = (float) (ii >> ") + num2str(logNPrev) + string(");\n");
|
{
|
||||||
|
kernelString += string(" angf = (float) (ii >> ") + num2str(logNPrev) + string(");\n");
|
||||||
|
}
|
||||||
else
|
else
|
||||||
kernelString += string(" angf = (float) ii;\n");
|
{
|
||||||
|
kernelString += string(" angf = (float) ii;\n");
|
||||||
|
}
|
||||||
}
|
}
|
||||||
else
|
else
|
||||||
{
|
{
|
||||||
if (Nprev > 1)
|
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
|
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++)
|
for (k = 1; k < Nr; k++)
|
||||||
@ -596,30 +668,41 @@ insertTwiddleKernel(string &kernelString, int Nr, int numIter, int Nprev, int le
|
|||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
static int
|
static int
|
||||||
getPadding(int numWorkItemsPerXForm, int Nprev, int numWorkItemsReq, int numXFormsPerWG, int Nr, int numBanks, int *offset, int *midPad)
|
getPadding(int numWorkItemsPerXForm, int Nprev, int numWorkItemsReq, int numXFormsPerWG, int Nr, int numBanks, int *offset, int *midPad)
|
||||||
{
|
{
|
||||||
if ((numWorkItemsPerXForm <= Nprev) || (Nprev >= numBanks))
|
if ((numWorkItemsPerXForm <= Nprev) || (Nprev >= numBanks))
|
||||||
*offset = 0;
|
{
|
||||||
|
*offset = 0;
|
||||||
|
}
|
||||||
else
|
else
|
||||||
{
|
{
|
||||||
int numRowsReq = ((numWorkItemsPerXForm < numBanks) ? numWorkItemsPerXForm : numBanks) / Nprev;
|
int numRowsReq = ((numWorkItemsPerXForm < numBanks) ? numWorkItemsPerXForm : numBanks) / Nprev;
|
||||||
int numColsReq = 1;
|
int numColsReq = 1;
|
||||||
if (numRowsReq > Nr)
|
if (numRowsReq > Nr)
|
||||||
numColsReq = numRowsReq / Nr;
|
{
|
||||||
|
numColsReq = numRowsReq / Nr;
|
||||||
|
}
|
||||||
numColsReq = Nprev * numColsReq;
|
numColsReq = Nprev * numColsReq;
|
||||||
*offset = numColsReq;
|
*offset = numColsReq;
|
||||||
}
|
}
|
||||||
|
|
||||||
if (numWorkItemsPerXForm >= numBanks || numXFormsPerWG == 1)
|
if (numWorkItemsPerXForm >= numBanks || numXFormsPerWG == 1)
|
||||||
*midPad = 0;
|
{
|
||||||
|
*midPad = 0;
|
||||||
|
}
|
||||||
else
|
else
|
||||||
{
|
{
|
||||||
int bankNum = ((numWorkItemsReq + *offset) * Nr) & (numBanks - 1);
|
int bankNum = ((numWorkItemsReq + *offset) * Nr) & (numBanks - 1);
|
||||||
if (bankNum >= numWorkItemsPerXForm)
|
if (bankNum >= numWorkItemsPerXForm)
|
||||||
*midPad = 0;
|
{
|
||||||
|
*midPad = 0;
|
||||||
|
}
|
||||||
else
|
else
|
||||||
*midPad = numWorkItemsPerXForm - bankNum;
|
{
|
||||||
|
*midPad = numWorkItemsPerXForm - bankNum;
|
||||||
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
int lMemSize = (numWorkItemsReq + *offset) * Nr * numXFormsPerWG + *midPad * (numXFormsPerWG - 1);
|
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");
|
kernelString += string(" barrier(CLK_LOCAL_MEM_FENCE);\n");
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
static void
|
static void
|
||||||
insertLocalLoads(string &kernelString, int n, int Nr, int Nrn, int Nprev, int Ncurr, int numWorkItemsPerXForm, int numWorkItemsReq, int offset, string &comp)
|
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");
|
kernelString += string(" barrier(CLK_LOCAL_MEM_FENCE);\n");
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
static void
|
static void
|
||||||
insertLocalLoadIndexArithmatic(string &kernelString, int Nprev, int Nr, int numWorkItemsReq, int numWorkItemsPerXForm, int numXFormsPerWG, int offset, int midPad)
|
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 (Ncurr < numWorkItemsPerXForm)
|
||||||
{
|
{
|
||||||
if (Nprev == 1)
|
if (Nprev == 1)
|
||||||
kernelString += string(" j = ii & ") + num2str(Ncurr - 1) + string(";\n");
|
{
|
||||||
|
kernelString += string(" j = ii & ") + num2str(Ncurr - 1) + string(";\n");
|
||||||
|
}
|
||||||
else
|
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)
|
if (Nprev == 1)
|
||||||
kernelString += string(" i = ii >> ") + num2str(logNcurr) + string(";\n");
|
{
|
||||||
|
kernelString += string(" i = ii >> ") + num2str(logNcurr) + string(";\n");
|
||||||
|
}
|
||||||
else
|
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
|
else
|
||||||
{
|
{
|
||||||
if (Nprev == 1)
|
if (Nprev == 1)
|
||||||
kernelString += string(" j = ii;\n");
|
{
|
||||||
|
kernelString += string(" j = ii;\n");
|
||||||
|
}
|
||||||
else
|
else
|
||||||
kernelString += string(" j = ii >> ") + num2str(logNprev) + string(";\n");
|
{
|
||||||
|
kernelString += string(" j = ii >> ") + num2str(logNprev) + string(";\n");
|
||||||
|
}
|
||||||
if (Nprev == 1)
|
if (Nprev == 1)
|
||||||
kernelString += string(" i = 0;\n");
|
{
|
||||||
|
kernelString += string(" i = 0;\n");
|
||||||
|
}
|
||||||
else
|
else
|
||||||
kernelString += string(" i = ii & ") + num2str(Nprev - 1) + string(";\n");
|
{
|
||||||
|
kernelString += string(" i = ii & ") + num2str(Nprev - 1) + string(";\n");
|
||||||
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
if (numXFormsPerWG > 1)
|
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");
|
kernelString += string(" lMemLoad = sMem + mad24(j, ") + num2str(numWorkItemsReq + offset) + string(", i);\n");
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
static void
|
static void
|
||||||
insertLocalStoreIndexArithmatic(string &kernelString, int numWorkItemsReq, int numXFormsPerWG, int Nr, int offset, int midPad)
|
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");
|
assert(numRadix > 0 && "no radix array supplied\n");
|
||||||
|
|
||||||
if (n / radixArray[0] > plan->max_work_item_per_workgroup)
|
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(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");
|
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);
|
insertHeader(*kernelString, kernelName, dataFormat);
|
||||||
*kernelString += string("{\n");
|
*kernelString += string("{\n");
|
||||||
if ((*kInfo)->lmem_size)
|
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 += localString;
|
||||||
*kernelString += string("}\n");
|
*kernelString += string("}\n");
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
// For n larger than what can be computed using local memory fft, global transposes
|
// 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
|
// 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
|
// 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
|
// in this example. Users can play with difference base radices and difference
|
||||||
// decompositions of base radices to generates different kernels and see which gives
|
// 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
|
// 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)
|
void getGlobalRadixInfo(int n, int *radix, int *R1, int *R2, int *numRadices)
|
||||||
{
|
{
|
||||||
int baseRadix = min(n, 128);
|
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++)
|
for (int i = 0; i < numR; i++)
|
||||||
radix[i] = baseRadix;
|
{
|
||||||
|
radix[i] = baseRadix;
|
||||||
|
}
|
||||||
|
|
||||||
radix[numR] = N;
|
radix[numR] = N;
|
||||||
numR++;
|
numR++;
|
||||||
@ -906,6 +1016,7 @@ void getGlobalRadixInfo(int n, int *radix, int *R1, int *R2, int *numRadices)
|
|||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
static void
|
static void
|
||||||
createGlobalFFTKernelString(cl_fft_plan *plan, int n, int BS, cl_fft_kernel_dir dir, int vertBS)
|
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;
|
int strideI = Rinit;
|
||||||
for (i = 0; i < numPasses; i++)
|
for (i = 0; i < numPasses; i++)
|
||||||
if (i != passNum)
|
{
|
||||||
strideI *= radixArr[i];
|
if (i != passNum)
|
||||||
|
{
|
||||||
|
strideI *= radixArr[i];
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
int strideO = Rinit;
|
int strideO = Rinit;
|
||||||
for (i = 0; i < passNum; i++)
|
for (i = 0; i < passNum; i++)
|
||||||
strideO *= radixArr[i];
|
{
|
||||||
|
strideO *= radixArr[i];
|
||||||
|
}
|
||||||
|
|
||||||
int threadsPerXForm = R2;
|
int threadsPerXForm = R2;
|
||||||
batchSize = R2 == 1 ? plan->max_work_item_per_workgroup : batchSize;
|
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 numBlocksPerXForm = strideI / batchSize;
|
||||||
int numBlocks = numBlocksPerXForm;
|
int numBlocks = numBlocksPerXForm;
|
||||||
if (!vertical)
|
if (!vertical)
|
||||||
numBlocks *= BS;
|
{
|
||||||
|
numBlocks *= BS;
|
||||||
|
}
|
||||||
else
|
else
|
||||||
numBlocks *= vertBS;
|
{
|
||||||
|
numBlocks *= vertBS;
|
||||||
|
}
|
||||||
|
|
||||||
kernelName = string("fft") + num2str(kCount);
|
kernelName = string("fft") + num2str(kCount);
|
||||||
*kInfo = (cl_fft_kernel_info *)malloc(sizeof(cl_fft_kernel_info));
|
*kInfo = (cl_fft_kernel_info *)malloc(sizeof(cl_fft_kernel_info));
|
||||||
(*kInfo)->kernel = nullptr;
|
(*kInfo)->kernel = nullptr;
|
||||||
if (R2 == 1)
|
if (R2 == 1)
|
||||||
(*kInfo)->lmem_size = 0;
|
{
|
||||||
|
(*kInfo)->lmem_size = 0;
|
||||||
|
}
|
||||||
else
|
else
|
||||||
{
|
{
|
||||||
if (strideO == 1)
|
if (strideO == 1)
|
||||||
(*kInfo)->lmem_size = (radix + 1) * batchSize;
|
{
|
||||||
|
(*kInfo)->lmem_size = (radix + 1) * batchSize;
|
||||||
|
}
|
||||||
else
|
else
|
||||||
(*kInfo)->lmem_size = threadsPerBlock * R1;
|
{
|
||||||
|
(*kInfo)->lmem_size = threadsPerBlock * R1;
|
||||||
|
}
|
||||||
}
|
}
|
||||||
(*kInfo)->num_workgroups = numBlocks;
|
(*kInfo)->num_workgroups = numBlocks;
|
||||||
(*kInfo)->num_xforms_per_workgroup = 1;
|
(*kInfo)->num_xforms_per_workgroup = 1;
|
||||||
(*kInfo)->num_workitems_per_workgroup = threadsPerBlock;
|
(*kInfo)->num_workitems_per_workgroup = threadsPerBlock;
|
||||||
(*kInfo)->dir = dir;
|
(*kInfo)->dir = dir;
|
||||||
if ((passNum == (numPasses - 1)) && (numPasses & 1))
|
if ((passNum == (numPasses - 1)) && (numPasses & 1))
|
||||||
(*kInfo)->in_place_possible = 1;
|
{
|
||||||
|
(*kInfo)->in_place_possible = 1;
|
||||||
|
}
|
||||||
else
|
else
|
||||||
(*kInfo)->in_place_possible = 0;
|
{
|
||||||
|
(*kInfo)->in_place_possible = 0;
|
||||||
|
}
|
||||||
(*kInfo)->next = nullptr;
|
(*kInfo)->next = nullptr;
|
||||||
(*kInfo)->kernel_name = (char *)malloc(sizeof(char) * (kernelName.size() + 1));
|
(*kInfo)->kernel_name = (char *)malloc(sizeof(char) * (kernelName.size() + 1));
|
||||||
snprintf((*kInfo)->kernel_name, sizeof((*kInfo)->kernel_name), kernelName.c_str());
|
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");
|
localString += string("j = tid & ") + num2str(strideO - 1) + string(";\n");
|
||||||
int stride = radix * Rinit;
|
int stride = radix * Rinit;
|
||||||
for (i = 0; i < passNum; i++)
|
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("indexOut = mad24(i, ") + num2str(stride) + string(", j + ") + string("(xNum << ") + num2str((int)log2(n * BS)) + string("));\n");
|
||||||
localString += string("bNum = groupId;\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");
|
localString += string("j = tid & ") + num2str(strideO - 1) + string(";\n");
|
||||||
int stride = radix * Rinit;
|
int stride = radix * Rinit;
|
||||||
for (i = 0; i < passNum; i++)
|
for (i = 0; i < passNum; i++)
|
||||||
stride *= radixArr[i];
|
{
|
||||||
|
stride *= radixArr[i];
|
||||||
|
}
|
||||||
localString += string("indexOut = mad24(i, ") + num2str(stride) + string(", j);\n");
|
localString += string("indexOut = mad24(i, ") + num2str(stride) + string(", j);\n");
|
||||||
localString += string("indexIn += (xNum << ") + num2str(m) + string(");\n");
|
localString += string("indexIn += (xNum << ") + num2str(m) + string(");\n");
|
||||||
localString += string("indexOut += (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_real += indexIn;\n");
|
||||||
localString += string("in_imag += indexIn;\n");
|
localString += string("in_imag += indexIn;\n");
|
||||||
for (j = 0; j < R1; j++)
|
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++)
|
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
|
else
|
||||||
{
|
{
|
||||||
localString += string("in += indexIn;\n");
|
localString += string("in += indexIn;\n");
|
||||||
for (j = 0; j < R1; j++)
|
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");
|
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("lMemStore = sMem + tid;\n");
|
||||||
localString += string("lMemLoad = sMem + indexIn;\n");
|
localString += string("lMemLoad = sMem + indexIn;\n");
|
||||||
for (k = 0; k < R1; k++)
|
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");
|
localString += string("barrier(CLK_LOCAL_MEM_FENCE);\n");
|
||||||
for (k = 0; k < numIter; k++)
|
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");
|
localString += string("barrier(CLK_LOCAL_MEM_FENCE);\n");
|
||||||
for (k = 0; k < R1; k++)
|
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");
|
localString += string("barrier(CLK_LOCAL_MEM_FENCE);\n");
|
||||||
for (k = 0; k < numIter; k++)
|
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");
|
localString += string("barrier(CLK_LOCAL_MEM_FENCE);\n");
|
||||||
|
|
||||||
for (j = 0; j < numIter; j++)
|
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
|
// 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");
|
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 (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");
|
localString += string("barrier(CLK_LOCAL_MEM_FENCE);\n");
|
||||||
if (threadsPerBlock >= radix)
|
if (threadsPerBlock >= radix)
|
||||||
{
|
{
|
||||||
for (i = 0; i < R1; i++)
|
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
|
else
|
||||||
{
|
{
|
||||||
int innerIter = radix / threadsPerBlock;
|
int innerIter = radix / threadsPerBlock;
|
||||||
int outerIter = R1 / innerIter;
|
int outerIter = R1 / innerIter;
|
||||||
for (i = 0; i < outerIter; i++)
|
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");
|
localString += string("barrier(CLK_LOCAL_MEM_FENCE);\n");
|
||||||
|
|
||||||
for (i = 0; i < R1 / R2; i++)
|
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");
|
localString += string("barrier(CLK_LOCAL_MEM_FENCE);\n");
|
||||||
if (threadsPerBlock >= radix)
|
if (threadsPerBlock >= radix)
|
||||||
{
|
{
|
||||||
for (i = 0; i < R1; i++)
|
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
|
else
|
||||||
{
|
{
|
||||||
int innerIter = radix / threadsPerBlock;
|
int innerIter = radix / threadsPerBlock;
|
||||||
int outerIter = R1 / innerIter;
|
int outerIter = R1 / innerIter;
|
||||||
for (i = 0; i < outerIter; i++)
|
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");
|
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_real += indexOut;\n");
|
||||||
localString += string("out_imag += indexOut;\n");
|
localString += string("out_imag += indexOut;\n");
|
||||||
for (k = 0; k < R1; k++)
|
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++)
|
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
|
else
|
||||||
{
|
{
|
||||||
localString += string("out += indexOut;\n");
|
localString += string("out += indexOut;\n");
|
||||||
for (k = 0; k < R1; k++)
|
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
|
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_real += indexOut;\n");
|
||||||
localString += string("out_imag += indexOut;\n");
|
localString += string("out_imag += indexOut;\n");
|
||||||
for (k = 0; k < R1; k++)
|
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++)
|
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
|
else
|
||||||
{
|
{
|
||||||
localString += string("out += indexOut;\n");
|
localString += string("out += indexOut;\n");
|
||||||
for (k = 0; k < R1; k++)
|
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);
|
insertHeader(*kernelString, kernelName, dataFormat);
|
||||||
*kernelString += string("{\n");
|
*kernelString += string("{\n");
|
||||||
if ((*kInfo)->lmem_size)
|
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 += localString;
|
||||||
*kernelString += string("}\n");
|
*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)
|
void FFT1D(cl_fft_plan *plan, cl_fft_kernel_dir dir)
|
||||||
{
|
{
|
||||||
unsigned int radixArray[10];
|
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);
|
getRadixArray(plan->n.x, radixArray, &numRadix, plan->max_radix);
|
||||||
if (plan->n.x / radixArray[0] <= plan->max_work_item_per_workgroup)
|
if (plan->n.x / radixArray[0] <= plan->max_work_item_per_workgroup)
|
||||||
createLocalMemfftKernelString(plan);
|
{
|
||||||
|
createLocalMemfftKernelString(plan);
|
||||||
|
}
|
||||||
else
|
else
|
||||||
createGlobalFFTKernelString(plan, plan->n.x, 1, cl_fft_kernel_x, 1);
|
{
|
||||||
|
createGlobalFFTKernelString(plan, plan->n.x, 1, cl_fft_kernel_x, 1);
|
||||||
|
}
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
break;
|
break;
|
||||||
|
|
||||||
case cl_fft_kernel_y:
|
case cl_fft_kernel_y:
|
||||||
if (plan->n.y > 1)
|
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;
|
break;
|
||||||
|
|
||||||
case cl_fft_kernel_z:
|
case cl_fft_kernel_z:
|
||||||
if (plan->n.z > 1)
|
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:
|
default:
|
||||||
return;
|
return;
|
||||||
}
|
}
|
||||||
|
@ -31,9 +31,13 @@ getBlockConfigAndKernelString(cl_fft_plan *plan)
|
|||||||
*plan->kernel_string += baseKernels;
|
*plan->kernel_string += baseKernels;
|
||||||
|
|
||||||
if (plan->format == clFFT_SplitComplexFormat)
|
if (plan->format == clFFT_SplitComplexFormat)
|
||||||
*plan->kernel_string += twistKernelPlannar;
|
{
|
||||||
|
*plan->kernel_string += twistKernelPlannar;
|
||||||
|
}
|
||||||
else
|
else
|
||||||
*plan->kernel_string += twistKernelInterleaved;
|
{
|
||||||
|
*plan->kernel_string += twistKernelInterleaved;
|
||||||
|
}
|
||||||
|
|
||||||
switch (plan->dim)
|
switch (plan->dim)
|
||||||
{
|
{
|
||||||
@ -72,13 +76,18 @@ deleteKernelInfo(cl_fft_kernel_info *kInfo)
|
|||||||
if (kInfo)
|
if (kInfo)
|
||||||
{
|
{
|
||||||
if (kInfo->kernel_name)
|
if (kInfo->kernel_name)
|
||||||
free(kInfo->kernel_name);
|
{
|
||||||
|
free(kInfo->kernel_name);
|
||||||
|
}
|
||||||
if (kInfo->kernel)
|
if (kInfo->kernel)
|
||||||
clReleaseKernel(kInfo->kernel);
|
{
|
||||||
|
clReleaseKernel(kInfo->kernel);
|
||||||
|
}
|
||||||
free(kInfo);
|
free(kInfo);
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
static void
|
static void
|
||||||
destroy_plan(cl_fft_plan *Plan)
|
destroy_plan(cl_fft_plan *Plan)
|
||||||
{
|
{
|
||||||
@ -125,6 +134,7 @@ destroy_plan(cl_fft_plan *Plan)
|
|||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
static int
|
static int
|
||||||
createKernelList(cl_fft_plan *plan)
|
createKernelList(cl_fft_plan *plan)
|
||||||
{
|
{
|
||||||
@ -136,21 +146,30 @@ createKernelList(cl_fft_plan *plan)
|
|||||||
{
|
{
|
||||||
kernel_info->kernel = clCreateKernel(program, kernel_info->kernel_name, &err);
|
kernel_info->kernel = clCreateKernel(program, kernel_info->kernel_name, &err);
|
||||||
if (!kernel_info->kernel || err != CL_SUCCESS)
|
if (!kernel_info->kernel || err != CL_SUCCESS)
|
||||||
return err;
|
{
|
||||||
|
return err;
|
||||||
|
}
|
||||||
kernel_info = kernel_info->next;
|
kernel_info = kernel_info->next;
|
||||||
}
|
}
|
||||||
|
|
||||||
if (plan->format == clFFT_SplitComplexFormat)
|
if (plan->format == clFFT_SplitComplexFormat)
|
||||||
plan->twist_kernel = clCreateKernel(program, "clFFT_1DTwistSplit", &err);
|
{
|
||||||
|
plan->twist_kernel = clCreateKernel(program, "clFFT_1DTwistSplit", &err);
|
||||||
|
}
|
||||||
else
|
else
|
||||||
plan->twist_kernel = clCreateKernel(program, "clFFT_1DTwistInterleaved", &err);
|
{
|
||||||
|
plan->twist_kernel = clCreateKernel(program, "clFFT_1DTwistInterleaved", &err);
|
||||||
|
}
|
||||||
|
|
||||||
if (!plan->twist_kernel || err)
|
if (!plan->twist_kernel || err)
|
||||||
return err;
|
{
|
||||||
|
return err;
|
||||||
|
}
|
||||||
|
|
||||||
return CL_SUCCESS;
|
return CL_SUCCESS;
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
int getMaxKernelWorkGroupSize(cl_fft_plan *plan, unsigned int *max_wg_size, unsigned int num_devices, cl_device_id *devices)
|
int getMaxKernelWorkGroupSize(cl_fft_plan *plan, unsigned int *max_wg_size, unsigned int num_devices, cl_device_id *devices)
|
||||||
{
|
{
|
||||||
int reg_needed = 0;
|
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);
|
err = clGetKernelWorkGroupInfo(kInfo->kernel, devices[i], CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &wg_size, nullptr);
|
||||||
if (err != CL_SUCCESS)
|
if (err != CL_SUCCESS)
|
||||||
return -1;
|
{
|
||||||
|
return -1;
|
||||||
|
}
|
||||||
|
|
||||||
if (wg_size < kInfo->num_workitems_per_workgroup)
|
if (wg_size < kInfo->num_workitems_per_workgroup)
|
||||||
reg_needed |= 1;
|
{
|
||||||
|
reg_needed |= 1;
|
||||||
|
}
|
||||||
|
|
||||||
if (*max_wg_size > wg_size)
|
if (*max_wg_size > wg_size)
|
||||||
*max_wg_size = wg_size;
|
{
|
||||||
|
*max_wg_size = wg_size;
|
||||||
|
}
|
||||||
|
|
||||||
kInfo = kInfo->next;
|
kInfo = kInfo->next;
|
||||||
}
|
}
|
||||||
@ -181,6 +206,7 @@ int getMaxKernelWorkGroupSize(cl_fft_plan *plan, unsigned int *max_wg_size, unsi
|
|||||||
return reg_needed;
|
return reg_needed;
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
#define ERR_MACRO(err) \
|
#define ERR_MACRO(err) \
|
||||||
{ \
|
{ \
|
||||||
if ((err) != CL_SUCCESS) \
|
if ((err) != CL_SUCCESS) \
|
||||||
@ -192,6 +218,7 @@ int getMaxKernelWorkGroupSize(cl_fft_plan *plan, unsigned int *max_wg_size, unsi
|
|||||||
} \
|
} \
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
clFFT_Plan
|
clFFT_Plan
|
||||||
clFFT_CreatePlan(cl_context context, clFFT_Dim3 n, clFFT_Dimension dim, clFFT_DataFormat dataFormat, cl_int *error_code)
|
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)
|
if (error_code)
|
||||||
*error_code = CL_SUCCESS;
|
{
|
||||||
|
*error_code = CL_SUCCESS;
|
||||||
|
}
|
||||||
|
|
||||||
return (clFFT_Plan)plan;
|
return (clFFT_Plan)plan;
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
void clFFT_DestroyPlan(clFFT_Plan plan)
|
void clFFT_DestroyPlan(clFFT_Plan plan)
|
||||||
{
|
{
|
||||||
auto *Plan = (cl_fft_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)
|
void clFFT_DumpPlan(clFFT_Plan Plan, FILE *file)
|
||||||
{
|
{
|
||||||
size_t gDim;
|
size_t gDim;
|
||||||
size_t lDim;
|
size_t lDim;
|
||||||
FILE *out;
|
FILE *out;
|
||||||
if (!file)
|
if (!file)
|
||||||
out = stdout;
|
{
|
||||||
|
out = stdout;
|
||||||
|
}
|
||||||
else
|
else
|
||||||
out = file;
|
{
|
||||||
|
out = file;
|
||||||
|
}
|
||||||
|
|
||||||
auto *plan = (cl_fft_plan *)Plan;
|
auto *plan = (cl_fft_plan *)Plan;
|
||||||
cl_fft_kernel_info *kInfo = plan->kernel_info;
|
cl_fft_kernel_info *kInfo = plan->kernel_info;
|
||||||
|
@ -5,7 +5,7 @@
|
|||||||
# SPDX-License-Identifier: BSD-3-Clause
|
# 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>
|
# Usage: cmake -DCMAKE_TOOLCHAIN_FILE=<this file> <source directory>
|
||||||
########################################################################
|
########################################################################
|
||||||
set(CMAKE_CXX_COMPILER g++)
|
set(CMAKE_CXX_COMPILER g++)
|
||||||
|
@ -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
|
enable_ovf_check_buffer_monitor_active_ = false; // check buffer overflow and buffer monitor disabled by default
|
||||||
|
|
||||||
// some basic checks
|
// 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 << "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";
|
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;
|
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 << "Configuration parameter gain_mode_rx1 should take one of these values:\n";
|
||||||
std::cout << " manual, slow_attack, fast_attack, hybrid\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;
|
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 << "Configuration parameter gain_mode_rx2 should take one of these values:\n";
|
||||||
std::cout << " manual, slow_attack, fast_attack, hybrid\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 (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 << "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";
|
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 (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 << "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";
|
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 << "Configuration parameter filter_source should take one of these values:\n";
|
||||||
std::cout << " Off: Disable filter\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";
|
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 << "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";
|
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
|
// LOCAL OSCILLATOR DDS GENERATOR FOR DUAL FREQUENCY OPERATION
|
||||||
if (enable_dds_lo_ == true)
|
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 << "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";
|
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;
|
tx_bandwidth_ = 500000;
|
||||||
LOG(WARNING) << "Invalid configuration value for tx_bandwidth parameter. Set to 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 << "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";
|
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()
|
gr::basic_block_sptr Ad9361FpgaSignalSource::get_left_block()
|
||||||
{
|
{
|
||||||
LOG(WARNING) << "Trying to get signal source 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()
|
gr::basic_block_sptr Ad9361FpgaSignalSource::get_right_block()
|
||||||
{
|
{
|
||||||
return gr::basic_block_sptr();
|
return {};
|
||||||
}
|
}
|
||||||
|
@ -65,6 +65,14 @@ FileSourceBase::FileSourceBase(ConfigurationInterface const* configuration, std:
|
|||||||
{
|
{
|
||||||
filename_ = FLAGS_s;
|
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));
|
||||||
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
|
@ -91,7 +91,7 @@ Fmcomms2SignalSource::Fmcomms2SignalSource(const ConfigurationInterface *configu
|
|||||||
item_size_ = sizeof(gr_complex);
|
item_size_ = sizeof(gr_complex);
|
||||||
|
|
||||||
// some basic checks
|
// 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 << "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";
|
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";
|
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 << "Configuration parameter gain_mode_rx1 should take one of these values:\n";
|
||||||
std::cout << " manual, slow_attack, fast_attack, hybrid\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;
|
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 << "Configuration parameter gain_mode_rx2 should take one of these values:\n";
|
||||||
std::cout << " manual, slow_attack, fast_attack, hybrid\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;
|
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 << "Configuration parameter filter_source should take one of these values:\n";
|
||||||
std::cout << " Off: Disable filter\n";
|
std::cout << " Off: Disable filter\n";
|
||||||
@ -138,7 +138,7 @@ Fmcomms2SignalSource::Fmcomms2SignalSource(const ConfigurationInterface *configu
|
|||||||
|
|
||||||
if (gain_mode_rx1_ == "manual")
|
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 << "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";
|
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 (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 << "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";
|
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 << "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";
|
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 (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 !";
|
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
|
// configure LO
|
||||||
if (enable_dds_lo_ == true)
|
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 << "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";
|
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;
|
tx_bandwidth_ = 500000;
|
||||||
LOG(WARNING) << "Invalid configuration value for tx_bandwidth parameter. Set to 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 << "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";
|
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)
|
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 !";
|
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
|
// configure LO
|
||||||
if (enable_dds_lo_ == true)
|
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 << "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";
|
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;
|
tx_bandwidth_ = 500000;
|
||||||
LOG(WARNING) << "Invalid configuration value for tx_bandwidth parameter. Set to 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 << "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";
|
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()
|
gr::basic_block_sptr Fmcomms2SignalSource::get_left_block()
|
||||||
{
|
{
|
||||||
LOG(WARNING) << "Trying to get signal source left block.";
|
LOG(WARNING) << "Trying to get signal source left block.";
|
||||||
return gr::basic_block_sptr();
|
return {};
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
|
@ -231,7 +231,7 @@ void LimesdrSignalSource::disconnect(gr::top_block_sptr top_block)
|
|||||||
gr::basic_block_sptr LimesdrSignalSource::get_left_block()
|
gr::basic_block_sptr LimesdrSignalSource::get_left_block()
|
||||||
{
|
{
|
||||||
LOG(WARNING) << "Trying to get signal source left block.";
|
LOG(WARNING) << "Trying to get signal source left block.";
|
||||||
return gr::basic_block_sptr();
|
return {};
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
|
@ -210,7 +210,7 @@ void OsmosdrSignalSource::disconnect(gr::top_block_sptr top_block)
|
|||||||
gr::basic_block_sptr OsmosdrSignalSource::get_left_block()
|
gr::basic_block_sptr OsmosdrSignalSource::get_left_block()
|
||||||
{
|
{
|
||||||
LOG(WARNING) << "Trying to get signal source left block.";
|
LOG(WARNING) << "Trying to get signal source left block.";
|
||||||
return gr::basic_block_sptr();
|
return {};
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
|
@ -71,7 +71,7 @@ PlutosdrSignalSource::PlutosdrSignalSource(const ConfigurationInterface* configu
|
|||||||
}
|
}
|
||||||
|
|
||||||
// basic check
|
// 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 << "Configuration parameter gain_mode should take one of these values:\n";
|
||||||
std::cout << " manual, slow_attack, fast_attack, hybrid\n";
|
std::cout << " manual, slow_attack, fast_attack, hybrid\n";
|
||||||
@ -83,7 +83,7 @@ PlutosdrSignalSource::PlutosdrSignalSource(const ConfigurationInterface* configu
|
|||||||
|
|
||||||
if (gain_mode_ == "manual")
|
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 << "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";
|
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 << "Configuration parameter filter_source should take one of these values:\n";
|
||||||
std::cout << " Off: Disable filter\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";
|
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 << "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";
|
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()
|
gr::basic_block_sptr PlutosdrSignalSource::get_left_block()
|
||||||
{
|
{
|
||||||
LOG(WARNING) << "Trying to get signal source left block.";
|
LOG(WARNING) << "Trying to get signal source left block.";
|
||||||
return gr::basic_block_sptr();
|
return {};
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
|
@ -84,6 +84,7 @@ kf_vtl_tracking::kf_vtl_tracking(const Kf_Conf &conf_)
|
|||||||
d_trk_parameters(conf_),
|
d_trk_parameters(conf_),
|
||||||
d_acquisition_gnss_synchro(nullptr),
|
d_acquisition_gnss_synchro(nullptr),
|
||||||
d_signal_type(d_trk_parameters.signal),
|
d_signal_type(d_trk_parameters.signal),
|
||||||
|
d_code_chip_rate(0.0),
|
||||||
d_acq_code_phase_samples(0.0),
|
d_acq_code_phase_samples(0.0),
|
||||||
d_acq_carrier_doppler_hz(0.0),
|
d_acq_carrier_doppler_hz(0.0),
|
||||||
d_current_correlation_time_s(0.0),
|
d_current_correlation_time_s(0.0),
|
||||||
|
@ -1207,6 +1207,7 @@ void ControlThread::keyboard_listener()
|
|||||||
{
|
{
|
||||||
std::cout << "Quit keystroke order received, stopping GNSS-SDR !!\n";
|
std::cout << "Quit keystroke order received, stopping GNSS-SDR !!\n";
|
||||||
control_queue_->push(pmt::make_any(command_event_make(200, 0)));
|
control_queue_->push(pmt::make_any(command_event_make(200, 0)));
|
||||||
|
stop_ = true;
|
||||||
read_keys = false;
|
read_keys = false;
|
||||||
}
|
}
|
||||||
else
|
else
|
||||||
|
File diff suppressed because it is too large
Load Diff
@ -183,36 +183,18 @@ private:
|
|||||||
int connect_tracking_monitor();
|
int connect_tracking_monitor();
|
||||||
int connect_navdata_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
|
#if ENABLE_FPGA
|
||||||
int connect_fpga_flowgraph();
|
int connect_fpga_flowgraph();
|
||||||
int disconnect_fpga_flowgraph();
|
|
||||||
int connect_fpga_sample_counter();
|
int connect_fpga_sample_counter();
|
||||||
int disconnect_fpga_sample_counter();
|
|
||||||
#endif
|
#endif
|
||||||
|
|
||||||
void assign_channels();
|
int assign_channels();
|
||||||
void check_signal_conditioners();
|
void check_signal_conditioners();
|
||||||
|
|
||||||
void set_signals_list();
|
void set_signals_list();
|
||||||
void set_channels_state(); // Initializes the channels state (start acquisition or keep standby)
|
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)
|
// using the configuration parameters (number of channels and max channels in acquisition)
|
||||||
Gnss_Signal search_next_signal(const std::string& searched_signal,
|
Gnss_Signal search_next_signal(const std::string& searched_signal,
|
||||||
const bool pop,
|
|
||||||
bool& is_primary_frequency,
|
bool& is_primary_frequency,
|
||||||
bool& assistance_available,
|
bool& assistance_available,
|
||||||
float& estimated_doppler,
|
float& estimated_doppler,
|
||||||
|
@ -30,7 +30,7 @@ if(NOT GOOGLETEST_FOUND)
|
|||||||
"--build" "${CMAKE_BINARY_DIR}/gtest-${GNSSSDR_GTEST_LOCAL_VERSION}"
|
"--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>
|
"--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)
|
# Parallel building of gtest causes problems in some environments (e.g. Macports buildings)
|
||||||
set(GTEST_BUILD_COMMAND ${GTEST_BUILD_COMMAND} "--parallel 1")
|
set(GTEST_BUILD_COMMAND ${GTEST_BUILD_COMMAND} "--parallel 1")
|
||||||
endif()
|
endif()
|
||||||
|
@ -22,46 +22,46 @@
|
|||||||
|
|
||||||
|
|
||||||
Spirent_Motion_Csv_Dump_Reader::Spirent_Motion_Csv_Dump_Reader()
|
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;
|
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
|
@ -71,7 +71,7 @@ private:
|
|||||||
GpsL1CaPcpsAcquisitionTest_msg_rx();
|
GpsL1CaPcpsAcquisitionTest_msg_rx();
|
||||||
|
|
||||||
public:
|
public:
|
||||||
int rx_message;
|
int rx_message{0};
|
||||||
~GpsL1CaPcpsAcquisitionTest_msg_rx() override; //!< Default destructor
|
~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->message_port_register_in(pmt::mp("events"));
|
||||||
this->set_msg_handler(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));
|
boost::bind(&GpsL1CaPcpsAcquisitionTest_msg_rx::msg_handler_channel_events, this, _1));
|
||||||
#endif
|
#endif
|
||||||
#endif
|
#endif
|
||||||
rx_message = 0;
|
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
@ -123,12 +124,11 @@ class GpsL1CaPcpsAcquisitionTest : public ::testing::Test
|
|||||||
{
|
{
|
||||||
protected:
|
protected:
|
||||||
GpsL1CaPcpsAcquisitionTest()
|
GpsL1CaPcpsAcquisitionTest()
|
||||||
|
: item_size(sizeof(gr_complex))
|
||||||
|
|
||||||
{
|
{
|
||||||
config = std::make_shared<InMemoryConfiguration>();
|
config = std::make_shared<InMemoryConfiguration>();
|
||||||
item_size = sizeof(gr_complex);
|
|
||||||
gnss_synchro = Gnss_Synchro();
|
gnss_synchro = Gnss_Synchro();
|
||||||
doppler_max = 5000;
|
|
||||||
doppler_step = 100;
|
|
||||||
}
|
}
|
||||||
|
|
||||||
~GpsL1CaPcpsAcquisitionTest() override = default;
|
~GpsL1CaPcpsAcquisitionTest() override = default;
|
||||||
@ -140,8 +140,8 @@ protected:
|
|||||||
std::shared_ptr<InMemoryConfiguration> config;
|
std::shared_ptr<InMemoryConfiguration> config;
|
||||||
Gnss_Synchro gnss_synchro{};
|
Gnss_Synchro gnss_synchro{};
|
||||||
size_t item_size;
|
size_t item_size;
|
||||||
unsigned int doppler_max;
|
unsigned int doppler_max{5000};
|
||||||
unsigned int doppler_step;
|
unsigned int doppler_step{100};
|
||||||
};
|
};
|
||||||
|
|
||||||
|
|
||||||
|
@ -53,9 +53,9 @@ public:
|
|||||||
|
|
||||||
|
|
||||||
DataTypeAdapter::DataTypeAdapter()
|
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<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};
|
std::array<int16_t, 8> input_shorts{2, 23, -1, 127, -127, 0, 255, 255};
|
||||||
|
|
||||||
|
@ -44,10 +44,9 @@ DEFINE_int32(filter_test_nsamples, 1000000, "Number of samples to filter in the
|
|||||||
class FirFilterTest : public ::testing::Test
|
class FirFilterTest : public ::testing::Test
|
||||||
{
|
{
|
||||||
protected:
|
protected:
|
||||||
FirFilterTest()
|
FirFilterTest() : item_size(sizeof(gr_complex))
|
||||||
{
|
{
|
||||||
queue = std::make_shared<Concurrent_Queue<pmt::pmt_t>>();
|
queue = std::make_shared<Concurrent_Queue<pmt::pmt_t>>();
|
||||||
item_size = sizeof(gr_complex);
|
|
||||||
config = std::make_shared<InMemoryConfiguration>();
|
config = std::make_shared<InMemoryConfiguration>();
|
||||||
}
|
}
|
||||||
~FirFilterTest() override = default;
|
~FirFilterTest() override = default;
|
||||||
|
@ -43,12 +43,10 @@ DEFINE_int32(notch_filter_lite_test_nsamples, 1000000, "Number of samples to fil
|
|||||||
class NotchFilterLiteTest : public ::testing::Test
|
class NotchFilterLiteTest : public ::testing::Test
|
||||||
{
|
{
|
||||||
protected:
|
protected:
|
||||||
NotchFilterLiteTest()
|
NotchFilterLiteTest() : item_size(sizeof(gr_complex)), nsamples(FLAGS_notch_filter_lite_test_nsamples)
|
||||||
{
|
{
|
||||||
queue = std::make_shared<Concurrent_Queue<pmt::pmt_t>>();
|
queue = std::make_shared<Concurrent_Queue<pmt::pmt_t>>();
|
||||||
item_size = sizeof(gr_complex);
|
|
||||||
config = std::make_shared<InMemoryConfiguration>();
|
config = std::make_shared<InMemoryConfiguration>();
|
||||||
nsamples = FLAGS_notch_filter_lite_test_nsamples;
|
|
||||||
}
|
}
|
||||||
~NotchFilterLiteTest() override = default;
|
~NotchFilterLiteTest() override = default;
|
||||||
|
|
||||||
|
@ -43,12 +43,10 @@ DEFINE_int32(notch_filter_test_nsamples, 1000000, "Number of samples to filter i
|
|||||||
class NotchFilterTest : public ::testing::Test
|
class NotchFilterTest : public ::testing::Test
|
||||||
{
|
{
|
||||||
protected:
|
protected:
|
||||||
NotchFilterTest()
|
NotchFilterTest() : item_size(sizeof(gr_complex)), nsamples(FLAGS_notch_filter_test_nsamples)
|
||||||
{
|
{
|
||||||
queue = std::make_shared<Concurrent_Queue<pmt::pmt_t>>();
|
queue = std::make_shared<Concurrent_Queue<pmt::pmt_t>>();
|
||||||
item_size = sizeof(gr_complex);
|
|
||||||
config = std::make_shared<InMemoryConfiguration>();
|
config = std::make_shared<InMemoryConfiguration>();
|
||||||
nsamples = FLAGS_notch_filter_test_nsamples;
|
|
||||||
}
|
}
|
||||||
~NotchFilterTest() override = default;
|
~NotchFilterTest() override = default;
|
||||||
|
|
||||||
|
@ -43,12 +43,10 @@ DEFINE_int32(pb_filter_test_nsamples, 1000000, "Number of samples to filter in t
|
|||||||
class PulseBlankingFilterTest : public ::testing::Test
|
class PulseBlankingFilterTest : public ::testing::Test
|
||||||
{
|
{
|
||||||
protected:
|
protected:
|
||||||
PulseBlankingFilterTest()
|
PulseBlankingFilterTest() : item_size(sizeof(gr_complex)), nsamples(FLAGS_pb_filter_test_nsamples)
|
||||||
{
|
{
|
||||||
queue = std::make_shared<Concurrent_Queue<pmt::pmt_t>>();
|
queue = std::make_shared<Concurrent_Queue<pmt::pmt_t>>();
|
||||||
item_size = sizeof(gr_complex);
|
|
||||||
config = std::make_shared<InMemoryConfiguration>();
|
config = std::make_shared<InMemoryConfiguration>();
|
||||||
nsamples = FLAGS_pb_filter_test_nsamples;
|
|
||||||
}
|
}
|
||||||
~PulseBlankingFilterTest() override = default;
|
~PulseBlankingFilterTest() override = default;
|
||||||
bool stop = false;
|
bool stop = false;
|
||||||
|
@ -196,21 +196,12 @@ Acquisition_Dump_Reader::Acquisition_Dump_Reader(const std::string& basename,
|
|||||||
unsigned int samples_per_code,
|
unsigned int samples_per_code,
|
||||||
int channel,
|
int channel,
|
||||||
int execution)
|
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)
|
if (d_doppler_step == 0)
|
||||||
{
|
{
|
||||||
d_doppler_step = 1;
|
d_doppler_step = 1;
|
||||||
|
@ -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->message_port_register_in(pmt::mp("events"));
|
||||||
this->set_msg_handler(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));
|
boost::bind(&Acquisition_msg_rx::msg_handler_channel_events, this, _1));
|
||||||
#endif
|
#endif
|
||||||
#endif
|
#endif
|
||||||
rx_message = 0;
|
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
|
@ -104,8 +104,8 @@ void Observables_Dump_Reader::close_obs_file()
|
|||||||
|
|
||||||
|
|
||||||
Observables_Dump_Reader::Observables_Dump_Reader(int n_channels_)
|
Observables_Dump_Reader::Observables_Dump_Reader(int n_channels_)
|
||||||
|
: n_channels(n_channels_)
|
||||||
{
|
{
|
||||||
n_channels = n_channels_;
|
|
||||||
RX_time = std::vector<double>(n_channels);
|
RX_time = std::vector<double>(n_channels);
|
||||||
TOW_at_current_symbol_s = std::vector<double>(n_channels);
|
TOW_at_current_symbol_s = std::vector<double>(n_channels);
|
||||||
Carrier_Doppler_hz = std::vector<double>(n_channels);
|
Carrier_Doppler_hz = std::vector<double>(n_channels);
|
||||||
|
@ -25,7 +25,7 @@
|
|||||||
class TransitionModel : public ModelFunction
|
class TransitionModel : public ModelFunction
|
||||||
{
|
{
|
||||||
public:
|
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; };
|
arma::vec operator()(const arma::vec& input) override { return coeff_mat * input; };
|
||||||
|
|
||||||
private:
|
private:
|
||||||
@ -35,7 +35,7 @@ private:
|
|||||||
class MeasurementModel : public ModelFunction
|
class MeasurementModel : public ModelFunction
|
||||||
{
|
{
|
||||||
public:
|
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; };
|
arma::vec operator()(const arma::vec& input) override { return coeff_mat * input; };
|
||||||
|
|
||||||
private:
|
private:
|
||||||
|
@ -41,12 +41,10 @@ class GalileoE1DllPllVemlTrackingInternalTest : public ::testing::Test
|
|||||||
{
|
{
|
||||||
protected:
|
protected:
|
||||||
GalileoE1DllPllVemlTrackingInternalTest()
|
GalileoE1DllPllVemlTrackingInternalTest()
|
||||||
|
: item_size(sizeof(gr_complex))
|
||||||
{
|
{
|
||||||
factory = std::make_shared<GNSSBlockFactory>();
|
factory = std::make_shared<GNSSBlockFactory>();
|
||||||
config = std::make_shared<InMemoryConfiguration>();
|
config = std::make_shared<InMemoryConfiguration>();
|
||||||
item_size = sizeof(gr_complex);
|
|
||||||
stop = false;
|
|
||||||
message = 0;
|
|
||||||
gnss_synchro = Gnss_Synchro();
|
gnss_synchro = Gnss_Synchro();
|
||||||
}
|
}
|
||||||
|
|
||||||
@ -60,8 +58,8 @@ protected:
|
|||||||
std::shared_ptr<InMemoryConfiguration> config;
|
std::shared_ptr<InMemoryConfiguration> config;
|
||||||
Gnss_Synchro gnss_synchro{};
|
Gnss_Synchro gnss_synchro{};
|
||||||
size_t item_size;
|
size_t item_size;
|
||||||
bool stop;
|
bool stop{false};
|
||||||
int message;
|
int message{0};
|
||||||
};
|
};
|
||||||
|
|
||||||
|
|
||||||
|
@ -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->message_port_register_in(pmt::mp("events"));
|
||||||
this->set_msg_handler(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));
|
boost::bind(&GpsL1CADllPllTrackingTest_msg_rx::msg_handler_channel_events, this, _1));
|
||||||
#endif
|
#endif
|
||||||
#endif
|
#endif
|
||||||
rx_message = 0;
|
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
@ -166,11 +167,10 @@ public:
|
|||||||
double& rmse);
|
double& rmse);
|
||||||
|
|
||||||
bool save_mat_xy(std::vector<double>& x, std::vector<double>& y, std::string filename);
|
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>();
|
factory = std::make_shared<GNSSBlockFactory>();
|
||||||
config = std::make_shared<InMemoryConfiguration>();
|
config = std::make_shared<InMemoryConfiguration>();
|
||||||
item_size = sizeof(gr_complex);
|
|
||||||
gnss_synchro = Gnss_Synchro();
|
gnss_synchro = Gnss_Synchro();
|
||||||
}
|
}
|
||||||
|
|
||||||
|
@ -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",
|
GpsL1CADllPllTrackingTestFpga_msg_rx::GpsL1CADllPllTrackingTestFpga_msg_rx()
|
||||||
gr::io_signature::make(0, 0, 0),
|
: gr::block("GpsL1CADllPllTrackingTestFpga_msg_rx",
|
||||||
gr::io_signature::make(0, 0, 0))
|
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->message_port_register_in(pmt::mp("events"));
|
||||||
this->set_msg_handler(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));
|
boost::bind(&GpsL1CADllPllTrackingTestFpga_msg_rx::msg_handler_channel_events, this, _1));
|
||||||
#endif
|
#endif
|
||||||
#endif
|
#endif
|
||||||
rx_message = 0;
|
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
@ -245,10 +246,10 @@ public:
|
|||||||
void check_results_codephase(arma::vec &true_time_s, arma::vec &true_value,
|
void check_results_codephase(arma::vec &true_time_s, arma::vec &true_value,
|
||||||
arma::vec &meas_time_s, arma::vec &meas_value);
|
arma::vec &meas_time_s, arma::vec &meas_value);
|
||||||
|
|
||||||
GpsL1CADllPllTrackingTestFpga()
|
GpsL1CADllPllTrackingTestFpga() : item_size(sizeof(gr_complex))
|
||||||
{
|
{
|
||||||
config = std::make_shared<InMemoryConfiguration>();
|
config = std::make_shared<InMemoryConfiguration>();
|
||||||
item_size = sizeof(gr_complex);
|
|
||||||
gnss_synchro = Gnss_Synchro();
|
gnss_synchro = Gnss_Synchro();
|
||||||
}
|
}
|
||||||
|
|
||||||
|
@ -117,7 +117,7 @@ private:
|
|||||||
FrontEndCal_msg_rx();
|
FrontEndCal_msg_rx();
|
||||||
|
|
||||||
public:
|
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->message_port_register_in(pmt::mp("events"));
|
||||||
this->set_msg_handler(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));
|
boost::bind(&FrontEndCal_msg_rx::msg_handler_channel_events, this, _1));
|
||||||
#endif
|
#endif
|
||||||
#endif
|
#endif
|
||||||
rx_message = 0;
|
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
|
Loading…
Reference in New Issue
Block a user