diff --git a/psrdada_cpp/effelsberg/edd/test/src/GatedSpectrometerTest.cu b/psrdada_cpp/effelsberg/edd/test/src/GatedSpectrometerTest.cu
index 911cd27cd41abe8e5bdfba2359992f45921cf204..769393d14a880a3678b7502929af61ba5e654487 100644
--- a/psrdada_cpp/effelsberg/edd/test/src/GatedSpectrometerTest.cu
+++ b/psrdada_cpp/effelsberg/edd/test/src/GatedSpectrometerTest.cu
@@ -1,5 +1,7 @@
+#include "psrdada_cpp/effelsberg/edd/DadaBufferLayout.hpp"
#include "psrdada_cpp/effelsberg/edd/GatedSpectrometer.cuh"
+#include "psrdada_cpp/dada_db.hpp"
#include "psrdada_cpp/dada_null_sink.hpp"
#include "psrdada_cpp/multilog.hpp"
#include "gtest/gtest.h"
@@ -79,70 +81,72 @@ TEST(GatedSpectrometer, stokes_IQUV)
-void testParam(size_t nchans, size_t naccumulate){
-
- CUDA_ERROR_CHECK(cudaDeviceSynchronize());
-
- thrust::device_vector<float2> P0(nchans * naccumulate);
- thrust::device_vector<float2> P1(nchans * naccumulate);
- thrust::fill(P0.begin(), P0.end(), (float2){0, 0});
- thrust::fill(P1.begin(), P1.end(), (float2){0, 0});
- thrust::device_vector<float> I(nchans);
- thrust::device_vector<float> Q(nchans);
- thrust::device_vector<float> U(nchans);
- thrust::device_vector<float> V(nchans);
- thrust::fill(I.begin(), I.end(), 0);
- thrust::fill(Q.begin(), Q.end(), 0);
- thrust::fill(U.begin(), U.end(), 0);
- thrust::fill(V.begin(), V.end(), 0);
-
- // This channel should be left circular polarized
- size_t idx0 = 23;
- for (int k = 0; k< naccumulate; k++)
- {
- size_t idx = idx0 + k * nchans;
- P0[idx] = (float2){0.0f, 1.0f/std::sqrt(2)};
- P1[idx] = (float2){1.0f/std::sqrt(2),0.0f};
- }
-
- psrdada_cpp::effelsberg::edd::stokes_accumulate<<<1024, 1024>>>(
- thrust::raw_pointer_cast(P0.data()),
- thrust::raw_pointer_cast(P1.data()),
- thrust::raw_pointer_cast(I.data()),
- thrust::raw_pointer_cast(Q.data()),
- thrust::raw_pointer_cast(U.data()),
- thrust::raw_pointer_cast(V.data()),
- nchans,
- naccumulate
- );
-
- CUDA_ERROR_CHECK(cudaDeviceSynchronize());
- thrust::pair<thrust::device_vector<float>::iterator, thrust::device_vector<float>::iterator> minmax;
-
- minmax = thrust::minmax_element(I.begin(), I.end());
- EXPECT_FLOAT_EQ(*minmax.first, 0);
- EXPECT_FLOAT_EQ(*minmax.second, naccumulate);
-
- minmax = thrust::minmax_element(Q.begin(), Q.end());
- EXPECT_FLOAT_EQ(*minmax.first, 0);
- EXPECT_FLOAT_EQ(*minmax.second, 0);
-
- minmax = thrust::minmax_element(U.begin(), U.end());
- EXPECT_FLOAT_EQ(*minmax.first, 0);
- EXPECT_FLOAT_EQ(*minmax.second, 0);
-
- minmax = thrust::minmax_element(V.begin(), V.end());
- EXPECT_FLOAT_EQ(*minmax.first, -1. * naccumulate);
- EXPECT_FLOAT_EQ(*minmax.second, 0);
- };
+void testStokesAccumulateParam(size_t nchans, size_t naccumulate){
+ // Test the stokes accumulate kernel for different channel / accumualte
+ // numbers
+
+ CUDA_ERROR_CHECK(cudaDeviceSynchronize());
+
+ thrust::device_vector<float2> P0(nchans * naccumulate);
+ thrust::device_vector<float2> P1(nchans * naccumulate);
+ thrust::fill(P0.begin(), P0.end(), (float2){0, 0});
+ thrust::fill(P1.begin(), P1.end(), (float2){0, 0});
+ thrust::device_vector<float> I(nchans);
+ thrust::device_vector<float> Q(nchans);
+ thrust::device_vector<float> U(nchans);
+ thrust::device_vector<float> V(nchans);
+ thrust::fill(I.begin(), I.end(), 0);
+ thrust::fill(Q.begin(), Q.end(), 0);
+ thrust::fill(U.begin(), U.end(), 0);
+ thrust::fill(V.begin(), V.end(), 0);
+
+ // This channel should be left circular polarized
+ size_t idx0 = 23;
+ for (int k = 0; k< naccumulate; k++)
+ {
+ size_t idx = idx0 + k * nchans;
+ P0[idx] = (float2){0.0f, 1.0f/std::sqrt(2)};
+ P1[idx] = (float2){1.0f/std::sqrt(2),0.0f};
+ }
+ psrdada_cpp::effelsberg::edd::stokes_accumulate<<<1024, 1024>>>(
+ thrust::raw_pointer_cast(P0.data()),
+ thrust::raw_pointer_cast(P1.data()),
+ thrust::raw_pointer_cast(I.data()),
+ thrust::raw_pointer_cast(Q.data()),
+ thrust::raw_pointer_cast(U.data()),
+ thrust::raw_pointer_cast(V.data()),
+ nchans,
+ naccumulate
+ );
+
+ CUDA_ERROR_CHECK(cudaDeviceSynchronize());
+ thrust::pair<thrust::device_vector<float>::iterator, thrust::device_vector<float>::iterator> minmax;
-TEST(GatedSpectrometer, stokes_accumulate)
-{
+ minmax = thrust::minmax_element(I.begin(), I.end());
+ EXPECT_FLOAT_EQ(*minmax.first, 0);
+ EXPECT_FLOAT_EQ(*minmax.second, naccumulate);
+
+ minmax = thrust::minmax_element(Q.begin(), Q.end());
+ EXPECT_FLOAT_EQ(*minmax.first, 0);
+ EXPECT_FLOAT_EQ(*minmax.second, 0);
+
+ minmax = thrust::minmax_element(U.begin(), U.end());
+ EXPECT_FLOAT_EQ(*minmax.first, 0);
+ EXPECT_FLOAT_EQ(*minmax.second, 0);
+ minmax = thrust::minmax_element(V.begin(), V.end());
+ EXPECT_FLOAT_EQ(*minmax.first, -1. * naccumulate);
+ EXPECT_FLOAT_EQ(*minmax.second, 0);
+};
- testParam(8 * 1024 * 1024 + 1, 5);
- testParam(1024 + 1, 5);
+
+TEST(GatedSpectrometer, stokes_accumulate)
+{
+ testStokesAccumulateParam(8 * 1024 * 1024 + 1, 5);
+ testStokesAccumulateParam(8 * 1024 * 1024 + 1, 32);
+ testStokesAccumulateParam(1024 + 1, 5);
+ testStokesAccumulateParam(1024 + 1, 64);
}
@@ -270,3 +274,112 @@ TEST(GatedSpectrometer, array_sum) {
EXPECT_EQ(size_t(blr[0]), inputLength);
}
+
+
+struct gated_params
+{
+ std::size_t fft_length;
+ std::size_t naccumulate;
+ std::size_t nbits;
+ std::size_t nHeaps;
+ std::string msg;
+};
+
+
+class ExecutionTests: public testing::TestWithParam<gated_params> {
+// Test that the spectrometers execute without error in certain parameter
+// settings
+
+};
+
+
+TEST_P(ExecutionTests, SinglePolOutput)
+{
+ gated_params params = GetParam();
+
+ const size_t heapSize = 4096 * params.nbits / 8;
+ const size_t inputBufferSize = params.nHeaps * (heapSize + 64 / 8);
+
+ psrdada_cpp::DadaDB buffer(5, inputBufferSize, 1, 4096);
+ buffer.create();
+
+ psrdada_cpp::effelsberg::edd::DadaBufferLayout bufferLayout(buffer.key(), heapSize, 1);
+
+ psrdada_cpp::NullSink sink;
+
+ psrdada_cpp::effelsberg::edd::GatedSpectrometer<
+ psrdada_cpp::NullSink,
+ psrdada_cpp::effelsberg::edd::SinglePolarizationInput,
+ psrdada_cpp::effelsberg::edd::GatedPowerSpectrumOutput>
+ spectrometer(bufferLayout,
+ 0, 0,
+ params.fft_length,
+ params.naccumulate, params.nbits,
+ sink);
+
+ char *raw_buffer = new char[inputBufferSize];
+
+ psrdada_cpp::RawBytes buff(raw_buffer, inputBufferSize);
+
+ EXPECT_NO_THROW(spectrometer.init(buff)) << params.msg;
+ for (int i = 0; i < 5; i++)
+ {
+ EXPECT_NO_THROW(spectrometer(buff)) << params.msg;
+ }
+
+ delete [] raw_buffer;
+}
+
+
+TEST_P(ExecutionTests, FullStokesOutput)
+{
+ gated_params params = GetParam();
+
+ const size_t heapSize = 4096 * params.nbits / 8;
+ const size_t inputBufferSize = params.nHeaps * (heapSize + 64 / 8);
+
+ psrdada_cpp::DadaDB buffer(5, inputBufferSize, 1, 4096);
+ buffer.create();
+
+ psrdada_cpp::effelsberg::edd::DadaBufferLayout bufferLayout(buffer.key(), heapSize, 1);
+
+ psrdada_cpp::NullSink sink;
+
+ psrdada_cpp::effelsberg::edd::GatedSpectrometer<
+ psrdada_cpp::NullSink,
+ psrdada_cpp::effelsberg::edd::DualPolarizationInput,
+ psrdada_cpp::effelsberg::edd::GatedFullStokesOutput>
+ spectrometer(bufferLayout,
+ 0, 0,
+ params.fft_length,
+ params.naccumulate, params.nbits,
+ sink);
+
+ char *raw_buffer = new char[inputBufferSize];
+
+ psrdada_cpp::RawBytes buff(raw_buffer, inputBufferSize);
+
+ EXPECT_NO_THROW(spectrometer.init(buff)) << params.msg;
+ for (int i = 0; i < 5; i++)
+ {
+ EXPECT_NO_THROW(spectrometer(buff)) << params.msg;
+ }
+
+ delete [] raw_buffer;
+}
+
+
+INSTANTIATE_TEST_CASE_P (GatedSpectrometer,
+ ExecutionTests,
+ testing::Values(
+ gated_params({2*1024, 2, 8, 1024, "1k Channel"}),
+ gated_params({2*1024, 2, 10, 1024, "1k Channel"}),
+ gated_params({2*1024, 2, 12, 1024, "1k Channel"}),
+ gated_params({2*1024 * 1024, 2, 8, 8*1024, "1M Channel, buffer larger than spectrum"}),
+ gated_params({2*1024 * 1024, 2, 10, 8*1024, "1M Channel, buffer larger than spectrum"}),
+ gated_params({2*1024 * 1024, 2, 12, 8*1024, "1M Channel, buffer larger tahn spectrum"}),
+
+ gated_params({2*1024 * 1024, 4096, 8, 1024, "1M Channel, integration larger than buffer"}),
+ gated_params({2*1024, 4096, 8, 1024, "1k Channel, integration larger than buffer"})
+ )
+ );