From 8b85f602f74b7993c6f45fc6e4bfb5bda319b6ed Mon Sep 17 00:00:00 2001
From: Tobias Winchen <tobias.winchen@rwth-aachen.de>
Date: Tue, 5 Mar 2019 13:14:40 +0000
Subject: [PATCH] Remove stream synchronization from methods to allow
memcpy/kernel overlap
---
psrdada_cpp/effelsberg/edd/src/DetectorAccumulator.cu | 1 -
psrdada_cpp/effelsberg/edd/src/Unpacker.cu | 2 --
2 files changed, 3 deletions(-)
diff --git a/psrdada_cpp/effelsberg/edd/src/DetectorAccumulator.cu b/psrdada_cpp/effelsberg/edd/src/DetectorAccumulator.cu
index e4c1b08f..527bfb95 100644
--- a/psrdada_cpp/effelsberg/edd/src/DetectorAccumulator.cu
+++ b/psrdada_cpp/effelsberg/edd/src/DetectorAccumulator.cu
@@ -57,7 +57,6 @@ void DetectorAccumulator::detect(InputType const& input, OutputType& output)
int8_t* output_ptr = thrust::raw_pointer_cast(output.data());
kernels::detect_and_accumulate<<<1024, 1024, 0, _stream>>>(
input_ptr, output_ptr, _nchans, nsamps, _tscrunch, _scale, _offset);
- CUDA_ERROR_CHECK(cudaStreamSynchronize(_stream));
}
} //namespace edd
diff --git a/psrdada_cpp/effelsberg/edd/src/Unpacker.cu b/psrdada_cpp/effelsberg/edd/src/Unpacker.cu
index 35065838..11c0ac7b 100644
--- a/psrdada_cpp/effelsberg/edd/src/Unpacker.cu
+++ b/psrdada_cpp/effelsberg/edd/src/Unpacker.cu
@@ -133,7 +133,6 @@ void Unpacker::unpack<12>(InputType const& input, OutputType& output)
OutputType::value_type* output_ptr = thrust::raw_pointer_cast(output.data());
kernels::unpack_edd_12bit_to_float32<<< nblocks, EDD_NTHREADS_UNPACK, 0, _stream>>>(
input_ptr, output_ptr, input.size());
- CUDA_ERROR_CHECK(cudaStreamSynchronize(_stream));
}
template <>
@@ -148,7 +147,6 @@ void Unpacker::unpack<8>(InputType const& input, OutputType& output)
OutputType::value_type* output_ptr = thrust::raw_pointer_cast(output.data());
kernels::unpack_edd_8bit_to_float32<<< nblocks, EDD_NTHREADS_UNPACK, 0, _stream>>>(
input_ptr, output_ptr, input.size());
- CUDA_ERROR_CHECK(cudaStreamSynchronize(_stream));
}
} //namespace edd
--
GitLab