Added execution tests

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

+#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"})
+            )
+        );