Detect accumulate fix

Fix detect accumulate kernel.

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

@@ -99,6 +99,112 @@ TEST_F(DetectorAccumulatorTester, noise_test)
     compare_against_host(gpu_output, host_output);
 }
 
+
+// parameters for the detector accumulator test
+struct detect_accumulate_params
+{
+    size_t nchan;       // umber of channels
+    size_t nspectra;    // number of output spectra
+    size_t naccumulate; // number of spectra to accumulate
+};
+
+
+class detect_and_accumulate32bit: public testing::TestWithParam<detect_accumulate_params> {
+// Test that the kernel executes in certain parameter
+// settings
+};
+
+
+TEST_P(detect_and_accumulate32bit, no_stride)
+{
+
+    thrust::device_vector<float2> input;
+    thrust::device_vector<float> output;
+
+    detect_accumulate_params params = GetParam();
+
+    input.resize(params.nspectra * params.nchan * params.naccumulate);
+    output.resize(params.nspectra * params.nchan);
+    float2 v;
+    v.x = 1.0;
+    v.y = 1.0;
+    thrust::fill(input.begin(), input.end(), v);
+    thrust::fill(output.begin(), output.end(), 0.);
+
+    kernels::detect_and_accumulate<float> <<<1024, 1024>>>(
+            thrust::raw_pointer_cast(input.data()),
+            thrust::raw_pointer_cast(output.data()),
+            params.nchan,
+            input.size(),
+            params.naccumulate,
+            1, 0., 1, 0);
+
+    cudaDeviceSynchronize();
+
+    thrust::host_vector<float> output_host = output;
+
+    for (size_t i =0; i < output.size(); i++)
+    {
+        ASSERT_FLOAT_EQ(output_host[i], (v.x * v.x + v.y * v.y) * params.naccumulate) << "i = " << i << " for nchan = " << params.nchan << ", nspectra = " << params.nspectra << ", naccumulate = " << params.naccumulate;
+    }
+}
+
+TEST_P(detect_and_accumulate32bit, w_stride)
+{
+
+    thrust::device_vector<float2> input;
+    thrust::device_vector<float> output;
+
+    detect_accumulate_params params = GetParam();
+
+    input.resize(params.nspectra * params.nchan * params.naccumulate);
+    output.resize(params.nspectra * params.nchan * 2);
+    float2 v;
+    v.x = 1.0;
+    v.y = 1.0;
+    thrust::fill(input.begin(), input.end(), v);
+    thrust::fill(output.begin(), output.end(), 0.);
+
+    kernels::detect_and_accumulate<float> <<<1024, 1024>>>(
+            thrust::raw_pointer_cast(input.data()),
+            thrust::raw_pointer_cast(output.data()),
+            params.nchan,
+            input.size(),
+            params.naccumulate,
+            1, 0., 2, 0);
+
+    cudaDeviceSynchronize();
+
+    thrust::host_vector<float> output_host = output;
+
+    for (size_t i =0; i < params.nchan * params.nspectra; i++)
+    {
+        size_t current_spectrum = i / params.nchan;
+        ASSERT_FLOAT_EQ(output_host[i + current_spectrum * params.nchan], (v.x * v.x + v.y * v.y) * params.naccumulate) << "i = " << i << " for nchan = " << params.nchan << ", nspectra = " << params.nspectra << ", naccumulate = " << params.naccumulate;
+    }
+}
+
+
+
+
+
+
+INSTANTIATE_TEST_CASE_P (DetectorAccumulatorTester,
+        detect_and_accumulate32bit,
+        testing::Values(
+          // nchan; nspectra; naccumulate;
+          detect_accumulate_params({1024, 1,  128}),
+          detect_accumulate_params({1024, 1,  1024}),
+          detect_accumulate_params({1024, 1,  2048}),
+          detect_accumulate_params({1024, 1,  16384}),
+          detect_accumulate_params({1024, 16,  1024}),
+          detect_accumulate_params({8388608, 2,  16}),
+          detect_accumulate_params({8388608, 1,  32})
+            )
+        );
+
+
+
 } //namespace test
 } //namespace edd
 } //namespace meerkat