Commit 48258b57 authored by Tobias Winchen's avatar Tobias Winchen
Browse files

Fix execution with multiple buffers per block + multiple blocks per output

parent 9ad5d074
......@@ -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))
......
......@@ -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);
}
......
......@@ -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);
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment