Commit 8b85f602 authored by Tobias Winchen's avatar Tobias Winchen
Browse files

Remove stream synchronization from methods to allow memcpy/kernel overlap

parent 020f3c10
......@@ -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
......
......@@ -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
......
Supports Markdown
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment