diff --git a/psrdada_cpp/effelsberg/edd/detail/GatedSpectrometer.cu b/psrdada_cpp/effelsberg/edd/detail/GatedSpectrometer.cu index 4f35e680d922c23db94bb5c8fe125410ec8198a3..5e7ca6725e9be10ec49b242d9c9ceb15a70e8d2b 100644 --- a/psrdada_cpp/effelsberg/edd/detail/GatedSpectrometer.cu +++ b/psrdada_cpp/effelsberg/edd/detail/GatedSpectrometer.cu @@ -366,8 +366,6 @@ 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()), @@ -386,13 +384,13 @@ void GatedSpectrometer<HandlerType, InputType, OutputType>::process(SinglePolari 1, 0., 1, 0); // count saturated samples - for(size_t output_block_number = 0; output_block_number <outputDataStream->G0._noOfOverflowed.size(); output_block_number++) + for(size_t output_block_number = 0; output_block_number < outputDataStream->G0._noOfOverflowed.size(); output_block_number++) { 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; + const int heaps_per_output_spectra = inputDataStream->_sideChannelData_h.size() / outputDataStream->G0._noOfOverflowed.size(); 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)) @@ -428,15 +426,6 @@ 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++) { @@ -469,7 +458,7 @@ void GatedSpectrometer<HandlerType, InputType, OutputType>::process(DualPolariza size_t lostHeaps = 0; - const int heaps_per_output_spectra = inputDataStream->polarization0._sideChannelData_h.size() / _naccumulate / _nBlocks; + const int heaps_per_output_spectra = inputDataStream->polarization0._sideChannelData_h.size() / outputDataStream->G0._noOfBitSets.size(); 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->polarization0._sideChannelData_h.a().data()[j], 3) || TEST_BIT(inputDataStream->polarization1._sideChannelData_h.a().data()[j], 3)) diff --git a/psrdada_cpp/effelsberg/edd/src/GatedSpectrometer.cu b/psrdada_cpp/effelsberg/edd/src/GatedSpectrometer.cu index 09a88327d90296a47458b76a1985ea424548e7f9..c84cd8f7ac685492aeadfb13234b885176320606 100644 --- a/psrdada_cpp/effelsberg/edd/src/GatedSpectrometer.cu +++ b/psrdada_cpp/effelsberg/edd/src/GatedSpectrometer.cu @@ -350,6 +350,7 @@ void PowerSpectrumOutput::swap(cudaStream_t &_proc_stream) _noOfOverflowed.swap(); thrust::fill(thrust::cuda::par.on(_proc_stream), data.a().begin(), data.a().end(), 0.); thrust::fill(thrust::cuda::par.on(_proc_stream), _noOfBitSets.a().begin(), _noOfBitSets.a().end(), 0L); + thrust::fill(_noOfOverflowed.a().begin(), _noOfOverflowed.a().end(), 0L); } @@ -435,6 +436,8 @@ void FullStokesOutput::swap(cudaStream_t &_proc_stream) thrust::fill(thrust::cuda::par.on(_proc_stream), U.a().begin(), U.a().end(), 0.); thrust::fill(thrust::cuda::par.on(_proc_stream), V.a().begin(), V.a().end(), 0.); thrust::fill(thrust::cuda::par.on(_proc_stream), _noOfBitSets.a().begin(), _noOfBitSets.a().end(), 0L); + thrust::fill( _noOfOverflowed.a().begin(), _noOfOverflowed.a().end(), 0L); + } diff --git a/psrdada_cpp/effelsberg/edd/test/src/GatedSpectrometerTest.cu b/psrdada_cpp/effelsberg/edd/test/src/GatedSpectrometerTest.cu index a60a751c544c543c54c019802e29f621124fd305..52b6f0d5e4b700041fc2852cc04eb85b1d31b563 100644 --- a/psrdada_cpp/effelsberg/edd/test/src/GatedSpectrometerTest.cu +++ b/psrdada_cpp/effelsberg/edd/test/src/GatedSpectrometerTest.cu @@ -9,7 +9,6 @@ #include "thrust/device_vector.h" #include "thrust/extrema.h" - TEST(GatedSpectrometer, BitManipulationMacros) { for (int i = 0; i < 64; i++) { uint64_t v = 0; @@ -509,6 +508,7 @@ TEST_P(ExecutionTests, SinglePolOutput) sink); char *raw_buffer = new char[inputBufferSize]; + memset(raw_buffer, 0, inputBufferSize); psrdada_cpp::RawBytes buff(raw_buffer, inputBufferSize, inputBufferSize); @@ -549,6 +549,7 @@ TEST_P(ExecutionTests, FullStokesOutput) sink); char *raw_buffer = new char[inputBufferSize]; + memset(raw_buffer, 0, inputBufferSize); psrdada_cpp::RawBytes buff(raw_buffer, inputBufferSize, inputBufferSize);