From ec0262770ab120acf038dd125d0b178cbf624d4f Mon Sep 17 00:00:00 2001
From: Marcel Loose <loose@astron.nl>
Date: Tue, 19 Nov 2013 14:30:08 +0000
Subject: [PATCH] Task #4589: Added possibility to enable/disable bandpass
correction in the new BandPassCorrection kernel. This is really useful for
testing. Note that this commit does not really belong to issue #4589, but
this was by far the easiest way to add it.
---
.../GPUProc/share/gpu/kernels/BandPassCorrection.cu | 13 ++++++++-----
.../src/cuda/Kernels/BandPassCorrectionKernel.cc | 7 ++++++-
.../src/cuda/Kernels/BandPassCorrectionKernel.h | 1 +
.../Cobalt/GPUProc/test/cuda/tBandPassCorrection.cc | 1 +
4 files changed, 16 insertions(+), 6 deletions(-)
diff --git a/RTCP/Cobalt/GPUProc/share/gpu/kernels/BandPassCorrection.cu b/RTCP/Cobalt/GPUProc/share/gpu/kernels/BandPassCorrection.cu
index 61a95f5aa26..5668b7423bc 100644
--- a/RTCP/Cobalt/GPUProc/share/gpu/kernels/BandPassCorrection.cu
+++ b/RTCP/Cobalt/GPUProc/share/gpu/kernels/BandPassCorrection.cu
@@ -92,9 +92,11 @@ __global__ void bandPassCorrection( fcomplex * outputDataPtr,
OutputDataType outputData = (OutputDataType) outputDataPtr;
InputDataType inputData = (InputDataType) inputDataPtr;
+#if defined DO_BANDPASS_CORRECTION
// Band pass to apply to the channels
BandPassFactorsType bandPassFactors = (BandPassFactorsType) bandPassFactorsPtr;
-
+#endif
+
// fasted dims
unsigned chan2 = blockIdx.x * blockDim.x + threadIdx.x;
unsigned sample = blockIdx.y * blockDim.y + threadIdx.y;
@@ -109,17 +111,18 @@ __global__ void bandPassCorrection( fcomplex * outputDataPtr,
for (unsigned idx_channel1 = 0; idx_channel1 < NR_CHANNELS_1; ++idx_channel1)
{
- // idx_channel1 steps with NR_CHANNELS_2 tru the channel weights
- float weight((*bandPassFactors)[idx_channel1 * NR_CHANNELS_2 + chan2]);
-
// Read from global memory in the quickest dimension (optimal)
fcomplex sampleX = (*inputData)[station][0][idx_channel1][sample][chan2];
fcomplex sampleY = (*inputData)[station][1][idx_channel1][sample][chan2];
-
+
+#if defined DO_BANDPASS_CORRECTION
+ // idx_channel1 steps with NR_CHANNELS_2 tru the channel weights
+ float weight((*bandPassFactors)[idx_channel1 * NR_CHANNELS_2 + chan2]);
sampleX.x *= weight;
sampleX.y *= weight;
sampleY.x *= weight;
sampleY.y *= weight;
+#endif
// Write the data to shared memory
tmp[threadIdx.y][threadIdx.x][0] = sampleX;
diff --git a/RTCP/Cobalt/GPUProc/src/cuda/Kernels/BandPassCorrectionKernel.cc b/RTCP/Cobalt/GPUProc/src/cuda/Kernels/BandPassCorrectionKernel.cc
index 7c707acd0c0..ee638a8ced6 100644
--- a/RTCP/Cobalt/GPUProc/src/cuda/Kernels/BandPassCorrectionKernel.cc
+++ b/RTCP/Cobalt/GPUProc/src/cuda/Kernels/BandPassCorrectionKernel.cc
@@ -48,7 +48,10 @@ namespace LOFAR
Kernel::Parameters(ps),
nrBitsPerSample(ps.settings.nrBitsPerSample),
nrBytesPerComplexSample(ps.nrBytesPerComplexSample()),
- nrSAPs(ps.settings.SAPs.size())
+ nrSAPs(ps.settings.SAPs.size()),
+ nrChannels1(64), // TODO: Must be read from parset?
+ nrChannels2(64), // TODO: Must be read from parset?
+ correctBandPass(ps.settings.corrections.bandPass)
{
dumpBuffers =
ps.getBool("Cobalt.Kernels.BandPassCorrectionKernel.dumpOutput", false);
@@ -125,6 +128,8 @@ namespace LOFAR
lexical_cast<string>(itsParameters.nrChannels1);
defs["NR_CHANNELS_2"] =
lexical_cast<string>(itsParameters.nrChannels2);
+ if (itsParameters.correctBandPass)
+ defs["DO_BANDPASS_CORRECTION"] = "1";
return defs;
}
}
diff --git a/RTCP/Cobalt/GPUProc/src/cuda/Kernels/BandPassCorrectionKernel.h b/RTCP/Cobalt/GPUProc/src/cuda/Kernels/BandPassCorrectionKernel.h
index 2e8dbf86b59..556ba756442 100644
--- a/RTCP/Cobalt/GPUProc/src/cuda/Kernels/BandPassCorrectionKernel.h
+++ b/RTCP/Cobalt/GPUProc/src/cuda/Kernels/BandPassCorrectionKernel.h
@@ -48,6 +48,7 @@ namespace LOFAR
size_t nrSAPs;
size_t nrChannels1;
size_t nrChannels2;
+ bool correctBandPass;
};
enum BufferType
diff --git a/RTCP/Cobalt/GPUProc/test/cuda/tBandPassCorrection.cc b/RTCP/Cobalt/GPUProc/test/cuda/tBandPassCorrection.cc
index 5625c13993b..126caca50e3 100644
--- a/RTCP/Cobalt/GPUProc/test/cuda/tBandPassCorrection.cc
+++ b/RTCP/Cobalt/GPUProc/test/cuda/tBandPassCorrection.cc
@@ -127,6 +127,7 @@ CompileDefinitions getDefaultCompileDefinitions()
boost::lexical_cast<string>(NR_POLARIZATIONS);
defs["NR_BITS_PER_SAMPLE"] =
boost::lexical_cast<string>(NR_BITS_PER_SAMPLE);
+ defs["DO_BANDPASS_CORRECTION"] = "1";
return defs;
}
--
GitLab