Commit b80c8722 authored by Tobias Winchen's avatar Tobias Winchen
Browse files

Added possiibility to skip calculation of one gate.

parent a3bff557
Pipeline #121470 passed with stages
in 11 minutes and 8 seconds
......@@ -14,7 +14,7 @@
#include "cublas_v2.h"
#include <bitset>
#include <iostream>
#include <iomanip>
#include <cstring>
......@@ -295,7 +295,9 @@ struct GatedSpectrometerInputParameters
size_t fft_length;
size_t naccumulate;
unsigned int nbits;
// size_t disable_gate;
std::bitset<2> active_gates;
GatedSpectrometerInputParameters(): selectedSideChannel(0), speadHeapSize(0), nSideChannels(0), selectedBit(0), fft_length(0), naccumulate(0), nbits(0), active_gates(3) {}; // Default both gates active
};
......@@ -371,6 +373,7 @@ private:
std::size_t _selectedBit;
std::size_t _batch;
std::size_t _nsamps_per_heap;
std::bitset<2> _active_gates;
HandlerType &_handler;
cufftHandle _fft_plan;
......
......@@ -66,7 +66,7 @@ __global__ void update_baselines(float* __restrict__ baseLineG0,
template <class HandlerType, class InputType, class OutputType>
GatedSpectrometer<HandlerType, InputType, OutputType>::GatedSpectrometer(
const GatedSpectrometerInputParameters &ip, HandlerType &handler) : _dadaBufferLayout(ip.dadaBufferLayout),
_selectedSideChannel(ip.selectedSideChannel), _selectedBit(ip.selectedBit),
_selectedSideChannel(ip.selectedSideChannel), _selectedBit(ip.selectedBit), _active_gates(ip.active_gates),
_fft_length(ip.fft_length), _naccumulate(ip.naccumulate),
_handler(handler), _fft_plan(0), _call_count(0), _nsamps_per_heap(4096)
{
......@@ -263,23 +263,29 @@ void GatedSpectrometer<HandlerType, InputType, OutputType>::gated_fft(
_noOfBitSetsIn_G0.size()
);
BOOST_LOG_TRIVIAL(debug) << "Performing FFT 1";
BOOST_LOG_TRIVIAL(debug) << "Accessing unpacked voltage";
UnpackedVoltageType *_unpacked_voltage_ptr =
thrust::raw_pointer_cast(_unpacked_voltage_G0.data());
BOOST_LOG_TRIVIAL(debug) << "Accessing channelized voltage";
ChannelisedVoltageType *_channelised_voltage_ptr =
thrust::raw_pointer_cast(data._channelised_voltage_G0.data());
CUFFT_ERROR_CHECK(cufftExecR2C(_fft_plan, (cufftReal *)_unpacked_voltage_ptr,
(cufftComplex *)_channelised_voltage_ptr));
if (_active_gates[0])
{
BOOST_LOG_TRIVIAL(debug) << "Performing FFT 2";
_unpacked_voltage_ptr = thrust::raw_pointer_cast(_unpacked_voltage_G1.data());
_channelised_voltage_ptr = thrust::raw_pointer_cast(data._channelised_voltage_G1.data());
CUFFT_ERROR_CHECK(cufftExecR2C(_fft_plan, (cufftReal *)_unpacked_voltage_ptr,
(cufftComplex *)_channelised_voltage_ptr));
BOOST_LOG_TRIVIAL(debug) << "Performing FFT 1";
BOOST_LOG_TRIVIAL(debug) << "Accessing unpacked voltage";
UnpackedVoltageType *_unpacked_voltage_ptr = thrust::raw_pointer_cast(_unpacked_voltage_G0.data());
BOOST_LOG_TRIVIAL(debug) << "Accessing channelized voltage";
ChannelisedVoltageType *_channelised_voltage_ptr =
thrust::raw_pointer_cast(data._channelised_voltage_G0.data());
CUFFT_ERROR_CHECK(cufftExecR2C(_fft_plan, (cufftReal *)_unpacked_voltage_ptr,
(cufftComplex *)_channelised_voltage_ptr));
}
if (_active_gates[1])
{
BOOST_LOG_TRIVIAL(debug) << "Performing FFT 2";
BOOST_LOG_TRIVIAL(debug) << "Accessing unpacked voltage";
UnpackedVoltageType *_unpacked_voltage_ptr = thrust::raw_pointer_cast(_unpacked_voltage_G1.data());
BOOST_LOG_TRIVIAL(debug) << "Accessing channelized voltage";
ChannelisedVoltageType *_channelised_voltage_ptr = thrust::raw_pointer_cast(data._channelised_voltage_G1.data());
CUFFT_ERROR_CHECK(cufftExecR2C(_fft_plan, (cufftReal *)_unpacked_voltage_ptr,
(cufftComplex *)_channelised_voltage_ptr));
}
// CUDA_ERROR_CHECK(cudaStreamSynchronize(_proc_stream));
// BOOST_LOG_TRIVIAL(debug) << "Exit processing";
} // process
......@@ -364,21 +370,27 @@ void GatedSpectrometer<HandlerType, InputType, OutputType>::process(SinglePolari
{
gated_fft(*inputDataStream, outputDataStream->G0._noOfBitSets.a(), outputDataStream->G1._noOfBitSets.a());
kernels::detect_and_accumulate<IntegratedPowerType> <<<1024, 1024, 0, _proc_stream>>>(
thrust::raw_pointer_cast(inputDataStream->_channelised_voltage_G0.data()),
thrust::raw_pointer_cast(outputDataStream->G0.data.a().data()),
_nchans,
inputDataStream->_channelised_voltage_G0.size(),
_naccumulate / _nBlocks,
1, 0., 1, 0);
kernels::detect_and_accumulate<IntegratedPowerType> <<<1024, 1024, 0, _proc_stream>>>(
thrust::raw_pointer_cast(inputDataStream->_channelised_voltage_G1.data()),
thrust::raw_pointer_cast(outputDataStream->G1.data.a().data()),
_nchans,
inputDataStream->_channelised_voltage_G1.size(),
_naccumulate / _nBlocks,
1, 0., 1, 0);
if (_active_gates[0])
{
kernels::detect_and_accumulate<IntegratedPowerType> <<<1024, 1024, 0, _proc_stream>>>(
thrust::raw_pointer_cast(inputDataStream->_channelised_voltage_G0.data()),
thrust::raw_pointer_cast(outputDataStream->G0.data.a().data()),
_nchans,
inputDataStream->_channelised_voltage_G0.size(),
_naccumulate / _nBlocks,
1, 0., 1, 0);
}
if (_active_gates[1])
{
kernels::detect_and_accumulate<IntegratedPowerType> <<<1024, 1024, 0, _proc_stream>>>(
thrust::raw_pointer_cast(inputDataStream->_channelised_voltage_G1.data()),
thrust::raw_pointer_cast(outputDataStream->G1.data.a().data()),
_nchans,
inputDataStream->_channelised_voltage_G1.size(),
_naccumulate / _nBlocks,
1, 0., 1, 0);
}
// count saturated samples
for(size_t output_block_number = 0; output_block_number < outputDataStream->G0._noOfOverflowed.size(); output_block_number++)
......@@ -429,6 +441,9 @@ void GatedSpectrometer<HandlerType, InputType, OutputType>::process(DualPolariza
size_t input_offset = output_block_number * inputDataStream->polarization0._channelised_voltage_G0.size() / outputDataStream->G0._noOfBitSets.size();
size_t output_offset = output_block_number * outputDataStream->G0.I.a().size() / outputDataStream->G0._noOfBitSets.size();
BOOST_LOG_TRIVIAL(debug) << "Accumulating data for output block " << output_block_number << " with input offset " << input_offset << " and output_offset " << output_offset;
if (_active_gates[0])
{
stokes_accumulate<<<1024, 1024, 0, _proc_stream>>>(
thrust::raw_pointer_cast(inputDataStream->polarization0._channelised_voltage_G0.data() + input_offset),
thrust::raw_pointer_cast(inputDataStream->polarization1._channelised_voltage_G0.data() + input_offset),
......@@ -438,16 +453,19 @@ void GatedSpectrometer<HandlerType, InputType, OutputType>::process(DualPolariza
thrust::raw_pointer_cast(outputDataStream->G0.V.a().data() + output_offset),
_nchans, _naccumulate / _nBlocks
);
}
stokes_accumulate<<<1024, 1024, 0, _proc_stream>>>(
thrust::raw_pointer_cast(inputDataStream->polarization0._channelised_voltage_G1.data() + input_offset),
thrust::raw_pointer_cast(inputDataStream->polarization1._channelised_voltage_G1.data() + input_offset),
thrust::raw_pointer_cast(outputDataStream->G1.I.a().data() + output_offset),
thrust::raw_pointer_cast(outputDataStream->G1.Q.a().data() + output_offset),
thrust::raw_pointer_cast(outputDataStream->G1.U.a().data() + output_offset),
thrust::raw_pointer_cast(outputDataStream->G1.V.a().data() + output_offset),
_nchans, _naccumulate / _nBlocks
);
if (_active_gates[1]){
stokes_accumulate<<<1024, 1024, 0, _proc_stream>>>(
thrust::raw_pointer_cast(inputDataStream->polarization0._channelised_voltage_G1.data() + input_offset),
thrust::raw_pointer_cast(inputDataStream->polarization1._channelised_voltage_G1.data() + input_offset),
thrust::raw_pointer_cast(outputDataStream->G1.I.a().data() + output_offset),
thrust::raw_pointer_cast(outputDataStream->G1.Q.a().data() + output_offset),
thrust::raw_pointer_cast(outputDataStream->G1.U.a().data() + output_offset),
thrust::raw_pointer_cast(outputDataStream->G1.V.a().data() + output_offset),
_nchans, _naccumulate / _nBlocks
);
}
// count saturated samples
outputDataStream->G0._noOfOverflowed.a().data()[output_block_number] = 0;
......
......@@ -177,10 +177,11 @@ int main(int argc, char **argv) {
desc.add_options()("naccumulate,a",
po::value<size_t>(&ip.naccumulate)->required(),
"The number of samples to integrate in each channel");
// desc.add_options()("disable_gate,d",
// po::value<size_t>(&ip.disable_gate)->default_value(-1),
// "Disable processing of ND state 0,1. Select -1 (default) to process both.");
//
desc.add_options()("disable_gate,d",
po::value<uint8_t>()->notifier(
[&ip](size_t in) { ip.active_gates.set(in, false); }),
"Disable processing of ND state 0 or 1.");
desc.add_options()(
"log_level", po::value<std::string>()->default_value("info")->notifier(
[](std::string level) { set_log_level(level); }),
......
Supports Markdown
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