added channeliser tester

psrdada_cpp/effelsberg/edd/Channeliser.cuh

@@ -2,7 +2,7 @@
 #define PSRDADA_CPP_EFFELSBERG_EDD_CHANNELISER_HPP
 
 #include "psrdada_cpp/effelsberg/edd/Unpacker.cuh"
-#include "psrdada_cpp/effelsberg/edd/DetectorAccumulator.cuh"
+#include "psrdada_cpp/effelsberg/edd/ScaledTransposeTFtoTFT.cuh"
 #include "psrdada_cpp/raw_bytes.hpp"
 #include "psrdada_cpp/double_device_buffer.cuh"
 #include "psrdada_cpp/double_host_buffer.cuh"
@@ -66,7 +66,7 @@ private:
     int _nchans;
     int _call_count;
     std::unique_ptr<Unpacker> _unpacker;
-    std::unique_ptr<DetectorAccumulator> _detector;
+    std::unique_ptr<ScaledTransposeTFtoTFT> _transposer;
     DoubleDeviceBuffer<RawVoltageType> _raw_voltage_db;
     DoubleDeviceBuffer<PackedChannelisedVoltageType> _packed_channelised_voltage;
     thrust::device_vector<UnpackedVoltageType> _unpacked_voltage;

psrdada_cpp/effelsberg/edd/detail/Channeliser.cu

@@ -60,6 +60,7 @@ Channeliser<HandlerType>::Channeliser(
     CUDA_ERROR_CHECK(cudaStreamCreate(&_d2h_stream));
     CUFFT_ERROR_CHECK(cufftSetStream(_fft_plan, _proc_stream));
     _unpacker.reset(new Unpacker(_proc_stream));
+    _transposer.reset(new ScaledTransposeTFtoTFT(_nchans, 8192, scale, 0.0, _proc_stream));
 }
 
 template <class HandlerType>
@@ -99,9 +100,7 @@ void Channeliser<HandlerType>::process(
         (cufftReal*) _unpacked_voltage_ptr,
         (cufftComplex*) _channelised_voltage_ptr));
     CUDA_ERROR_CHECK(cudaStreamSynchronize(_proc_stream));
-    
-
-
+    _transposer.transpose(_channelised_voltage, _packed_channelised_voltage);
 }
 
 template <class HandlerType>
@@ -125,8 +124,8 @@ bool Channeliser<HandlerType>::operator()(RawBytes& block)
 
     // Synchronize all streams
     CUDA_ERROR_CHECK(cudaStreamSynchronize(_proc_stream));
-    _power_db.swap();
-    process(_raw_voltage_db.b(), _power_db.a());
+    _packed_channelised_voltage.swap();
+    process(_raw_voltage_db.b(), _packed_channelised_voltage.a());
 
     if (_call_count == 2)
     {
@@ -134,11 +133,11 @@ bool Channeliser<HandlerType>::operator()(RawBytes& block)
     }
 
     CUDA_ERROR_CHECK(cudaStreamSynchronize(_d2h_stream));
-    _host_power_db.swap();
+    _host_packed_channelised_voltage.swap();
     CUDA_ERROR_CHECK(cudaMemcpyAsync(
-        static_cast<void*>(_host_power_db.a_ptr()),
-        static_cast<void*>(_power_db.b_ptr()),
-        _power_db.size() * sizeof(IntegratedPowerType),
+        static_cast<void*>(_host_packed_channelised_voltage.a_ptr()),
+        static_cast<void*>(_packed_channelised_voltage.b_ptr()),
+        _packed_channelised_voltage.size() * sizeof(PackedChannelisedVoltageType),
         cudaMemcpyDeviceToHost,
         _d2h_stream));
     
@@ -148,9 +147,9 @@ bool Channeliser<HandlerType>::operator()(RawBytes& block)
     }   
 
     //Wrap _detected_host_previous in a RawBytes object here;
-    RawBytes bytes(reinterpret_cast<char*>(_host_power_db.b_ptr()),
-        _host_power_db.size() * sizeof(IntegratedPowerType),
-        _host_power_db.size() * sizeof(IntegratedPowerType));
+    RawBytes bytes(reinterpret_cast<char*>(_host_packed_channelised_voltage.b_ptr()),
+        _host_packed_channelised_voltage.size() * sizeof(PackedChannelisedVoltageType),
+        _host_packed_channelised_voltage.size() * sizeof(PackedChannelisedVoltageType));
     BOOST_LOG_TRIVIAL(debug) << "Calling handler";
     // The handler can't do anything asynchronously without a copy here 
     // as it would be unsafe (given that it does not own the memory it 

psrdada_cpp/effelsberg/edd/test/CMakeLists.txt

@@ -4,6 +4,7 @@ link_directories(${GTEST_LIBRARY_DIR})
 
 set(
     gtest_edd_src
+    src/ChanneliserTester.cu
     src/DetectorAccumulatorTester.cu
     src/FftSpectrometerTester.cu
     src/UnpackerTester.cu

psrdada_cpp/effelsberg/edd/test/ChanneliserTester.cuh

+#ifndef PSRDADA_CPP_EFFELSBERG_EDD_CHANNELISERTESTER_CUH
+#define PSRDADA_CPP_EFFELSBERG_EDD_CHANNELISERTESTER_CUH
+
+#include "psrdada_cpp/effelsberg/edd/Channeliser.cuh"
+#include "psrdada_cpp/dada_db.hpp"
+#include "thrust/host_vector.h"
+#include <gtest/gtest.h>
+#include <vector>
+
+namespace psrdada_cpp {
+namespace effelsberg {
+namespace edd {
+namespace test {
+
+class ChanneliserTester: public ::testing::Test
+{
+protected:
+    void SetUp() override;
+    void TearDown() override;
+
+public:
+    ChanneliserTester();
+    ~ChanneliserTester();
+    void performance_test(std::size_t nchans, std::size_t nsamps_per_packet, std::size_t nbits);
+};
+
+} //namespace test
+} //namespace edd
+} //namespace meerkat
+} //namespace psrdada_cpp
+
+#endif //PSRDADA_CPP_EFFELSBERG_EDD_CHANNELISERTESTER_CUH

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

+#include "psrdada_cpp/effelsberg/edd/test/ChanneliserTester.cuh"
+#include "psrdada_cpp/effelsberg/edd/Channeliser.cuh"
+#include "psrdada_cpp/dada_null_sink.hpp"
+#include "psrdada_cpp/double_host_buffer.cuh"
+#include "psrdada_cpp/raw_bytes.hpp"
+#include "psrdada_cpp/cuda_utils.hpp"
+#include <vector>
+
+namespace psrdada_cpp {
+namespace effelsberg {
+namespace edd {
+namespace test {
+
+ChanneliserTester::ChanneliserTester()
+    : ::testing::Test()
+{
+
+}
+
+ChanneliserTester::~ChanneliserTester()
+{
+
+}
+
+void ChanneliserTester::SetUp()
+{
+}
+
+void ChanneliserTester::TearDown()
+{
+}
+
+
+       std::size_t buffer_bytes,
+        std::size_t fft_length,
+        std::size_t nbits,
+        float input_level,
+        float output_level,
+        HandlerType& handler
+
+void ChanneliserTester::performance_test(std::size_t fft_length, std::size_t nbits)
+{
+    std::size_t input_block_bytes = fft_length * 8192 * 1024 * nbits / 8;
+   
+    DoublePinnedHostBuffer<char> input_block;
+    input_block.resize(input_block_bytes);	
+    RawBytes input_raw_bytes(input_block.a_ptr(), input_block_bytes, input_block_bytes);
+    std::vector<char> header_block(4096);
+    RawBytes header_raw_bytes(header_block.data(), 4096, 4096);
+    NullSink null_sink;
+    Channeliser<NullSink> channeliser(input_block_bytes, fft_length, nbits, 16.0f, 16.0f, null_sink);
+    spectrometer.init(header_raw_bytes);
+    for (int ii = 0; ii < 100; ++ii)
+    {
+        channeliser(input_raw_bytes);
+    }
+}
+
+
+TEST_F(ChanneliserTester, simple_exec_test)
+{
+    performance_test(1024, 16, 128, 12);
+}
+
+} //namespace test
+} //namespace edd
+} //namespace meerkat
+} //namespace psrdada_cpp