GatedSpectrometer.cuh 6.57 KB
Newer Older
1
2
3
4
5
6
7
8
9
#ifndef PSRDADA_CPP_EFFELSBERG_EDD_GATEDSPECTROMETER_HPP
#define PSRDADA_CPP_EFFELSBERG_EDD_GATEDSPECTROMETER_HPP

#include "psrdada_cpp/effelsberg/edd/Unpacker.cuh"
#include "psrdada_cpp/raw_bytes.hpp"
#include "psrdada_cpp/cuda_utils.hpp"
#include "psrdada_cpp/double_device_buffer.cuh"
#include "psrdada_cpp/double_host_buffer.cuh"
#include "psrdada_cpp/effelsberg/edd/DetectorAccumulator.cuh"
10
#include "psrdada_cpp/effelsberg/edd/DadaBufferLayout.hpp"
11
12
13
14

#include "thrust/device_vector.h"
#include "cufft.h"

Tobias Winchen's avatar
Tobias Winchen committed
15
16
#include "cublas_v2.h"

17
18
19
20
21
22
23
24
25
26
namespace psrdada_cpp {
namespace effelsberg {
namespace edd {


#define BIT_MASK(bit) (1L << (bit))
#define SET_BIT(value, bit) ((value) |= BIT_MASK(bit))
#define CLEAR_BIT(value, bit) ((value) &= ~BIT_MASK(bit))
#define TEST_BIT(value, bit) (((value)&BIT_MASK(bit)) ? 1 : 0)

27
28
typedef unsigned long long int uint64_cu;
static_assert(sizeof(uint64_cu) == sizeof(uint64_t), "Long long int not of 64 bit! This is problematic for CUDA!");
29
30
31
32
33
34
/**
 @class GatedSpectrometer
 @brief Split data into two streams and create integrated spectra depending on
 bit set in side channel data.

 */
35
template <class HandlerType, typename IntegratedPowerType> class GatedSpectrometer {
36
37
38
39
public:
  typedef uint64_t RawVoltageType;
  typedef float UnpackedVoltageType;
  typedef float2 ChannelisedVoltageType;
40

41
//  typedef float IntegratedPowerType;
Tobias Winchen's avatar
Tobias Winchen committed
42
  //typedef int8_t IntegratedPowerType;
43
44
45
46
47
48
49
50
51

public:
  /**
   * @brief      Constructor
   *
   * @param      buffer_bytes A RawBytes object wrapping a DADA header buffer
   * @param      nSideChannels Number of side channel items in the data stream,
   * @param      selectedSideChannel Side channel item used for gating
   * @param      selectedBit bit of side channel item used for gating
52
53
54
55
56
57
58
59
60
   * @param      speadHeapSize Size of the spead heap block.
   * @param      fftLength Size of the FFT
   * @param      naccumulate Number of samples to integrate in the individual
   *             FFT bins
   * @param      nbits Bit depth of the sampled signal
   * @param      input_level Normalization level of the input signal
   * @param      output_level Normalization level of the output signal
   * @param      handler Output handler
   *
61
   */
62
  GatedSpectrometer(const DadaBufferLayout &bufferLayout,
63
                    std::size_t selectedSideChannel, std::size_t selectedBit,
64
                     std::size_t fft_length,
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
                    std::size_t naccumulate, std::size_t nbits,
                    float input_level, float output_level,
                    HandlerType &handler);
  ~GatedSpectrometer();

  /**
   * @brief      A callback to be called on connection
   *             to a ring buffer.
   *
   * @detail     The first available header block in the
   *             in the ring buffer is provided as an argument.
   *             It is here that header parameters could be read
   *             if desired.
   *
   * @param      block  A RawBytes object wrapping a DADA header buffer
   */
  void init(RawBytes &block);

  /**
   * @brief      A callback to be called on acqusition of a new
   *             data block.
   *
   * @param      block  A RawBytes object wrapping a DADA data buffer output
88
   *             are the integrated specttra with/without bit set.
89
90
91
92
93
   */
  bool operator()(RawBytes &block);

private:
  void process(thrust::device_vector<RawVoltageType> const &digitiser_raw,
94
               thrust::device_vector<uint64_t> const &sideChannelData,
95
               thrust::device_vector<IntegratedPowerType> &detected,
96
97
               thrust::device_vector<uint64_cu> &noOfBitSetsIn_G0,
               thrust::device_vector<uint64_cu> &noOfBitSetsIn_G1);
98
99

private:
100
  DadaBufferLayout _dadaBufferLayout;
101
102
103
104
105
  std::size_t _fft_length;
  std::size_t _naccumulate;
  std::size_t _nbits;
  std::size_t _selectedSideChannel;
  std::size_t _selectedBit;
106
  std::size_t _batch;
107
108
  std::size_t _nsamps_per_output_spectra;
  std::size_t _nsamps_per_buffer;
109
  std::size_t _nsamps_per_heap;
110
111
112

  HandlerType &_handler;
  cufftHandle _fft_plan;
Tobias Winchen's avatar
Tobias Winchen committed
113
114
  uint64_t _nchans;
  uint64_t _call_count;
115
116
  double _processing_efficiency;

117
  std::unique_ptr<Unpacker> _unpacker;
Tobias Winchen's avatar
Tobias Winchen committed
118
  std::unique_ptr<DetectorAccumulator<IntegratedPowerType> > _detector;
119

120
  // Input data
121
  DoubleDeviceBuffer<RawVoltageType> _raw_voltage_db;
122
  DoubleDeviceBuffer<uint64_t> _sideChannelData_db;
123
124
125

  // Output data
  DoubleDeviceBuffer<IntegratedPowerType> _power_db;
126
127
128
129


  DoubleDeviceBuffer<uint64_cu> _noOfBitSetsIn_G0;
  DoubleDeviceBuffer<uint64_cu> _noOfBitSetsIn_G1;
130
  DoublePinnedHostBuffer<char> _host_power_db;
131

132
  // Intermediate process steps
133
134
  thrust::device_vector<UnpackedVoltageType> _unpacked_voltage_G0;
  thrust::device_vector<UnpackedVoltageType> _unpacked_voltage_G1;
135
  thrust::device_vector<ChannelisedVoltageType> _channelised_voltage;
Tobias Winchen's avatar
Tobias Winchen committed
136
137
  thrust::device_vector<UnpackedVoltageType> _baseLineN;

138
139
140
141
142
143
  cudaStream_t _h2d_stream;
  cudaStream_t _proc_stream;
  cudaStream_t _d2h_stream;
};


144

145
146
147
148
149
/**
   * @brief      Splits the input data depending on a bit set into two arrays.
   *
   * @detail     The resulting gaps are filled with zeros in the other stream.
   *
Tobias Winchen's avatar
Tobias Winchen committed
150
   * @param      GO Input data. Data is set to the baseline value if corresponding
151
   *             sideChannelData bit at bitpos os set.
Tobias Winchen's avatar
Tobias Winchen committed
152
   * @param      G1 Data in this array is set to the baseline value if corresponding
153
154
155
156
157
158
159
160
161
162
163
   *             sideChannelData bit at bitpos is not set.
   * @param      sideChannelData noOfSideChannels items per block of heapSize
   *             bytes in the input data.
   * @param      N lebgth of the input/output arrays G0.
   * @param      heapsize Size of the blocks for which there is an entry in the
                 sideChannelData.
   * @param      bitpos Position of the bit to evaluate for processing.
   * @param      noOfSideChannels Number of side channels items per block of
   *             data.
   * @param      selectedSideChannel No. of side channel item to be eveluated.
                 0 <= selectedSideChannel < noOfSideChannels.
164
165
166
167
   * @param      stats_G0 No. of sampels contributing to G0, accounting also
   *             for loat heaps
   * @param      stats_G1 No. of sampels contributing to G1, accounting also
   *             for loat heaps
168
   */
169
__global__ void gating(float *G0, float *G1, const int64_t *sideChannelData,
170
                       size_t N, size_t heapSize, size_t bitpos,
171
172
173
                       size_t noOfSideChannels, size_t selectedSideChannel,
                       const float *_baseLine, uint64_cu* stats_G0, uint64_cu*
                       stats_G1);
Tobias Winchen's avatar
Tobias Winchen committed
174
175


176
177
178
179
180
181
182
183


} // edd
} // effelsberg
} // psrdada_cpp

#include "psrdada_cpp/effelsberg/edd/detail/GatedSpectrometer.cu"
#endif //PSRDADA_CPP_EFFELSBERG_EDD_GATEDSPECTROMETER_HPP