diff --git a/psrdada_cpp/effelsberg/edd/DetectorAccumulator.cuh b/psrdada_cpp/effelsberg/edd/DetectorAccumulator.cuh index c95b3e0ca5dec13c062d9cf55350609bdb5ec62f..82e7d3cdf89136face34313de39b94907dae6343 100644 --- a/psrdada_cpp/effelsberg/edd/DetectorAccumulator.cuh +++ b/psrdada_cpp/effelsberg/edd/DetectorAccumulator.cuh @@ -30,7 +30,7 @@ void detect_and_accumulate(float2 const* __restrict__ in, int8_t* __restrict__ o double y = tmp.y * tmp.y; sum += x + y; } - size_t toff = out_offset * nchans + currentOutputSpectra * nchans; + size_t toff = out_offset * nchans + currentOutputSpectra * nchans *stride; out[toff + i] = (int8_t) ((sum - offset)/scale); } diff --git a/psrdada_cpp/effelsberg/edd/test/src/DetectorAccumulatorTester.cu b/psrdada_cpp/effelsberg/edd/test/src/DetectorAccumulatorTester.cu index d371f22daf319abfbf1ebd9d6ce9a7b67139ea7f..fdcacff20a361028cb758a81326e0bf79d33ba02 100644 --- a/psrdada_cpp/effelsberg/edd/test/src/DetectorAccumulatorTester.cu +++ b/psrdada_cpp/effelsberg/edd/test/src/DetectorAccumulatorTester.cu @@ -90,10 +90,12 @@ TEST_F(DetectorAccumulatorTester, noise_test) } DetectorAccumulator<int8_t>::InputType gpu_input = host_input; DetectorAccumulator<int8_t>::OutputType gpu_output; + gpu_output.resize(gpu_input.size() / tscrunch ); OutputType host_output; DetectorAccumulator<int8_t> detector(nchans, tscrunch, scale, 0.0, _stream); detector.detect(gpu_input, gpu_output); detect_c_reference(host_input, host_output, nchans, tscrunch, scale, 0.0); + CUDA_ERROR_CHECK(cudaStreamSynchronize(_stream)); compare_against_host(gpu_output, host_output); }