diff --git a/onnxruntime/contrib_ops/cuda/math/isfinite.cc b/onnxruntime/contrib_ops/cuda/math/isfinite.cc index 0ebe21b23d..cd8b9c9555 100644 --- a/onnxruntime/contrib_ops/cuda/math/isfinite.cc +++ b/onnxruntime/contrib_ops/cuda/math/isfinite.cc @@ -2,6 +2,7 @@ // Licensed under the MIT License. #include "contrib_ops/cuda/math/isfinite.h" +#include "isfinite_impl.h" using namespace ONNX_NAMESPACE; using namespace onnxruntime::common; diff --git a/onnxruntime/contrib_ops/cuda/math/isfinite.cuh b/onnxruntime/contrib_ops/cuda/math/isfinite.cuh index d27a6631c1..847addb6ad 100644 --- a/onnxruntime/contrib_ops/cuda/math/isfinite.cuh +++ b/onnxruntime/contrib_ops/cuda/math/isfinite.cuh @@ -1,6 +1,8 @@ // Copyright (c) Microsoft Corporation. All rights reserved. // Licensed under the MIT License. +#pragma once + #include #include "core/providers/cuda/cu_inc/common.cuh" #include "contrib_ops/cuda/math/isfinite.h" diff --git a/onnxruntime/contrib_ops/cuda/math/isfinite.h b/onnxruntime/contrib_ops/cuda/math/isfinite.h index f33de16d2a..6233e7cb90 100644 --- a/onnxruntime/contrib_ops/cuda/math/isfinite.h +++ b/onnxruntime/contrib_ops/cuda/math/isfinite.h @@ -4,7 +4,6 @@ #pragma once #include "core/common/common.h" #include "core/providers/cuda/cuda_kernel.h" -#include "core/providers/cuda/multi_tensor/common.cuh" namespace onnxruntime { namespace cuda { @@ -31,10 +30,5 @@ class IsAllFiniteOp final : public CudaKernel { bool isinf_only_, isnan_only_; }; -template -struct IsAllFiniteFunctor { - void operator()(cudaStream_t stream, ChunkGroup<1> chunks, bool* output, const bool isinf_only, const bool isnan_only); -}; - } // namespace cuda } // namespace onnxruntime \ No newline at end of file diff --git a/onnxruntime/contrib_ops/cuda/math/isfinite.cu b/onnxruntime/contrib_ops/cuda/math/isfinite_impl.cu similarity index 99% rename from onnxruntime/contrib_ops/cuda/math/isfinite.cu rename to onnxruntime/contrib_ops/cuda/math/isfinite_impl.cu index fe2f0e2af0..68d9662d45 100644 --- a/onnxruntime/contrib_ops/cuda/math/isfinite.cu +++ b/onnxruntime/contrib_ops/cuda/math/isfinite_impl.cu @@ -2,6 +2,7 @@ // Licensed under the MIT License. #include +#include "isfinite_impl.h" #include "core/providers/cuda/cu_inc/common.cuh" #include "contrib_ops/cuda/math/isfinite.cuh" diff --git a/onnxruntime/contrib_ops/cuda/math/isfinite_impl.h b/onnxruntime/contrib_ops/cuda/math/isfinite_impl.h new file mode 100644 index 0000000000..05f2f31428 --- /dev/null +++ b/onnxruntime/contrib_ops/cuda/math/isfinite_impl.h @@ -0,0 +1,18 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#pragma once + +#include +#include "core/providers/cuda/multi_tensor/common.cuh" + +namespace onnxruntime { +namespace cuda { + +template +struct IsAllFiniteFunctor { + void operator()(cudaStream_t stream, ChunkGroup<1> chunks, bool* output, const bool isinf_only, const bool isnan_only); +}; + +} +} // namespace onnxruntime \ No newline at end of file diff --git a/onnxruntime/core/providers/cuda/cuda_allocator.h b/onnxruntime/core/providers/cuda/cuda_allocator.h index 2c56a9444b..c81292cd26 100644 --- a/onnxruntime/core/providers/cuda/cuda_allocator.h +++ b/onnxruntime/core/providers/cuda/cuda_allocator.h @@ -3,7 +3,8 @@ #pragma once -#include + +#include "core/common/inlined_containers.h" #include "core/framework/allocator.h" #include "core/platform/ort_mutex.h" @@ -47,7 +48,7 @@ class CUDAExternalAllocator : public CUDAAllocator { ExternalAlloc alloc_; ExternalFree free_; ExternalEmptyCache empty_cache_; - std::unordered_set reserved_; + InlinedHashSet reserved_; }; //TODO: add a default constructor diff --git a/onnxruntime/core/providers/cuda/multi_tensor/common.cuh b/onnxruntime/core/providers/cuda/multi_tensor/common.cuh index 84cd10ad24..508cf56297 100644 --- a/onnxruntime/core/providers/cuda/multi_tensor/common.cuh +++ b/onnxruntime/core/providers/cuda/multi_tensor/common.cuh @@ -9,6 +9,8 @@ #pragma once #include +#include "core/common/common.h" + namespace onnxruntime { namespace cuda { // initial reference from: @@ -80,9 +82,9 @@ void launch_multi_tensor_functor( TMultiTensorFunctor multipleTensorKernel, TFunctorParams&&... kernelParams) { ORT_ENFORCE(tensor_sizes.size() > 0); - ORT_ENFORCE(tensor_sizes.size() < static_cast(std::numeric_limits::max())); + ORT_ENFORCE(tensor_sizes.size() < static_cast(INT_MAX)); ORT_ENFORCE(grouped_tensor_pointers.size() > 0); - ORT_ENFORCE(grouped_tensor_pointers.size() < static_cast(std::numeric_limits::max())); + ORT_ENFORCE(grouped_tensor_pointers.size() < static_cast(INT_MAX)); ORT_ENFORCE(chunk_size > 0); // Number of groups, for example, the number of updated weight tensors in Lamb optimizer. const int group_count = static_cast(grouped_tensor_pointers.size()); @@ -92,7 +94,7 @@ void launch_multi_tensor_functor( int block_index = 0; // Check if 32-bit integer is enough. - ORT_ENFORCE(tensor_sizes.size() < static_cast(std::numeric_limits::max())); + ORT_ENFORCE(tensor_sizes.size() < static_cast(INT_MAX)); ORT_ENFORCE(grouped_tensor_pointers.size() == tensor_sizes.size()); ORT_ENFORCE(group_size == ACTUAL_TENSOR_GROUP_SIZE[TensorGroupSize]); for (int i = 0; i < group_count; ++i) { diff --git a/orttraining/orttraining/training_ops/cuda/math/isfinite.cc b/orttraining/orttraining/training_ops/cuda/math/isfinite.cc index 9741ad2d54..6f58fca768 100644 --- a/orttraining/orttraining/training_ops/cuda/math/isfinite.cc +++ b/orttraining/orttraining/training_ops/cuda/math/isfinite.cc @@ -2,6 +2,7 @@ // Licensed under the MIT License. #include "orttraining/training_ops/cuda/math/isfinite.h" +#include "orttraining/training_ops/cuda/math/isfinite_impl.h" using namespace ONNX_NAMESPACE; using namespace onnxruntime::common; diff --git a/orttraining/orttraining/training_ops/cuda/math/isfinite.h b/orttraining/orttraining/training_ops/cuda/math/isfinite.h index 91e4c051c1..fda4a76412 100644 --- a/orttraining/orttraining/training_ops/cuda/math/isfinite.h +++ b/orttraining/orttraining/training_ops/cuda/math/isfinite.h @@ -2,9 +2,9 @@ // Licensed under the MIT License. #pragma once + #include "core/common/common.h" #include "core/providers/cuda/cuda_kernel.h" -#include "core/providers/cuda/multi_tensor/common.cuh" namespace onnxruntime { namespace cuda { @@ -18,8 +18,5 @@ class IsFiniteOp final : public CudaKernel { Status ComputeInternal(OpKernelContext* context) const override; }; -template -void IsFinite(cudaStream_t stream, const TSrc* input, bool* output, size_t N); - } // namespace cuda } // namespace onnxruntime \ No newline at end of file diff --git a/orttraining/orttraining/training_ops/cuda/math/isfinite.cu b/orttraining/orttraining/training_ops/cuda/math/isfinite_impl.cu similarity index 97% rename from orttraining/orttraining/training_ops/cuda/math/isfinite.cu rename to orttraining/orttraining/training_ops/cuda/math/isfinite_impl.cu index fb1e938ead..f4a1a95e2a 100644 --- a/orttraining/orttraining/training_ops/cuda/math/isfinite.cu +++ b/orttraining/orttraining/training_ops/cuda/math/isfinite_impl.cu @@ -1,10 +1,12 @@ // Copyright (c) Microsoft Corporation. All rights reserved. // Licensed under the MIT License. +#include "isfinite_impl.h" #include #include "core/providers/cuda/cu_inc/common.cuh" #include "contrib_ops/cuda/math/isfinite.cuh" + namespace onnxruntime { namespace cuda { diff --git a/orttraining/orttraining/training_ops/cuda/math/isfinite_impl.h b/orttraining/orttraining/training_ops/cuda/math/isfinite_impl.h new file mode 100644 index 0000000000..842be98775 --- /dev/null +++ b/orttraining/orttraining/training_ops/cuda/math/isfinite_impl.h @@ -0,0 +1,15 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#pragma once + +#include + +namespace onnxruntime { +namespace cuda { + +template +void IsFinite(cudaStream_t stream, const TSrc* input, bool* output, size_t N); + +} +} // namespace onnxruntime diff --git a/orttraining/orttraining/training_ops/cuda/math/mixed_precision_scale.cc b/orttraining/orttraining/training_ops/cuda/math/mixed_precision_scale.cc index 636cb2d8d5..9ebd7a03f0 100644 --- a/orttraining/orttraining/training_ops/cuda/math/mixed_precision_scale.cc +++ b/orttraining/orttraining/training_ops/cuda/math/mixed_precision_scale.cc @@ -2,6 +2,7 @@ // Licensed under the MIT License. #include "mixed_precision_scale.h" +#include "mixed_precision_scale_impl.h" using namespace ONNX_NAMESPACE; using namespace onnxruntime::common; diff --git a/orttraining/orttraining/training_ops/cuda/math/mixed_precision_scale.h b/orttraining/orttraining/training_ops/cuda/math/mixed_precision_scale.h index b5400cc30e..48195b50bc 100644 --- a/orttraining/orttraining/training_ops/cuda/math/mixed_precision_scale.h +++ b/orttraining/orttraining/training_ops/cuda/math/mixed_precision_scale.h @@ -8,14 +8,6 @@ namespace onnxruntime { namespace cuda { -template -void Impl_MixedPrecisionScale( - cudaStream_t stream, - const SrcT* input_data, - const float* scale_data, - DstT* output_data, - size_t count); - template class MixedPrecisionScale final : public CudaKernel { public: diff --git a/orttraining/orttraining/training_ops/cuda/math/mixed_precision_scale.cu b/orttraining/orttraining/training_ops/cuda/math/mixed_precision_scale_impl.cu similarity index 97% rename from orttraining/orttraining/training_ops/cuda/math/mixed_precision_scale.cu rename to orttraining/orttraining/training_ops/cuda/math/mixed_precision_scale_impl.cu index cc38593376..0a32141eea 100644 --- a/orttraining/orttraining/training_ops/cuda/math/mixed_precision_scale.cu +++ b/orttraining/orttraining/training_ops/cuda/math/mixed_precision_scale_impl.cu @@ -5,9 +5,9 @@ #pragma warning(disable : 4244) #endif +#include "mixed_precision_scale_impl.h" #include #include "core/providers/cuda/cu_inc/common.cuh" -#include "mixed_precision_scale.h" namespace onnxruntime { namespace cuda { diff --git a/orttraining/orttraining/training_ops/cuda/math/mixed_precision_scale_impl.h b/orttraining/orttraining/training_ops/cuda/math/mixed_precision_scale_impl.h new file mode 100644 index 0000000000..49a2b192bd --- /dev/null +++ b/orttraining/orttraining/training_ops/cuda/math/mixed_precision_scale_impl.h @@ -0,0 +1,19 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#pragma once + +#include + +namespace onnxruntime { +namespace cuda { +template +void Impl_MixedPrecisionScale( + cudaStream_t stream, + const SrcT* input_data, + const float* scale_data, + DstT* output_data, + size_t count); +} +} // namespace onnxruntime + diff --git a/orttraining/orttraining/training_ops/cuda/math/scale.cc b/orttraining/orttraining/training_ops/cuda/math/scale.cc index 8b2ebb45a3..9569131ed5 100644 --- a/orttraining/orttraining/training_ops/cuda/math/scale.cc +++ b/orttraining/orttraining/training_ops/cuda/math/scale.cc @@ -2,6 +2,7 @@ // Licensed under the MIT License. #include "orttraining/training_ops/cuda/math/scale.h" +#include "orttraining/training_ops/cuda/math/scale_impl.h" using namespace ONNX_NAMESPACE; using namespace onnxruntime::common; @@ -25,6 +26,15 @@ namespace cuda { .InputMemoryType(OrtMemTypeCPUInput, 1), \ Scale); +template +struct GetScaleValueImpl { + void operator()(const Tensor* scale, float& scale_value) const { + ORT_ENFORCE(scale->Shape().Size() == 1, "Scale input should have a single value."); + scale_value = static_cast(*(scale->template Data())); + ORT_ENFORCE(scale_value != 0.0f, "Scale value must not be 0."); + } +}; + template Scale::Scale(const OpKernelInfo& info) : CudaKernel(info) { int64_t scale_down; diff --git a/orttraining/orttraining/training_ops/cuda/math/scale.h b/orttraining/orttraining/training_ops/cuda/math/scale.h index 020f4efbdb..58270447fa 100644 --- a/orttraining/orttraining/training_ops/cuda/math/scale.h +++ b/orttraining/orttraining/training_ops/cuda/math/scale.h @@ -7,23 +7,6 @@ namespace onnxruntime { namespace cuda { -template -struct GetScaleValueImpl { - void operator()(const Tensor* scale, float& scale_value) const { - ORT_ENFORCE(scale->Shape().Size() == 1, "Scale input should have a single value."); - scale_value = static_cast(*(scale->template Data())); - ORT_ENFORCE(scale_value != 0.0f, "Scale value must not be 0."); - } -}; - -template -void Impl_Scale( - cudaStream_t stream, - const T* input_data, - const float scale_value, - T* output_data, - size_t count); - template class Scale final : public CudaKernel { public: diff --git a/orttraining/orttraining/training_ops/cuda/math/scale.cu b/orttraining/orttraining/training_ops/cuda/math/scale_impl.cu similarity index 96% rename from orttraining/orttraining/training_ops/cuda/math/scale.cu rename to orttraining/orttraining/training_ops/cuda/math/scale_impl.cu index b132665039..376a9696c7 100644 --- a/orttraining/orttraining/training_ops/cuda/math/scale.cu +++ b/orttraining/orttraining/training_ops/cuda/math/scale_impl.cu @@ -1,8 +1,8 @@ // Copyright (c) Microsoft Corporation. All rights reserved. // Licensed under the MIT License. +#include "orttraining/training_ops/cuda/math/scale_impl.h" #include "core/providers/cuda/cu_inc/common.cuh" -#include "orttraining/training_ops/cuda/math/scale.h" namespace onnxruntime { namespace cuda { diff --git a/orttraining/orttraining/training_ops/cuda/math/scale_impl.h b/orttraining/orttraining/training_ops/cuda/math/scale_impl.h new file mode 100644 index 0000000000..6f35a485b5 --- /dev/null +++ b/orttraining/orttraining/training_ops/cuda/math/scale_impl.h @@ -0,0 +1,19 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#pragma once + +#include + +namespace onnxruntime { +namespace cuda { + +template +void Impl_Scale( + cudaStream_t stream, + const T* input_data, + const float scale_value, + T* output_data, + size_t count); + } +} // namespace onnxruntime diff --git a/orttraining/orttraining/training_ops/cuda/math/softmax_grad.cc b/orttraining/orttraining/training_ops/cuda/math/softmax_grad.cc index 1a8af5045b..1a85f4a87b 100644 --- a/orttraining/orttraining/training_ops/cuda/math/softmax_grad.cc +++ b/orttraining/orttraining/training_ops/cuda/math/softmax_grad.cc @@ -2,6 +2,7 @@ // Licensed under the MIT License. #include "orttraining/training_ops/cuda/math/softmax_grad.h" +#include "orttraining/training_ops/cuda/math/softmax_grad_impl.h" #include "core/providers/common.h" #include "core/providers/cuda/cudnn_common.h" diff --git a/orttraining/orttraining/training_ops/cuda/math/softmax_grad.h b/orttraining/orttraining/training_ops/cuda/math/softmax_grad.h index 21543ca9a2..c58d24e1e7 100644 --- a/orttraining/orttraining/training_ops/cuda/math/softmax_grad.h +++ b/orttraining/orttraining/training_ops/cuda/math/softmax_grad.h @@ -8,9 +8,6 @@ namespace onnxruntime { namespace cuda { -template -void dispatch_softmax_backward(cudaStream_t stream, output_t* grad_input, const input_t* grad, const input_t* output, int softmax_elements, int softmax_elements_stride, int batch_count); - template class SoftmaxGrad final : public CudaKernel { public: diff --git a/orttraining/orttraining/training_ops/cuda/math/softmax_grad_impl.cu b/orttraining/orttraining/training_ops/cuda/math/softmax_grad_impl.cu index 17da25cc9d..608d4223f5 100644 --- a/orttraining/orttraining/training_ops/cuda/math/softmax_grad_impl.cu +++ b/orttraining/orttraining/training_ops/cuda/math/softmax_grad_impl.cu @@ -18,7 +18,7 @@ // The code below is mostly copied from Pytorch PersistentSoftmax.cuh -#include "orttraining/training_ops/cuda/math/softmax_grad.h" +#include "orttraining/training_ops/cuda/math/softmax_grad_impl.h" #include "core/providers/cuda/cu_inc/common.cuh" #include "core/providers/cuda/math/softmax_warpwise_impl.cuh" diff --git a/orttraining/orttraining/training_ops/cuda/math/softmax_grad_impl.h b/orttraining/orttraining/training_ops/cuda/math/softmax_grad_impl.h new file mode 100644 index 0000000000..ac4285f326 --- /dev/null +++ b/orttraining/orttraining/training_ops/cuda/math/softmax_grad_impl.h @@ -0,0 +1,14 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#pragma once + +#include + +namespace onnxruntime { +namespace cuda { +template +void dispatch_softmax_backward(cudaStream_t stream, output_t* grad_input, const input_t* grad, const input_t* output, + int softmax_elements, int softmax_elements_stride, int batch_count); +} +} // namespace onnxruntime diff --git a/orttraining/orttraining/training_ops/cuda/optimizer/adam.cc b/orttraining/orttraining/training_ops/cuda/optimizer/adam.cc index d8e424a439..8318ab5553 100644 --- a/orttraining/orttraining/training_ops/cuda/optimizer/adam.cc +++ b/orttraining/orttraining/training_ops/cuda/optimizer/adam.cc @@ -1,11 +1,13 @@ // Copyright (c) Microsoft Corporation. All rights reserved. // Licensed under the MIT License. +#include "orttraining/training_ops/cuda/optimizer/adam.h" +#include "orttraining/training_ops/cuda/optimizer/adam_impl.h" + #include "core/providers/cuda/cuda_allocator.h" #include "core/providers/cuda/reduction/reduction_functions.h" #include "core/providers/cuda/math/binary_elementwise_ops.h" #include "orttraining/training_ops/cuda/optimizer/common.h" -#include "orttraining/training_ops/cuda/optimizer/adam.h" namespace onnxruntime { namespace cuda { diff --git a/orttraining/orttraining/training_ops/cuda/optimizer/adam.h b/orttraining/orttraining/training_ops/cuda/optimizer/adam.h index 74f57e419f..0a894f318c 100644 --- a/orttraining/orttraining/training_ops/cuda/optimizer/adam.h +++ b/orttraining/orttraining/training_ops/cuda/optimizer/adam.h @@ -7,32 +7,6 @@ namespace onnxruntime { namespace cuda { - -template -void AdamOptimizerImpl( - cudaStream_t stream, - const T1* eta, - const T2 update_count, - const T3* weights, - const T_GRAD* grads, - const T4* moment_1, - const T4* moment_2, - const T3* loss_scale, - const T_GRAD_NORM* grad_norm, - const float alpha, - const float beta, - const float lambda, - const float epsilon, - const float max_norm, - const bool do_bias_correction, - const int64_t weight_decay_mode, - T4* moment_1_out, - T4* moment_2_out, - T3* weights_out, - T_GRAD* grads_out, - T_MIXED_PRECISION_FP* mixed_precision_weights_out, - size_t count); - template class AdamOptimizer final : public CudaKernel { public: diff --git a/orttraining/orttraining/training_ops/cuda/optimizer/adam.cu b/orttraining/orttraining/training_ops/cuda/optimizer/adam_impl.cu similarity index 99% rename from orttraining/orttraining/training_ops/cuda/optimizer/adam.cu rename to orttraining/orttraining/training_ops/cuda/optimizer/adam_impl.cu index d72b9f924e..95ee9d2dae 100644 --- a/orttraining/orttraining/training_ops/cuda/optimizer/adam.cu +++ b/orttraining/orttraining/training_ops/cuda/optimizer/adam_impl.cu @@ -1,10 +1,10 @@ // Copyright (c) Microsoft Corporation. All rights reserved. // Licensed under the MIT License. +#include "orttraining/training_ops/cuda/optimizer/adam_impl.h" #include "core/providers/cuda/cuda_common.h" #include "core/providers/cuda/cu_inc/common.cuh" #include "orttraining/training_ops/cuda/optimizer/common.cuh" -#include "orttraining/training_ops/cuda/optimizer/adam.h" #include "orttraining/training_ops/cuda/optimizer/common.h" namespace onnxruntime { diff --git a/orttraining/orttraining/training_ops/cuda/optimizer/adam_impl.h b/orttraining/orttraining/training_ops/cuda/optimizer/adam_impl.h new file mode 100644 index 0000000000..d575dfc78c --- /dev/null +++ b/orttraining/orttraining/training_ops/cuda/optimizer/adam_impl.h @@ -0,0 +1,37 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#pragma once + +#include +#include + +namespace onnxruntime { +namespace cuda { + +template +void AdamOptimizerImpl( + cudaStream_t stream, + const T1* eta, + const T2 update_count, + const T3* weights, + const T_GRAD* grads, + const T4* moment_1, + const T4* moment_2, + const T3* loss_scale, + const T_GRAD_NORM* grad_norm, + const float alpha, + const float beta, + const float lambda, + const float epsilon, + const float max_norm, + const bool do_bias_correction, + const int64_t weight_decay_mode, + T4* moment_1_out, + T4* moment_2_out, + T3* weights_out, + T_GRAD* grads_out, + T_MIXED_PRECISION_FP* mixed_precision_weights_out, + size_t count); +} +} // namespace onnxruntime diff --git a/orttraining/orttraining/training_ops/cuda/optimizer/gradient_control.cc b/orttraining/orttraining/training_ops/cuda/optimizer/gradient_control.cc index 5a7815ff60..d1ddeee3bb 100644 --- a/orttraining/orttraining/training_ops/cuda/optimizer/gradient_control.cc +++ b/orttraining/orttraining/training_ops/cuda/optimizer/gradient_control.cc @@ -1,11 +1,13 @@ // Copyright (c) Microsoft Corporation. All rights reserved. // Licensed under the MIT License. +#include "gradient_control.h" +#include "gradient_control_impl.h" + #include "core/providers/cuda/math/binary_elementwise_ops.h" #include "core/providers/cuda/reduction/reduction_functions.h" #include "core/providers/cuda/cuda_allocator.h" #include "common.h" -#include "gradient_control.h" namespace onnxruntime { namespace cuda { diff --git a/orttraining/orttraining/training_ops/cuda/optimizer/gradient_control.h b/orttraining/orttraining/training_ops/cuda/optimizer/gradient_control.h index c2a4f8e234..a4f60bc66b 100644 --- a/orttraining/orttraining/training_ops/cuda/optimizer/gradient_control.h +++ b/orttraining/orttraining/training_ops/cuda/optimizer/gradient_control.h @@ -22,14 +22,5 @@ class InPlaceAccumulator final : public CudaKernel { Status ComputeInternal(OpKernelContext* context) const override; }; -// Implementation can be found in cuda file, optimizers_impl.cu -template -void InPlaceAccumulatorImpl( - cudaStream_t stream, - const T* gradient_buffer, - const T_GRAD* gradient, - T* accumulated_gradient, - size_t count); - } // namespace cuda } // namespace onnxruntime \ No newline at end of file diff --git a/orttraining/orttraining/training_ops/cuda/optimizer/gradient_control.cu b/orttraining/orttraining/training_ops/cuda/optimizer/gradient_control_impl.cu similarity index 98% rename from orttraining/orttraining/training_ops/cuda/optimizer/gradient_control.cu rename to orttraining/orttraining/training_ops/cuda/optimizer/gradient_control_impl.cu index 48295f5a0c..e6cd44c215 100644 --- a/orttraining/orttraining/training_ops/cuda/optimizer/gradient_control.cu +++ b/orttraining/orttraining/training_ops/cuda/optimizer/gradient_control_impl.cu @@ -2,10 +2,10 @@ // Copyright (c) Microsoft Corporation. All rights reserved. // Licensed under the MIT License. +#include "gradient_control_impl.h" #include "core/providers/cuda/cuda_common.h" #include "core/providers/cuda/cu_inc/common.cuh" #include "core/providers/cuda/atomic/common.cuh" -#include "gradient_control.h" namespace onnxruntime { namespace cuda { diff --git a/orttraining/orttraining/training_ops/cuda/optimizer/gradient_control_impl.h b/orttraining/orttraining/training_ops/cuda/optimizer/gradient_control_impl.h new file mode 100644 index 0000000000..2864094f9a --- /dev/null +++ b/orttraining/orttraining/training_ops/cuda/optimizer/gradient_control_impl.h @@ -0,0 +1,20 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#pragma once + +#include + +namespace onnxruntime { +namespace cuda { +// Implementation can be found in cuda file +template +void InPlaceAccumulatorImpl( + cudaStream_t stream, + const T* gradient_buffer, + const T_GRAD* gradient, + T* accumulated_gradient, + size_t count); +} +} // namespace onnxruntime + diff --git a/orttraining/orttraining/training_ops/cuda/optimizer/lamb.cc b/orttraining/orttraining/training_ops/cuda/optimizer/lamb.cc index 17cff8db25..205db1640f 100644 --- a/orttraining/orttraining/training_ops/cuda/optimizer/lamb.cc +++ b/orttraining/orttraining/training_ops/cuda/optimizer/lamb.cc @@ -1,12 +1,15 @@ // Copyright (c) Microsoft Corporation. All rights reserved. // Licensed under the MIT License. -#include -#include "core/providers/cuda/cuda_allocator.h" + +#include "orttraining/training_ops/cuda/optimizer/lamb.h" +#include "orttraining/training_ops/cuda/optimizer/lamb_impl.h" + #include "core/providers/cuda/reduction/reduction_functions.h" #include "core/providers/cuda/math/binary_elementwise_ops.h" #include "orttraining/training_ops/cuda/optimizer/common.h" -#include "orttraining/training_ops/cuda/optimizer/lamb.h" + +#include namespace onnxruntime { namespace cuda { diff --git a/orttraining/orttraining/training_ops/cuda/optimizer/lamb.h b/orttraining/orttraining/training_ops/cuda/optimizer/lamb.h index 3be720fb8b..8ec0335465 100644 --- a/orttraining/orttraining/training_ops/cuda/optimizer/lamb.h +++ b/orttraining/orttraining/training_ops/cuda/optimizer/lamb.h @@ -4,7 +4,6 @@ #pragma once #include "core/common/common.h" #include "core/providers/cuda/cuda_kernel.h" -#include "core/providers/cuda/multi_tensor/common.cuh" namespace onnxruntime { namespace cuda { @@ -43,156 +42,5 @@ class LambOptimizer final : public CudaKernel { bool do_bias_correction_; }; -// Implementation can be found in cuda file, optimizers_impl.cu -// T1's precision should be higher than T2. It's used for -// large tensors. Small tensors should use multi-tensor version -// of this. -template -void LambComputeDirection( - cudaStream_t stream, - const T1* weights, - const T2* grads, - const T3* moment_1, - const T3* moment_2, - const T1* loss_scale, - const T_GRAD_NORM* grad_norm, - float alpha, - float beta, - float lambda, - float epsilon, - float max_norm, - float alpha_correction, - float beta_correction, - T2* update_direction, - T3* moment_1_out, - T3* moment_2_out, - size_t count); - -// Implementation can be found in cuda file, optimizers_impl.cu -// T2's precision should be higher than T1. It's used for -// large tensors. Small tensors should use multi-tensor version -// of this. -template -void LambUpdate( - cudaStream_t stream, - const T1* eta, - const float ratio_min, - const float ratio_max, - const T2* r_norm, - const T2* w_norm, - const T2* weights, - const T3* update_direction, - T2* weights_out, - T3* gradients_out, - T_MIXED_PRECISION_FP* mixed_precision_weights_out, - size_t count); - -// Lamb's stage 1 maps [w, g, m1, m2] to [d, m1_new, m2_new] where -// w: weight tensor -// g: gradient (reused to store update direction) -// m1: 1st momentum -// m2: 2nd momentum -// d: update direction -// m1_new: updated 1st momentum -// m2_new: updated 2nd momentum -// Because we reuse g to store d, there are only 6 tensors in total and -// therefore the type of chunk_group is ChunkGroup<6>. -// -// Tensor pointers associated with the i-th tensor in this chunk: -// w: chunk_group.tensor_ptrs[0][i] -// g (or d): chunk_group.tensor_ptrs[1][i] -// m1: chunk_group.tensor_ptrs[2][i] -// m2: chunk_group.tensor_ptrs[3][i] -// m1_new: chunk_group.tensor_ptrs[4][i] -// m2_new: chunk_group.tensor_ptrs[5][i] -template -struct LambMultiTensorComputeDirectionFunctor { - void operator()( - cudaStream_t stream, - ChunkGroup<6> chunk_group, - const T1* loss_scale, - const T_GRAD_NORM* grad_norm, - const float lambda, - const float alpha, - const float beta, - const float epsilon, - const float max_norm, - const float alpha_correction, - const float beta_correction); -}; - -// Lamb's reduction maps [w, d] to [w_norm, d_norm] where -// w: weight tensor -// d: update direction -// w_norm: norm of w -// d_norm: norm of d -// There are 4 distinct tensors in total and therefore the -// type of chunk_group is ChunkGroup<4>. -// -// Tensor pointers associated with the i-th tensor in this chunk: -// w: chunk_group.tensor_ptrs[0][i] -// d: chunk_group.tensor_ptrs[1][i] -// w_norm: chunk_group.tensor_ptrs[2][i] -// d_norm: chunk_group.tensor_ptrs[3][i] -template -struct LambMultiTensorReductionFunctor { - void operator()( - cudaStream_t stream, - ChunkGroup<4> chunk_group, - const CudaKernel& kernel, - void* reduction_buffer, - size_t reduction_buffer_size); -}; - -// Lamb's reduction mapping [w, d] to [w_norm, d_norm] spans multiples thread blocks -// -// This includes any block-index for which it holds -// i-th tensor-index == chunk_group.block_index_to_tensor_group_index[ block-index ] -// and where i-th tensor-index corresponds to tensor group w(i), d(i), w_norm(i), d_norm(i) -// (see above) -// -// The above span of blocks corresponding i-th tensor will be contiguous. -// To perform an ORDERED reduction across the thread blocks for i-th tensor, -// the following struct is passed for every tensor. -// It consists of fields: -// 'leading_block' := lowest block-index corresponding i-th tensor -// 'number_blocks' := number block-index " -// 'completed_blocks' := initialized to zero (for internal use) -// Note 'completed_blocks' prevents inter-block reduction until intra-block reduction is complete. -struct LambMultiTensorSyncRangeAndLock { - int leading_block; - int number_blocks; - int completed_blocks; -}; - -// Lamb's stage 2 maps [w_norm, w_norm, w, d] to [w_new, g_new, w_mixed_precision_new] where -// w_norm: norm of w -// d_norm: norm of d -// w: weight tensor -// d: update direction -// w_new: updated weight tensor -// g_new: updated gradient tensor -// w_mixed_precision_new: updated weight tensor of mixed-precision type -// There are 7 distinct tensors in total and therefore the -// type of chunk_group is ChunkGroup<7>. -// -// Tensor pointers associated with the i-th tensor in this chunk: -// w_norm: chunk_group.tensor_ptrs[0][i] -// d_norm: chunk_group.tensor_ptrs[1][i] -// w: chunk_group.tensor_ptrs[2][i] -// d: chunk_group.tensor_ptrs[3][i] -// w_new: chunk_group.tensor_ptrs[4][i] -// g_new: chunk_group.tensor_ptrs[5][i] -// w_mixed_precision_new: chunk_group.tensor_ptrs[6][i] -template -struct LambMultiTensorUpdateFunctor { - void operator()( - cudaStream_t stream, - ChunkGroup<7> chunk_group, - const T1* eta, - const float ratio_min, - const float ratio_max); -}; - } // namespace cuda } // namespace onnxruntime diff --git a/orttraining/orttraining/training_ops/cuda/optimizer/lamb.cu b/orttraining/orttraining/training_ops/cuda/optimizer/lamb_impl.cu similarity index 99% rename from orttraining/orttraining/training_ops/cuda/optimizer/lamb.cu rename to orttraining/orttraining/training_ops/cuda/optimizer/lamb_impl.cu index c7c9936ed1..683bb23910 100644 --- a/orttraining/orttraining/training_ops/cuda/optimizer/lamb.cu +++ b/orttraining/orttraining/training_ops/cuda/optimizer/lamb_impl.cu @@ -1,15 +1,13 @@ // Copyright (c) Microsoft Corporation. All rights reserved. // Licensed under the MIT License. +#include "orttraining/training_ops/cuda/optimizer/lamb_impl.h" #include "core/providers/cuda/cu_inc/common.cuh" -#include "core/providers/cuda/cuda_allocator.h" #include "core/providers/cuda/cuda_common.h" #include "core/providers/cuda/atomic/common.cuh" #include "core/providers/cuda/reduction/reduction_utils.cuh" #include "contrib_ops/cuda/math/isfinite.cuh" -#include "orttraining/training_ops/cuda/optimizer/common.h" #include "orttraining/training_ops/cuda/optimizer/common.cuh" -#include "orttraining/training_ops/cuda/optimizer/lamb.h" namespace onnxruntime { namespace cuda { diff --git a/orttraining/orttraining/training_ops/cuda/optimizer/lamb_impl.h b/orttraining/orttraining/training_ops/cuda/optimizer/lamb_impl.h new file mode 100644 index 0000000000..c1a6e19f5b --- /dev/null +++ b/orttraining/orttraining/training_ops/cuda/optimizer/lamb_impl.h @@ -0,0 +1,164 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#pragma once + +#include +#include "core/providers/cuda/cuda_kernel.h" +#include "core/providers/cuda/multi_tensor/common.cuh" + +namespace onnxruntime { +namespace cuda { +// Implementation can be found in cuda file. +// T1's precision should be higher than T2. It's used for +// large tensors. Small tensors should use multi-tensor version +// of this. +template +void LambComputeDirection( + cudaStream_t stream, + const T1* weights, + const T2* grads, + const T3* moment_1, + const T3* moment_2, + const T1* loss_scale, + const T_GRAD_NORM* grad_norm, + float alpha, + float beta, + float lambda, + float epsilon, + float max_norm, + float alpha_correction, + float beta_correction, + T2* update_direction, + T3* moment_1_out, + T3* moment_2_out, + size_t count); + +// Implementation can be found in cuda file. +// T2's precision should be higher than T1. It's used for +// large tensors. Small tensors should use multi-tensor version +// of this. +template +void LambUpdate( + cudaStream_t stream, + const T1* eta, + const float ratio_min, + const float ratio_max, + const T2* r_norm, + const T2* w_norm, + const T2* weights, + const T3* update_direction, + T2* weights_out, + T3* gradients_out, + T_MIXED_PRECISION_FP* mixed_precision_weights_out, + size_t count); + +// Lamb's stage 1 maps [w, g, m1, m2] to [d, m1_new, m2_new] where +// w: weight tensor +// g: gradient (reused to store update direction) +// m1: 1st momentum +// m2: 2nd momentum +// d: update direction +// m1_new: updated 1st momentum +// m2_new: updated 2nd momentum +// Because we reuse g to store d, there are only 6 tensors in total and +// therefore the type of chunk_group is ChunkGroup<6>. +// +// Tensor pointers associated with the i-th tensor in this chunk: +// w: chunk_group.tensor_ptrs[0][i] +// g (or d): chunk_group.tensor_ptrs[1][i] +// m1: chunk_group.tensor_ptrs[2][i] +// m2: chunk_group.tensor_ptrs[3][i] +// m1_new: chunk_group.tensor_ptrs[4][i] +// m2_new: chunk_group.tensor_ptrs[5][i] +template +struct LambMultiTensorComputeDirectionFunctor { + void operator()( + cudaStream_t stream, + ChunkGroup<6> chunk_group, + const T1* loss_scale, + const T_GRAD_NORM* grad_norm, + const float lambda, + const float alpha, + const float beta, + const float epsilon, + const float max_norm, + const float alpha_correction, + const float beta_correction); +}; + +// Lamb's reduction maps [w, d] to [w_norm, d_norm] where +// w: weight tensor +// d: update direction +// w_norm: norm of w +// d_norm: norm of d +// There are 4 distinct tensors in total and therefore the +// type of chunk_group is ChunkGroup<4>. +// +// Tensor pointers associated with the i-th tensor in this chunk: +// w: chunk_group.tensor_ptrs[0][i] +// d: chunk_group.tensor_ptrs[1][i] +// w_norm: chunk_group.tensor_ptrs[2][i] +// d_norm: chunk_group.tensor_ptrs[3][i] +template +struct LambMultiTensorReductionFunctor { + void operator()( + cudaStream_t stream, + ChunkGroup<4> chunk_group, + const CudaKernel& kernel, + void* reduction_buffer, + size_t reduction_buffer_size); +}; + +// Lamb's reduction mapping [w, d] to [w_norm, d_norm] spans multiples thread blocks +// +// This includes any block-index for which it holds +// i-th tensor-index == chunk_group.block_index_to_tensor_group_index[ block-index ] +// and where i-th tensor-index corresponds to tensor group w(i), d(i), w_norm(i), d_norm(i) +// (see above) +// +// The above span of blocks corresponding i-th tensor will be contiguous. +// To perform an ORDERED reduction across the thread blocks for i-th tensor, +// the following struct is passed for every tensor. +// It consists of fields: +// 'leading_block' := lowest block-index corresponding i-th tensor +// 'number_blocks' := number block-index " +// 'completed_blocks' := initialized to zero (for internal use) +// Note 'completed_blocks' prevents inter-block reduction until intra-block reduction is complete. +struct LambMultiTensorSyncRangeAndLock { + int leading_block; + int number_blocks; + int completed_blocks; +}; + +// Lamb's stage 2 maps [w_norm, w_norm, w, d] to [w_new, g_new, w_mixed_precision_new] where +// w_norm: norm of w +// d_norm: norm of d +// w: weight tensor +// d: update direction +// w_new: updated weight tensor +// g_new: updated gradient tensor +// w_mixed_precision_new: updated weight tensor of mixed-precision type +// There are 7 distinct tensors in total and therefore the +// type of chunk_group is ChunkGroup<7>. +// +// Tensor pointers associated with the i-th tensor in this chunk: +// w_norm: chunk_group.tensor_ptrs[0][i] +// d_norm: chunk_group.tensor_ptrs[1][i] +// w: chunk_group.tensor_ptrs[2][i] +// d: chunk_group.tensor_ptrs[3][i] +// w_new: chunk_group.tensor_ptrs[4][i] +// g_new: chunk_group.tensor_ptrs[5][i] +// w_mixed_precision_new: chunk_group.tensor_ptrs[6][i] +template +struct LambMultiTensorUpdateFunctor { + void operator()( + cudaStream_t stream, + ChunkGroup<7> chunk_group, + const T1* eta, + const float ratio_min, + const float ratio_max); +}; + +} +} // namespace onnxruntime diff --git a/orttraining/orttraining/training_ops/cuda/optimizer/sg.cc b/orttraining/orttraining/training_ops/cuda/optimizer/sg.cc index e2133af960..5bfd5fa4c1 100644 --- a/orttraining/orttraining/training_ops/cuda/optimizer/sg.cc +++ b/orttraining/orttraining/training_ops/cuda/optimizer/sg.cc @@ -1,10 +1,11 @@ // Copyright (c) Microsoft Corporation. All rights reserved. // Licensed under the MIT License. -#include "core/providers/cuda/cuda_allocator.h" +#include "sg.h" +#include "sg_impl.h" + #include "core/providers/cuda/reduction/reduction_functions.h" #include "core/providers/cuda/math/binary_elementwise_ops.h" -#include "sg.h" namespace onnxruntime { namespace cuda { diff --git a/orttraining/orttraining/training_ops/cuda/optimizer/sg.h b/orttraining/orttraining/training_ops/cuda/optimizer/sg.h index 80d47a8fa0..aa5f5f2128 100644 --- a/orttraining/orttraining/training_ops/cuda/optimizer/sg.h +++ b/orttraining/orttraining/training_ops/cuda/optimizer/sg.h @@ -8,16 +8,6 @@ namespace onnxruntime { namespace cuda { -template -void SGDOptimizerImpl( - cudaStream_t stream, - const T* eta, - const T* weights, - const T* gradients, - T* weight_out, - T* gradients_out, - size_t count); - class SGDOptimizer final : public CudaKernel { public: SGDOptimizer(const OpKernelInfo& info) : CudaKernel(info) {} diff --git a/orttraining/orttraining/training_ops/cuda/optimizer/sg.cu b/orttraining/orttraining/training_ops/cuda/optimizer/sg_impl.cu similarity index 98% rename from orttraining/orttraining/training_ops/cuda/optimizer/sg.cu rename to orttraining/orttraining/training_ops/cuda/optimizer/sg_impl.cu index aeab19d5eb..1f23309e88 100644 --- a/orttraining/orttraining/training_ops/cuda/optimizer/sg.cu +++ b/orttraining/orttraining/training_ops/cuda/optimizer/sg_impl.cu @@ -1,10 +1,10 @@ // Copyright (c) Microsoft Corporation. All rights reserved. // Licensed under the MIT License. +#include "sg_impl.h" #include "core/providers/cuda/cuda_common.h" #include "core/providers/cuda/cu_inc/common.cuh" #include "core/providers/cuda/atomic/common.cuh" -#include "sg.h" namespace onnxruntime { namespace cuda { diff --git a/orttraining/orttraining/training_ops/cuda/optimizer/sg_impl.h b/orttraining/orttraining/training_ops/cuda/optimizer/sg_impl.h new file mode 100644 index 0000000000..2c37bf164a --- /dev/null +++ b/orttraining/orttraining/training_ops/cuda/optimizer/sg_impl.h @@ -0,0 +1,20 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#pragma once + +#include + +namespace onnxruntime { +namespace cuda { +template +void SGDOptimizerImpl( + cudaStream_t stream, + const T* eta, + const T* weights, + const T* gradients, + T* weight_out, + T* gradients_out, + size_t count); +} +} // namespace onnxruntime diff --git a/orttraining/orttraining/training_ops/cuda/reduction/all.cc b/orttraining/orttraining/training_ops/cuda/reduction/all.cc index 4178d2dd59..4454221847 100644 --- a/orttraining/orttraining/training_ops/cuda/reduction/all.cc +++ b/orttraining/orttraining/training_ops/cuda/reduction/all.cc @@ -2,6 +2,7 @@ // Licensed under the MIT License. #include "orttraining/training_ops/cuda/reduction/all.h" +#include "orttraining/training_ops/cuda/reduction/all_impl.h" namespace onnxruntime { namespace cuda { diff --git a/orttraining/orttraining/training_ops/cuda/reduction/all.h b/orttraining/orttraining/training_ops/cuda/reduction/all.h index f15f3fdff5..7fbff5434a 100644 --- a/orttraining/orttraining/training_ops/cuda/reduction/all.h +++ b/orttraining/orttraining/training_ops/cuda/reduction/all.h @@ -15,8 +15,5 @@ class All final : public CudaKernel { Status ComputeInternal(OpKernelContext* context) const override; }; -template -void LaunchAllKernel(cudaStream_t stream, const T* data, const int size, bool* output); - } // namespace cuda } // namespace onnxruntime diff --git a/orttraining/orttraining/training_ops/cuda/reduction/all.cu b/orttraining/orttraining/training_ops/cuda/reduction/all_impl.cu similarity index 92% rename from orttraining/orttraining/training_ops/cuda/reduction/all.cu rename to orttraining/orttraining/training_ops/cuda/reduction/all_impl.cu index 3647da6903..0da76b0a4b 100644 --- a/orttraining/orttraining/training_ops/cuda/reduction/all.cu +++ b/orttraining/orttraining/training_ops/cuda/reduction/all_impl.cu @@ -1,11 +1,12 @@ // Copyright (c) Microsoft Corporation. All rights reserved. // Licensed under the MIT License. -#include "orttraining/training_ops/cuda/reduction/all.h" +#include "orttraining/training_ops/cuda/reduction/all_impl.h" #include #include #include + #ifdef _WIN32 #pragma warning(disable : 4244) #endif diff --git a/orttraining/orttraining/training_ops/cuda/reduction/all_impl.h b/orttraining/orttraining/training_ops/cuda/reduction/all_impl.h new file mode 100644 index 0000000000..7c1a2fe8c8 --- /dev/null +++ b/orttraining/orttraining/training_ops/cuda/reduction/all_impl.h @@ -0,0 +1,14 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#pragma once + +#include + +namespace onnxruntime { +namespace cuda { +template +void LaunchAllKernel(cudaStream_t stream, const T* data, const int size, bool* output); +} +} // namespace onnxruntime + diff --git a/orttraining/orttraining/training_ops/cuda/reduction/reduction_all.cc b/orttraining/orttraining/training_ops/cuda/reduction/reduction_all.cc index 529206ad16..83e9a546e1 100644 --- a/orttraining/orttraining/training_ops/cuda/reduction/reduction_all.cc +++ b/orttraining/orttraining/training_ops/cuda/reduction/reduction_all.cc @@ -2,6 +2,7 @@ // Licensed under the MIT License. #include "orttraining/training_ops/cuda/reduction/reduction_all.h" +#include "orttraining/training_ops/cuda/reduction/reduction_all_impl.h" #include "core/providers/cuda/reduction/reduction_functions.h" #include "core/providers/cuda/shared_inc/accumulation_type.h" diff --git a/orttraining/orttraining/training_ops/cuda/reduction/reduction_all.h b/orttraining/orttraining/training_ops/cuda/reduction/reduction_all.h index 7de6e2ee9b..431e9005b2 100644 --- a/orttraining/orttraining/training_ops/cuda/reduction/reduction_all.h +++ b/orttraining/orttraining/training_ops/cuda/reduction/reduction_all.h @@ -3,7 +3,6 @@ #pragma once #include "core/providers/cuda/cuda_kernel.h" -#include "core/providers/cuda/multi_tensor/common.cuh" namespace onnxruntime { namespace cuda { @@ -16,13 +15,5 @@ class ReduceAllL2 final : public CudaKernel { Status ComputeInternal(OpKernelContext* context) const override; }; -template -struct MultiTensorReduceL2 { - void operator()(cudaStream_t stream, ChunkGroup<1> chunk_group, TOut* output); -}; - -template -void ScalarSqrt(cudaStream_t stream, Tin* input, Tout* output); - } // namespace cuda } // namespace onnxruntime diff --git a/orttraining/orttraining/training_ops/cuda/reduction/reduction_all.cu b/orttraining/orttraining/training_ops/cuda/reduction/reduction_all_impl.cu similarity index 98% rename from orttraining/orttraining/training_ops/cuda/reduction/reduction_all.cu rename to orttraining/orttraining/training_ops/cuda/reduction/reduction_all_impl.cu index 7ac6a2cdbd..36fb371b4d 100644 --- a/orttraining/orttraining/training_ops/cuda/reduction/reduction_all.cu +++ b/orttraining/orttraining/training_ops/cuda/reduction/reduction_all_impl.cu @@ -1,10 +1,11 @@ // Copyright (c) Microsoft Corporation. All rights reserved. // Licensed under the MIT License. -#include "orttraining/training_ops/cuda/reduction/reduction_all.h" +#include "reduction_all_impl.h" #include "core/providers/cuda/cu_inc/common.cuh" #include "core/providers/cuda/cuda_common.h" #include "core/providers/cuda/atomic/common.cuh" +#include "core/providers/cuda/multi_tensor/common.cuh" #include "core/providers/cuda/reduction/reduction_utils.cuh" #include "core/providers/cuda/shared_inc/accumulation_type.h" diff --git a/orttraining/orttraining/training_ops/cuda/reduction/reduction_all_impl.h b/orttraining/orttraining/training_ops/cuda/reduction/reduction_all_impl.h new file mode 100644 index 0000000000..32224453d6 --- /dev/null +++ b/orttraining/orttraining/training_ops/cuda/reduction/reduction_all_impl.h @@ -0,0 +1,20 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#pragma once + +#include +#include "core/providers/cuda/multi_tensor/common.cuh" + +namespace onnxruntime { +namespace cuda { +template +struct MultiTensorReduceL2 { + void operator()(cudaStream_t stream, ChunkGroup<1> chunk_group, TOut* output); +}; + +template +void ScalarSqrt(cudaStream_t stream, Tin* input, Tout* output); +} +} // namespace onnxruntime + diff --git a/orttraining/orttraining/training_ops/cuda/tensor/gather_nd_grad_impl.h b/orttraining/orttraining/training_ops/cuda/tensor/gather_nd_grad_impl.h index e00a3ed410..19dd3be218 100644 --- a/orttraining/orttraining/training_ops/cuda/tensor/gather_nd_grad_impl.h +++ b/orttraining/orttraining/training_ops/cuda/tensor/gather_nd_grad_impl.h @@ -2,7 +2,9 @@ // Licensed under the MIT License. #pragma once -#include "core/providers/cuda/shared_inc/cuda_utils.h" + +#include +#include namespace onnxruntime { namespace cuda { diff --git a/orttraining/orttraining/training_ops/rocm/math/softmax_grad.cc b/orttraining/orttraining/training_ops/rocm/math/softmax_grad.cc index 66e1376765..f7851b5adc 100644 --- a/orttraining/orttraining/training_ops/rocm/math/softmax_grad.cc +++ b/orttraining/orttraining/training_ops/rocm/math/softmax_grad.cc @@ -2,6 +2,7 @@ // Licensed under the MIT License. #include "orttraining/training_ops/rocm/math/softmax_grad.h" +#include "orttraining/training_ops/rocm/math/softmax_grad_impl.h" #include "core/providers/common.h" #include "core/providers/rocm/miopen_common.h" diff --git a/orttraining/orttraining/training_ops/rocm/math/softmax_grad_impl.cu b/orttraining/orttraining/training_ops/rocm/math/softmax_grad_impl.cu index 81a176fc73..b12fe031c3 100644 --- a/orttraining/orttraining/training_ops/rocm/math/softmax_grad_impl.cu +++ b/orttraining/orttraining/training_ops/rocm/math/softmax_grad_impl.cu @@ -19,7 +19,7 @@ // The code below is mostly copied from Pytorch PersistentSoftmax.cuh #include "hip/hip_runtime.h" -#include "orttraining/training_ops/rocm/math/softmax_grad.h" +#include "orttraining/training_ops/rocm/math/softmax_grad_impl.h" #include "core/providers/rocm/cu_inc/common.cuh" #include "core/providers/rocm/math/softmax_warpwise_impl.cuh" diff --git a/orttraining/orttraining/training_ops/rocm/reduction/reduction_all.cc b/orttraining/orttraining/training_ops/rocm/reduction/reduction_all.cc index 1577dafe9d..0543beb957 100644 --- a/orttraining/orttraining/training_ops/rocm/reduction/reduction_all.cc +++ b/orttraining/orttraining/training_ops/rocm/reduction/reduction_all.cc @@ -2,6 +2,7 @@ // Licensed under the MIT License. #include "orttraining/training_ops/rocm/reduction/reduction_all.h" +#include "orttraining/training_ops/rocm/reduction/reduction_all_impl.h" #include "core/providers/rocm/reduction/reduction_functions.h" #include "core/providers/rocm/shared_inc/accumulation_type.h"