diff --git a/onnxruntime/core/providers/cuda/math/variadic_elementwise_ops.cc b/onnxruntime/core/providers/cuda/math/variadic_elementwise_ops.cc index 3404dd9c95..419743ecf7 100644 --- a/onnxruntime/core/providers/cuda/math/variadic_elementwise_ops.cc +++ b/onnxruntime/core/providers/cuda/math/variadic_elementwise_ops.cc @@ -235,46 +235,37 @@ using MinOp = VariadicElementwiseOp; - -const DeleteOnUnloadPtr> k_uzilhfd_datatypes = new std::vector( - BuildKernelDefConstraints()); -const DeleteOnUnloadPtr> k_hfd_datatypes = - new std::vector(BuildKernelDefConstraints()); - } // namespace // kernel registration -#define REGISTER_KERNEL(name, impl_class, version, datatypes) \ - ONNX_OPERATOR_KERNEL_EX( \ - name, \ - kOnnxDomain, \ - version, \ - kCudaExecutionProvider, \ - (*KernelDefBuilder::Create()).TypeConstraint("T", datatypes), \ - impl_class) +#define REGISTER_KERNEL(name, impl_class, version, datatypes) \ + ONNX_OPERATOR_KERNEL_EX(name, kOnnxDomain, version, kCudaExecutionProvider, \ + (*KernelDefBuilder::Create()).TypeConstraint("T", BuildKernelDefConstraints()), \ + impl_class) #define REGISTER_VERSIONED_KERNEL(name, impl_class, start_version, end_version, datatypes) \ ONNX_OPERATOR_VERSIONED_KERNEL_EX( \ - name, \ - kOnnxDomain, \ - start_version, end_version, \ - kCudaExecutionProvider, \ - (*KernelDefBuilder::Create()).TypeConstraint("T", datatypes), \ - impl_class) + name, kOnnxDomain, start_version, end_version, kCudaExecutionProvider, \ + (*KernelDefBuilder::Create()).TypeConstraint("T", BuildKernelDefConstraints()), impl_class) -REGISTER_KERNEL(Sum, SumOp, 13, *k_hfd_datatypes) -REGISTER_VERSIONED_KERNEL(Sum, SumOp, 8, 12, *k_hfd_datatypes) -REGISTER_VERSIONED_KERNEL(Sum, SumOp, 6, 7, *k_hfd_datatypes) +#define UZILHFD_TYPES uint32_t, uint64_t, int32_t, int64_t, MLFloat16, float, double, BFloat16 +#define HFD_TYPES MLFloat16, float, double, BFloat16 -REGISTER_KERNEL(Min, MinOp, 13, *k_uzilhfd_datatypes) -REGISTER_VERSIONED_KERNEL(Min, MinOp, 12, 12, *k_uzilhfd_datatypes) -REGISTER_VERSIONED_KERNEL(Min, MinOp, 6, 11, *k_hfd_datatypes) +REGISTER_KERNEL(Sum, SumOp, 13, HFD_TYPES) +REGISTER_VERSIONED_KERNEL(Sum, SumOp, 8, 12, HFD_TYPES) +REGISTER_VERSIONED_KERNEL(Sum, SumOp, 6, 7, HFD_TYPES) -REGISTER_KERNEL(Max, MaxOp, 13, *k_uzilhfd_datatypes) -REGISTER_VERSIONED_KERNEL(Max, MaxOp, 12, 12, *k_uzilhfd_datatypes) -REGISTER_VERSIONED_KERNEL(Max, MaxOp, 6, 11, *k_hfd_datatypes) +REGISTER_KERNEL(Min, MinOp, 13, UZILHFD_TYPES) +REGISTER_VERSIONED_KERNEL(Min, MinOp, 12, 12, UZILHFD_TYPES) +REGISTER_VERSIONED_KERNEL(Min, MinOp, 6, 11, HFD_TYPES) +REGISTER_KERNEL(Max, MaxOp, 13, UZILHFD_TYPES) +REGISTER_VERSIONED_KERNEL(Max, MaxOp, 12, 12, UZILHFD_TYPES) +REGISTER_VERSIONED_KERNEL(Max, MaxOp, 6, 11, HFD_TYPES) + +#undef HFD_TYPES +#undef UZILHFD_TYPES #undef REGISTER_VERSIONED_KERNEL #undef REGISTER_KERNEL diff --git a/orttraining/orttraining/training_ops/cuda/nn/conv_grad.cc b/orttraining/orttraining/training_ops/cuda/nn/conv_grad.cc index 1cd8fdbe81..1d00c222c9 100644 --- a/orttraining/orttraining/training_ops/cuda/nn/conv_grad.cc +++ b/orttraining/orttraining/training_ops/cuda/nn/conv_grad.cc @@ -140,7 +140,7 @@ struct AlgoSearch { static constexpr int num_algos = CUDNN_CONVOLUTION_BWD_DATA_ALGO_COUNT; ORT_ENFORCE(sizeof(algos) / sizeof(algos[0]) == num_algos, "Missing cuDNN convolution backward data algorithms."); int perf_count; - std::unique_ptr candidates(new T_BwdDataPerf[num_algos]); + std::unique_ptr candidates = std::make_unique(num_algos); if (args.params.algo_mode == OrtCudnnConvAlgoSearchHeuristic) { CUDNN_RETURN_IF_ERROR(cudnnGetConvolutionBackwardDataAlgorithm_v7(args.handle, args.w_desc, args.y_tensor, args.conv_desc, args.x_tensor, num_algos, @@ -180,7 +180,7 @@ struct AlgoSearch { // NOTE: - 1 because ALGO_WINOGRAD is not implemented. static constexpr int num_algos = CUDNN_CONVOLUTION_BWD_FILTER_ALGO_COUNT - 1; ORT_ENFORCE(sizeof(algos) / sizeof(algos[0]) == num_algos, "Missing cuDNN convolution backward filter algorithms."); - std::unique_ptr candidates(new T_BwdFilterPerf[num_algos]); + std::unique_ptr candidates = std::make_unique(num_algos); int perf_count; if (args.params.algo_mode == OrtCudnnConvAlgoSearchHeuristic) { CUDNN_RETURN_IF_ERROR(cudnnGetConvolutionBackwardFilterAlgorithm_v7(args.handle, args.x_tensor, args.y_tensor, @@ -222,7 +222,7 @@ class AlgoIterator { Status TryAll(const CUDAExecutionProvider* provider, std::function f) { auto& cache = AlgoSearch::Cache(); - + if (T_Perf algo_perf; cache.Find(args_.params, &algo_perf) && f(algo_perf) == Status::OK()) { return Status::OK(); } diff --git a/orttraining/orttraining/training_ops/rocm/nn/conv_grad.cc b/orttraining/orttraining/training_ops/rocm/nn/conv_grad.cc index eecdf61e17..317cb1b0e7 100644 --- a/orttraining/orttraining/training_ops/rocm/nn/conv_grad.cc +++ b/orttraining/orttraining/training_ops/rocm/nn/conv_grad.cc @@ -140,7 +140,7 @@ struct AlgoSearch { static constexpr int num_algos = MIOPEN_CONVOLUTION_BWD_DATA_ALGO_COUNT; ORT_ENFORCE(sizeof(algos) / sizeof(algos[0]) == num_algos, "Missing MIOpen convolution backward data algorithms."); int perf_count; - std::unique_ptr candidates(new T_BwdDataPerf[num_algos]); + std::unique_ptr candidates = std::make_unique(num_algos); size_t max_workspace_size = provider->GetMiopenConvUseMaxWorkspace() ? GetMaxWorkspaceSize(args, algos, num_algos) : AlgoSearchWorkspaceSize; // Use GetTransientScratchBuffer() so the workspace can be freed instead of cached. @@ -169,7 +169,7 @@ struct AlgoSearch { static constexpr int num_algos = MIOPEN_CONVOLUTION_BWD_FILTER_ALGO_COUNT; ORT_ENFORCE(sizeof(algos) / sizeof(algos[0]) == num_algos, "Missing MIOpen convolution backward filter algorithms."); - std::unique_ptr candidates(new T_BwdFilterPerf[num_algos]); + std::unique_ptr candidates = std::make_unique(num_algos); int perf_count; size_t max_workspace_size = provider->GetMiopenConvUseMaxWorkspace() ? GetMaxWorkspaceSize(args, algos, num_algos) : AlgoSearchWorkspaceSize;