saymrwulf/onnxruntime
1
0
Fork 0
mirror of https://github.com/saymrwulf/onnxruntime.git synced 2026-07-30 20:18:08 +00:00
Code Issues Projects Releases Packages Wiki Activity Actions

Cuda pad optimize when no padding is needed. (#2625)

Browse source 
* Shortcut cuda Pad() when no padding is needed.
This commit is contained in:
Zhang Lei 2019-12-12 16:36:46 -08:00 committed by GitHub
parent 8631b70c73
commit e31be23c42
No known key found for this signature in database
GPG key ID: 4AEE18F83AFDEB23
1 changed files with 12 additions and 2 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

14
onnxruntime/core/providers/cuda/tensor/pad.cc
View file

@ -36,9 +36,9 @@ typename ToCudaType<T>::MappedType ToCudaValue(const T& value) {
return value;
}
template<>
template <>
typename ToCudaType<MLFloat16>::MappedType ToCudaValue<MLFloat16>(const MLFloat16& value) {
return *reinterpret_cast<const typename ToCudaType<MLFloat16>::MappedType *>(&value.val);
return *reinterpret_cast<const typename ToCudaType<MLFloat16>::MappedType*>(&value.val);
}
template <typename T>
@ -120,6 +120,16 @@ Status Pad<T>::ComputeInternal(OpKernelContext* ctx) const {
}
auto& output_tensor = *ctx->Output(0, output_shape);
if (std::all_of(p_pads->begin(), p_pads->end(), [](const int64_t v) { return v == 0; }) &&
std::all_of(p_slices->begin(), p_slices->end(), [](const int64_t v) { return v == 0; }) &&
output_shape.Size() > 0) {
CUDA_RETURN_IF_ERROR(cudaMemcpyAsync(
output_tensor.template MutableData<T>(), input_tensor.template Data<T>(),
sizeof(typename ToCudaType<T>::MappedType) * output_shape.Size(),
cudaMemcpyDeviceToDevice, 0));
return Status::OK();
}
ORT_ENFORCE(CalculateFdmStrides(fdm_output_strides.CpuSpan(), output_dims));
ORT_RETURN_IF_ERROR(input_dims.CopyToGpu());
ORT_RETURN_IF_ERROR(input_strides.CopyToGpu());