Skip to content
GitLab
Explore
Sign in
Primary navigation
Search or go to…
Project
LOFAR
Manage
Activity
Members
Labels
Plan
Issues
Wiki
Jira issues
Open Jira
Code
Merge requests
Repository
Branches
Commits
Tags
Repository graph
Compare revisions
Snippets
Locked files
Deploy
Releases
Package registry
Container Registry
Model registry
Operate
Environments
Terraform modules
Analyze
Value stream analytics
Contributor analytics
Repository analytics
Code review analytics
Insights
Model experiments
Help
Help
Support
GitLab documentation
Compare GitLab plans
Community forum
Contribute to GitLab
Provide feedback
Keyboard shortcuts
?
Snippets
Groups
Projects
Show more breadcrumbs
RadioObservatory
LOFAR
Commits
edf1405e
Commit
edf1405e
authored
4 years ago
by
Sarod Yatawatta
Browse files
Options
Downloads
Patches
Plain Diff
COB-53
: Only use delay gradient within a block for Doppler correction, the other terms are constant
parent
d5e1e945
No related branches found
No related tags found
1 merge request
!249
COB 153 - Doppler correction in correlator pipeline
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
RTCP/Cobalt/GPUProc/share/gpu/kernels/FIR_Filter.cu
+15
-11
15 additions, 11 deletions
RTCP/Cobalt/GPUProc/share/gpu/kernels/FIR_Filter.cu
with
15 additions
and
11 deletions
RTCP/Cobalt/GPUProc/share/gpu/kernels/FIR_Filter.cu
+
15
−
11
View file @
edf1405e
...
...
@@ -119,23 +119,19 @@ inline __device__ float sincos_f2f_select(float phi, int pol)
typedef
const
double
(
*
DelaysType
)[
1
][
NR_STABS
][
NR_POLARIZATIONS
];
// 2 Polarizations; in seconds
typedef
const
double
(
*
Phase0sType
)[
NR_STABS
][
NR_POLARIZATIONS
];
// 2 Polarizations; in radians
inline
__device__
float
correction_factor
(
unsigned
time
,
double
phi
AtBegin
,
double
phiAfterEnd
,
int
ri
,
int
pol
,
unsigned
channel
)
inline
__device__
float
correction_factor
(
unsigned
time
,
double
phi
Gradient
,
int
pol
,
unsigned
channel
)
{
// Offset of this sample between begin and end.
// Offset of this sample between begin and end.
= t/T fraction
const
double
blockOffset
=
double
(
time
*
NR_CHANNELS
+
channel
)
/
NR_SAMPLES_PER_SUBBAND
;
//const double blockOffset = double(channel * NR_SAMPLES_PER_CHANNEL + time) / (double)NR_SAMPLES_PER_SUBBAND;
// Interpolate the required phase rotation for this sample.
//
// Single precision angle + sincos is measured to be good enough (2013-11-20).
// Note that we do the interpolation in double precision still.
// We can afford to if we keep this kernel memory bound.
const
float
phi
=
phiAtBegin
*
(
1.0
-
blockOffset
)
+
phiAfterEnd
*
blockOffset
;
//if (phi>1e3) {
// printf("%d %d %e %e ",time,channel,blockOffset,phi);
//}
return
sincos_f2f_select
(
-
phi
,
pol
);
const
float
phi
=
phiGradient
*
blockOffset
;
return
sincos_f2f_select
(
phi
,
pol
);
}
#define FACTOR(time) correction_factor(time, phi
AtBegin, phiAfterEnd, ri
, pol, channel)
#define FACTOR(time) correction_factor(time, phi
Gradient
, pol, channel)
#endif
/* DOPPLER_CORRECTION */
/*!
...
...
@@ -221,8 +217,16 @@ __global__ void FIR_filter( void *filteredDataPtr,
//
// We need to undo the delay, so we rotate BACK, resulting in a negative constant factor.
// Note we use the sample frequency=CLOCK_MHZ/1024 (MHz) to normalize the frequency
const
double
phiAtBegin
=
-
2.0
*
M_PI
*
subbandFrequency
*
delayAtBegin
/
(
CLOCK_MHZ
*
1e6
/
1024.0
)
-
phase0
;
const
double
phiAfterEnd
=
-
2.0
*
M_PI
*
subbandFrequency
*
delayAfterEnd
/
(
CLOCK_MHZ
*
1e6
/
1024.0
)
-
phase0
;
// Let tau_0: delay at begin, tau_1: delay after end
// t : time within a block, T: total time (duration) of block,
// So, t in [0,T] and t/T=blockoffset in [0,1]
// then delay at t = tau_1 (t/T) + tau_0 (1 - t/T)
// exponent at frqeuency f = -j 2 pi (f tau)
// simplifying, exponent = -j 2 pi f tau_0 - j 2 pi f (tau_1 -tau_0) (t/T)
// We only need the term that changes with t, so discard the rest
// and only keep : 2 pi f (tau_1 - tau_0) (t/T) for the rest of calculations
// also replace f with f/f_s where f_s is sample frequency = clock/1024
const
double
phiGradient
=
2.0
*
M_PI
*
(
subbandFrequency
/
(
CLOCK_MHZ
*
1e6
/
1024.0
)
)
*
(
delayAfterEnd
-
delayAtBegin
);
#endif
//# const float16 weights = (*weightsData)[channel];
...
...
This diff is collapsed.
Click to expand it.
Preview
0%
Loading
Try again
or
attach a new file
.
Cancel
You are about to add
0
people
to the discussion. Proceed with caution.
Finish editing this message first!
Save comment
Cancel
Please
register
or
sign in
to comment