Task #4589: Added possibility to enable/disable bandpass correction in the new...

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.

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;

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;
     }
   }

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

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;
 }