Fix out of bounds access in single pol

psrdada_cpp/effelsberg/edd/GatedSpectrometer.cuh

@@ -356,7 +356,7 @@ private:
   HandlerType &_handler;
   cufftHandle _fft_plan;
   uint64_t _nchans;
-  uint64_t _nBlocks;
+  uint64_t _nBlocks;            // Number of dada blocks to integrate into one spectrum
   uint64_t _call_count;
 
   std::unique_ptr<Unpacker> _unpacker;

psrdada_cpp/effelsberg/edd/detail/GatedSpectrometer.cu

@@ -82,7 +82,7 @@ GatedSpectrometer<HandlerType, InputType, OutputType>::GatedSpectrometer(
   CUDA_ERROR_CHECK(cudaDeviceSynchronize());
 
   BOOST_LOG_TRIVIAL(info)
-      << "Creating new GatedSpectrometer instance with parameters: \n"
+      << "Creating new GatedSpectrometer instance with parameters:\n"
       << "  fft_length           " << _fft_length << "\n"
       << "  naccumulate          " << _naccumulate << "\n"
       << "  nSideChannels        " << _dadaBufferLayout.getNSideChannels() << "\n"
@@ -104,29 +104,32 @@ GatedSpectrometer<HandlerType, InputType, OutputType>::GatedSpectrometer(
   size_t nsamps_per_output_spectra = fft_length * naccumulate;
 
   size_t nsamps_per_pol = inputDataStream->getSamplesPerInputPolarization();
+  BOOST_LOG_TRIVIAL(debug) << "Samples per input polarization: " << nsamps_per_pol;
   if (nsamps_per_output_spectra <= nsamps_per_pol)
-  { // one buffer block is used for one or multiple output spectra
+  { //
+    BOOST_LOG_TRIVIAL(debug) << "One buffer block is used for one or multiple output spectra";
     size_t N = nsamps_per_pol / nsamps_per_output_spectra;
     // All data in one block has to be used
     assert(N * nsamps_per_output_spectra == nsamps_per_pol);
     _nBlocks = 1;
   }
   else
-  { // multiple blocks are integrated into one output
+  {
+    BOOST_LOG_TRIVIAL(debug) << "Multiple blocks are integrated into one output spectrum";
     size_t N =  nsamps_per_output_spectra /  nsamps_per_pol;
     // All data in multiple blocks has to be used
     assert(N * nsamps_per_pol == nsamps_per_output_spectra);
     _nBlocks = N;
   }
   BOOST_LOG_TRIVIAL(debug) << "Integrating  " << nsamps_per_output_spectra <<
-      " samples from " << _nBlocks << "blocks into one output spectrum.";
+      " samples from " << _nBlocks << " blocks into one output spectrum.";
 
 
   // plan the FFT
   size_t nsamps_per_buffer = _dadaBufferLayout.sizeOfData() * 8 / nbits;
   int batch = nsamps_per_pol / _fft_length;
   int n[] = {static_cast<int>(_fft_length)};
-  BOOST_LOG_TRIVIAL(debug) << "Generating FFT plan: \n"
+  BOOST_LOG_TRIVIAL(debug) << "Generating FFT plan:\n"
       << "   fft_length = " << _fft_length << "\n"
       << "   n[0] = " << n[0] << "\n"
       << "   _nchans = " << _nchans << "\n"
@@ -363,6 +366,8 @@ template <class HandlerType, class InputType, class OutputType>
 void GatedSpectrometer<HandlerType, InputType, OutputType>::process(SinglePolarizationInput *inputDataStream, GatedPowerSpectrumOutput *outputDataStream)
 {
   gated_fft(*inputDataStream, outputDataStream->G0._noOfBitSets.a(), outputDataStream->G1._noOfBitSets.a());
+  thrust::fill(outputDataStream->G0.data.a().begin(), outputDataStream->G0.data.a().end(), 0);
+  thrust::fill(outputDataStream->G1.data.a().begin(), outputDataStream->G1.data.a().end(), 0);
 
   kernels::detect_and_accumulate<IntegratedPowerType> <<<1024, 1024, 0, _proc_stream>>>(
             thrust::raw_pointer_cast(inputDataStream->_channelised_voltage_G0.data()),
@@ -381,31 +386,35 @@ void GatedSpectrometer<HandlerType, InputType, OutputType>::process(SinglePolari
             1, 0., 1, 0);
 
     // count saturated samples
-    for(size_t output_block_number = 0; output_block_number <_nBlocks; output_block_number++)
+    for(size_t output_block_number = 0; output_block_number <outputDataStream->G0._noOfOverflowed.size(); output_block_number++)
     {
-        outputDataStream->G0._noOfOverflowed.a()[output_block_number] = 0;
-        outputDataStream->G1._noOfOverflowed.a()[output_block_number] = 0;
+        outputDataStream->G0._noOfOverflowed.a().data()[output_block_number] = 0;
+        outputDataStream->G1._noOfOverflowed.a().data()[output_block_number] = 0;
 
+        size_t lostHeaps = 0;
         const int heaps_per_output_spectra = inputDataStream->_sideChannelData_h.size() / _naccumulate / _nBlocks;
         for (size_t j = output_block_number * heaps_per_output_spectra ; j < (output_block_number+1) * heaps_per_output_spectra * _dadaBufferLayout.getNSideChannels(); j+=_dadaBufferLayout.getNSideChannels())
         {
             if (TEST_BIT(inputDataStream->_sideChannelData_h.a().data()[j], 3))
             { // heap was lost
+                lostHeaps++;
                 continue;
             }
 
             uint16_t n_saturated = (inputDataStream->_sideChannelData_h.a().data()[j] >> 32);
             if (TEST_BIT(inputDataStream->_sideChannelData_h.a().data()[j], _selectedBit))
             {
-                outputDataStream->G1._noOfOverflowed.a()[output_block_number] += n_saturated;
+                outputDataStream->G1._noOfOverflowed.a().data()[output_block_number] += n_saturated;
             }
             else
             {
-                outputDataStream->G0._noOfOverflowed.a()[output_block_number] += n_saturated;
+                outputDataStream->G0._noOfOverflowed.a().data()[output_block_number] += n_saturated;
             }
         }
-        BOOST_LOG_TRIVIAL(debug) << "Number of saturated samples G0: " << outputDataStream->G0._noOfOverflowed.a()[output_block_number]
-                    << ", G1: " << outputDataStream->G1._noOfOverflowed.a()[output_block_number];
+        BOOST_LOG_TRIVIAL(debug) << "Number of saturated samples G0: " << outputDataStream->G0._noOfOverflowed.a().data()[output_block_number]
+                    << ", G1: " << outputDataStream->G1._noOfOverflowed.a().data()[output_block_number];
+
+        BOOST_LOG_TRIVIAL(debug) << "Lost heaps: " << lostHeaps;
     }
 
 }
@@ -419,6 +428,15 @@ void GatedSpectrometer<HandlerType, InputType, OutputType>::process(DualPolariza
 
   gated_fft(inputDataStream->polarization0, outputDataStream->G0._noOfBitSets.a(), outputDataStream->G1._noOfBitSets.a());
   gated_fft(inputDataStream->polarization1, outputDataStream->G0._noOfBitSets.a(), outputDataStream->G1._noOfBitSets.a());
+  thrust::fill(outputDataStream->G0.I.a().begin(), outputDataStream->G0.I.a().end(), 0);
+  thrust::fill(outputDataStream->G0.Q.a().begin(), outputDataStream->G0.Q.a().end(), 0);
+  thrust::fill(outputDataStream->G0.U.a().begin(), outputDataStream->G0.U.a().end(), 0);
+  thrust::fill(outputDataStream->G0.V.a().begin(), outputDataStream->G0.V.a().end(), 0);
+
+  thrust::fill(outputDataStream->G1.I.a().begin(), outputDataStream->G1.I.a().end(), 0);
+  thrust::fill(outputDataStream->G1.Q.a().begin(), outputDataStream->G1.Q.a().end(), 0);
+  thrust::fill(outputDataStream->G1.U.a().begin(), outputDataStream->G1.U.a().end(), 0);
+  thrust::fill(outputDataStream->G1.V.a().begin(), outputDataStream->G1.V.a().end(), 0);
 
   for(size_t output_block_number = 0; output_block_number < outputDataStream->G0._noOfBitSets.size(); output_block_number++)
   {
@@ -445,9 +463,9 @@ void GatedSpectrometer<HandlerType, InputType, OutputType>::process(DualPolariza
               _nchans, _naccumulate / _nBlocks
               );
 
-      // count saturated samples
-        outputDataStream->G0._noOfOverflowed.a()[output_block_number] = 0;
-        outputDataStream->G1._noOfOverflowed.a()[output_block_number] = 0;
+        // count saturated samples
+        outputDataStream->G0._noOfOverflowed.a().data()[output_block_number] = 0;
+        outputDataStream->G1._noOfOverflowed.a().data()[output_block_number] = 0;
 
         size_t lostHeaps = 0;
 
@@ -466,17 +484,17 @@ void GatedSpectrometer<HandlerType, InputType, OutputType>::process(DualPolariza
 
             if (TEST_BIT(inputDataStream->polarization0._sideChannelData_h.a().data()[j], _selectedBit))
             {
-                outputDataStream->G1._noOfOverflowed.a()[output_block_number] += n_saturated;
+                outputDataStream->G1._noOfOverflowed.a().data()[output_block_number] += n_saturated;
             }
             else
             {
-                outputDataStream->G0._noOfOverflowed.a()[output_block_number] += n_saturated;
+                outputDataStream->G0._noOfOverflowed.a().data()[output_block_number] += n_saturated;
             }
 
         }
 
-        BOOST_LOG_TRIVIAL(debug) << "Number of saturated samples G0: " << outputDataStream->G0._noOfOverflowed.a()[output_block_number]
-                    << ", G1: " << outputDataStream->G1._noOfOverflowed.a()[output_block_number];
+        BOOST_LOG_TRIVIAL(debug) << "Number of saturated samples G0: " << outputDataStream->G0._noOfOverflowed.a().data()[output_block_number]
+                    << ", G1: " << outputDataStream->G1._noOfOverflowed.a().data()[output_block_number];
 
         BOOST_LOG_TRIVIAL(debug) << "Lost heaps: " << lostHeaps;
 

psrdada_cpp/effelsberg/edd/src/DadaBufferLayout.cpp

@@ -2,6 +2,7 @@
 
 #include "psrdada_cpp/multilog.hpp"
 #include "psrdada_cpp/dada_client_base.hpp"
+#include <iomanip>
 
 namespace psrdada_cpp {
 namespace effelsberg {
@@ -30,7 +31,7 @@ void DadaBufferLayout::intitialize(key_t input_key, size_t heapSize, size_t nSid
   _gapSize = (_bufferSize - _nHeaps * totalHeapSize);
   _dataBlockBytes = _nHeaps * _heapSize;
 
-  BOOST_LOG_TRIVIAL(debug) << "Memory configuration of dada buffer '" << _input_key << "' with " << nSideChannels << " side channels items and heapsize " << heapSize << " byte: \n"
+  BOOST_LOG_TRIVIAL(debug) << "Memory configuration of dada buffer '" << std::hex << _input_key << std::dec << "'\n with " << nSideChannels << " side channels items and heapsize " << heapSize << " byte: \n"
                            << "  total size of buffer: " << _bufferSize << " byte\n"
                            << "  number of heaps per buffer: " << _nHeaps << "\n"
                            << "  datablock size in buffer: " << _dataBlockBytes << " byte\n"

psrdada_cpp/effelsberg/edd/test/src/GatedSpectrometerTest.cu

@@ -493,6 +493,8 @@ TEST_P(ExecutionTests, SinglePolOutput)
     buffer.create();
 
     psrdada_cpp::effelsberg::edd::DadaBufferLayout bufferLayout(buffer.key(), heapSize, 1);
+    // Test buffer consistency with input parameters
+    EXPECT_EQ(bufferLayout.getNHeaps(), params.nHeaps);
 
     psrdada_cpp::NullSink sink;
 
@@ -531,6 +533,8 @@ TEST_P(ExecutionTests, FullStokesOutput)
     buffer.create();
 
     psrdada_cpp::effelsberg::edd::DadaBufferLayout bufferLayout(buffer.key(), heapSize, 1);
+    // Test buffer consistency with input parameters
+    EXPECT_EQ(bufferLayout.getNHeaps(), params.nHeaps);
 
     psrdada_cpp::NullSink sink;
 
@@ -561,14 +565,15 @@ TEST_P(ExecutionTests, FullStokesOutput)
 INSTANTIATE_TEST_CASE_P (GatedSpectrometer,
         ExecutionTests,
         testing::Values(
-          gated_params({2*1024, 2, 8, 1024, "1k Channel"}),
-          gated_params({2*1024, 2, 10, 1024, "1k Channel"}),
-          gated_params({2*1024, 2, 12, 1024, "1k Channel"}),
-          gated_params({2*1024 * 1024, 2, 8, 8*1024, "1M Channel, buffer larger than spectrum"}),
-          gated_params({2*1024 * 1024, 2, 10, 8*1024, "1M Channel, buffer larger than spectrum"}),
-          gated_params({2*1024 * 1024, 2, 12, 8*1024, "1M Channel, buffer larger tahn spectrum"}),
-
-          gated_params({2*1024 * 1024, 4096, 8, 1024, "1M Channel, integration larger than buffer"}),
+
+          // fft_length; naccumulate; nbits; nHeaps; msg;
+          gated_params({2*1024, 2,  8, 1024, "1k Channel:  8 Bit"}),
+          gated_params({2*1024, 2, 10, 1024, "1k Channel: 10 Bit"}),
+          gated_params({2*1024, 2, 12, 1024, "1k Channel: 12 Bit"}),
+          gated_params({2*1024 * 1024, 2,  8, 8*1024, "1M Channel, buffer 8x larger than spectrum; 8 bit"}),
+          gated_params({2*1024 * 1024, 2, 10, 8*1024, "1M Channel, buffer 8x larger than spectrum; 10-bit"}),
+          gated_params({2*1024 * 1024, 2, 12, 8*1024, "1M Channel, buffer 8x larger than spectrum; 12-bit"}),
+          gated_params({2*1024 * 1024, 2048, 8, 1024, "1M Channel, integration larger than buffer"}),
           gated_params({2*1024, 4096, 8, 1024, "1k Channel, integration larger than buffer"})
             )
         );