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

More build options tested, converted the training ops over.

Browse source
This commit is contained in:
Ryan Hill 2021-04-12 14:02:08 -07:00
parent a6f7b58292
commit a841d17d06
214 changed files with 2532 additions and 2133 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

4
cmake/onnxruntime_framework.cmake
View file

@ -37,6 +37,10 @@ target_include_directories(onnxruntime_framework PRIVATE ${ONNXRUNTIME_ROOT} ${e
else()
target_include_directories(onnxruntime_framework PRIVATE ${ONNXRUNTIME_ROOT} ${eigen_INCLUDE_DIRS} PUBLIC ${CMAKE_CURRENT_BINARY_DIR})
endif()
# Needed for the provider interface, as it includes training headers when training is enabled
if (onnxruntime_ENABLE_TRAINING OR onnxruntime_ENABLE_TRAINING_OPS)
target_include_directories(onnxruntime_framework PRIVATE ${ORTTRAINING_ROOT})
endif()
onnxruntime_add_include_to_target(onnxruntime_framework onnxruntime_common onnx onnx_proto protobuf::libprotobuf flatbuffers)
set_target_properties(onnxruntime_framework PROPERTIES FOLDER "ONNXRuntime")
# need onnx to build to create headers that this project includes

21
cmake/onnxruntime_providers.cmake
View file

@ -218,6 +218,8 @@ if (onnxruntime_USE_CUDA)
file(GLOB_RECURSE onnxruntime_providers_cuda_cc_srcs CONFIGURE_DEPENDS
"${ONNXRUNTIME_ROOT}/core/providers/cuda/*.h"
"${ONNXRUNTIME_ROOT}/core/providers/cuda/*.cc"
)
file(GLOB_RECURSE onnxruntime_providers_cuda_shared_srcs CONFIGURE_DEPENDS
"${ONNXRUNTIME_ROOT}/core/providers/shared_library/*.h"
"${ONNXRUNTIME_ROOT}/core/providers/shared_library/*.cc"
)
@ -226,8 +228,8 @@ if (onnxruntime_USE_CUDA)
"${ONNXRUNTIME_ROOT}/core/providers/cuda/*.cuh"
)
source_group(TREE ${ONNXRUNTIME_ROOT}/core FILES ${onnxruntime_providers_cuda_cc_srcs} ${onnxruntime_providers_cuda_cu_srcs})
set(onnxruntime_providers_cuda_src ${onnxruntime_providers_cuda_cc_srcs} ${onnxruntime_providers_cuda_cu_srcs})
source_group(TREE ${ONNXRUNTIME_ROOT}/core FILES ${onnxruntime_providers_cuda_cc_srcs} ${onnxruntime_providers_cuda_shared_srcs} ${onnxruntime_providers_cuda_cu_srcs})
set(onnxruntime_providers_cuda_src ${onnxruntime_providers_cuda_cc_srcs} ${onnxruntime_providers_cuda_shared_srcs} ${onnxruntime_providers_cuda_cu_srcs})
# disable contrib ops conditionally
if(NOT onnxruntime_DISABLE_CONTRIB_OPS)
@ -420,6 +422,11 @@ if (onnxruntime_USE_DNNL)
target_compile_definitions(onnxruntime_providers_dnnl PRIVATE USE_DNNL_GPU_OCL=1)
endif()
# Needed for the provider interface, as it includes training headers when training is enabled
if (onnxruntime_ENABLE_TRAINING OR onnxruntime_ENABLE_TRAINING_OPS)
target_include_directories(onnxruntime_providers_dnnl PRIVATE ${ORTTRAINING_ROOT})
endif()
if(APPLE)
set_property(TARGET onnxruntime_providers_dnnl APPEND_STRING PROPERTY LINK_FLAGS "-Xlinker -exported_symbols_list ${ONNXRUNTIME_ROOT}/core/providers/dnnl/exported_symbols.lst")
target_link_libraries(onnxruntime_providers_dnnl PRIVATE nsync_cpp)
@ -507,6 +514,11 @@ if (onnxruntime_USE_TENSORRT)
target_compile_options(onnxruntime_providers_tensorrt INTERFACE /wd4267 /wd4244 /wd4996 /wd4456)
endif()
# Needed for the provider interface, as it includes training headers when training is enabled
if (onnxruntime_ENABLE_TRAINING OR onnxruntime_ENABLE_TRAINING_OPS)
target_include_directories(onnxruntime_providers_tensorrt PRIVATE ${ORTTRAINING_ROOT})
endif()
if(APPLE)
set_property(TARGET onnxruntime_providers_tensorrt APPEND_STRING PROPERTY LINK_FLAGS "-Xlinker -exported_symbols_list ${ONNXRUNTIME_ROOT}/core/providers/tensorrt/exported_symbols.lst")
target_link_libraries(onnxruntime_providers_tensorrt PRIVATE nsync_cpp)
@ -615,6 +627,11 @@ if (onnxruntime_USE_OPENVINO)
target_compile_options(onnxruntime_providers_openvino PUBLIC /wd4275 /wd4100 /wd4005 /wd4244 /wd4267)
endif()
# Needed for the provider interface, as it includes training headers when training is enabled
if (onnxruntime_ENABLE_TRAINING OR onnxruntime_ENABLE_TRAINING_OPS)
target_include_directories(onnxruntime_providers_openvino PRIVATE ${ORTTRAINING_ROOT})
endif()
if(APPLE)
set_property(TARGET onnxruntime_providers_openvino APPEND_STRING PROPERTY LINK_FLAGS "-Xlinker -exported_symbols_list ${ONNXRUNTIME_ROOT}/core/providers/openvino/exported_symbols.lst")
elseif(UNIX)

10
cmake/onnxruntime_python.cmake
View file

@ -404,6 +404,16 @@ if (onnxruntime_USE_OPENVINO)
)
endif()
if (onnxruntime_USE_CUDA)
add_custom_command(
TARGET onnxruntime_pybind11_state POST_BUILD
COMMAND ${CMAKE_COMMAND} -E copy
$<TARGET_FILE:onnxruntime_providers_cuda>
$<TARGET_FILE:onnxruntime_providers_shared>
$<TARGET_FILE_DIR:${build_output_target}>/onnxruntime/capi/
)
endif()
if (onnxruntime_USE_TVM)
add_custom_command(
TARGET onnxruntime_pybind11_state POST_BUILD

5
cmake/onnxruntime_unittests.cmake
View file

@ -382,7 +382,7 @@ endif()
set (onnxruntime_test_providers_dependencies ${onnxruntime_EXTERNAL_DEPENDENCIES})
if(onnxruntime_USE_CUDA)
list(APPEND onnxruntime_test_providers_dependencies onnxruntime_providers_cuda)
list(APPEND onnxruntime_test_providers_dependencies onnxruntime_providers_cuda onnxruntime_providers_shared)
endif()
if(onnxruntime_USE_DNNL)
@ -451,8 +451,7 @@ set(ONNXRUNTIME_TEST_LIBS
onnxruntime_session
${ONNXRUNTIME_INTEROP_TEST_LIBS}
${onnxruntime_libs}
${PROVIDERS_CUDA}
# TENSORRT, DNNL, and OpenVINO are explicitly linked at runtime
# CUDA, TENSORRT, DNNL, and OpenVINO are explicitly linked at runtime
${PROVIDERS_MIGRAPHX}
${PROVIDERS_NUPHAR}
${PROVIDERS_NNAPI}

7
include/onnxruntime/core/framework/kernel_def_builder.h
View file

@ -285,6 +285,13 @@ class KernelDefBuilder {
return *this;
}
KernelDefBuilder& InputMemoryType(OrtMemType type, const std::vector<int>& input_indexes) {
for (auto input_index : input_indexes) {
kernel_def_->input_memory_type_args_.insert(std::make_pair(input_index, type));
}
return *this;
}
/**
Specify that this kernel provides an output arg
in certain memory type (instead of the default, device memory).

2
include/onnxruntime/core/framework/ml_value.h
View file

@ -4,11 +4,13 @@
#pragma once
#include <string>
#ifndef SHARED_PROVIDER
#include "core/common/common.h"
#include "core/common/exceptions.h"
#include "core/framework/allocator.h"
#include "core/framework/data_types.h"
#include "core/framework/tensor.h"
#endif
namespace onnxruntime {
class SparseTensor;

23
include/onnxruntime/core/providers/cuda/cuda_provider_factory.h
View file

@ -4,6 +4,29 @@
#include "onnxruntime_c_api.h"
#ifdef __cplusplus
namespace onnxruntime {
class IAllocator;
class IDataTransfer;
} // namespace onnxruntime
struct ProviderInfo_CUDA {
virtual OrtStatus* SetCurrentGpuDeviceId(_In_ int device_id) = 0;
virtual OrtStatus* GetCurrentGpuDeviceId(_In_ int* device_id) = 0;
virtual std::unique_ptr<onnxruntime::IAllocator> CreateCUDAAllocator(int16_t device_id, const char* name) = 0;
virtual std::unique_ptr<onnxruntime::IAllocator> CreateCUDAPinnedAllocator(int16_t device_id, const char* name) = 0;
virtual std::unique_ptr<onnxruntime::IDataTransfer> CreateGPUDataTransfer(void* stream) = 0;
virtual void cuda__Impl_Cast(void* stream, const int64_t* input_data, int32_t* output_data, size_t count) = 0;
virtual void cuda__Impl_Cast(void* stream, const int32_t* input_data, int64_t* output_data, size_t count) = 0;
virtual bool CudaCall_false(int retCode, const char* exprString, const char* libName, int successCode, const char* msg) = 0;
virtual bool CudaCall_true(int retCode, const char* exprString, const char* libName, int successCode, const char* msg) = 0;
virtual void CopyGpuToCpu(void* dst_ptr, const void* src_ptr, const size_t size, const OrtMemoryInfo& dst_location, const OrtMemoryInfo& src_location) = 0;
virtual void cudaMemcpy_HostToDevice(void* dst, const void* src, size_t count) = 0;
};
extern "C" {
#endif

8
onnxruntime/contrib_ops/cpu/bert/attention.cc
View file

@ -40,14 +40,6 @@ ONNX_OPERATOR_TYPED_KERNEL_EX(
.TypeConstraint("T", DataTypeImpl::GetTensorType<float>()),
Attention<float>);
AttentionBase::AttentionBase(const OpKernelInfo& info) {
int64_t num_heads = 0;
ORT_ENFORCE(info.GetAttr("num_heads", &num_heads).IsOK() && num_heads > 0);
num_heads_ = static_cast<int>(num_heads);
is_unidirectional_ = info.GetAttrOrDefault<int64_t>("unidirectional", 0) == 1;
}
Status AttentionBase::CheckInputs(const TensorShape& input_shape,
const TensorShape& weights_shape,
const TensorShape& bias_shape,

25
onnxruntime/contrib_ops/cpu/bert/attention_base.h
View file

@ -10,15 +10,7 @@ namespace onnxruntime {
namespace contrib {
class AttentionBase {
protected:
AttentionBase(const OpKernelInfo& info);
Status CheckInputs(const TensorShape& input_shape,
const TensorShape& weights_shape,
const TensorShape& bias_shape,
const Tensor*& mask_index, // For dummy mask with shape (1, 1) or (batch_size, 1), it will be updated to nullptr.
const Tensor* past) const;
public:
// This check function is specifically used in cuda
Status CheckInputs(const TensorShape& input_shape,
const TensorShape& weights_shape,
@ -34,6 +26,21 @@ class AttentionBase {
int sequence_length,
int& past_sequence_length) const;
protected:
AttentionBase(const OpKernelInfo& info) {
int64_t num_heads = 0;
ORT_ENFORCE(info.GetAttr("num_heads", &num_heads).IsOK() && num_heads > 0);
num_heads_ = static_cast<int>(num_heads);
is_unidirectional_ = info.GetAttrOrDefault<int64_t>("unidirectional", 0) == 1;
}
Status CheckInputs(const TensorShape& input_shape,
const TensorShape& weights_shape,
const TensorShape& bias_shape,
const Tensor*& mask_index, // For dummy mask with shape (1, 1) or (batch_size, 1), it will be updated to nullptr.
const Tensor* past) const;
int num_heads_; // number of attention heads
bool is_unidirectional_; // whether every token can only attend to previous tokens.
};

10
onnxruntime/contrib_ops/cpu/bert/longformer_attention_base.cc
View file

@ -6,16 +6,6 @@
namespace onnxruntime {
namespace contrib {
LongformerAttentionBase::LongformerAttentionBase(const OpKernelInfo& info) {
int64_t num_heads = 0;
ORT_ENFORCE(info.GetAttr("num_heads", &num_heads).IsOK() && num_heads > 0);
num_heads_ = static_cast<int>(num_heads);
int64_t window = 0;
ORT_ENFORCE(info.GetAttr("window", &window).IsOK() && window > 0);
window_ = static_cast<int>(window);
}
Status LongformerAttentionBase::CheckInputs(const TensorShape& input_shape,
const TensorShape& weights_shape,
const TensorShape& bias_shape,

15
onnxruntime/contrib_ops/cpu/bert/longformer_attention_base.h
View file

@ -10,9 +10,7 @@ namespace onnxruntime {
namespace contrib {
class LongformerAttentionBase {
protected:
LongformerAttentionBase(const OpKernelInfo& info);
public:
Status CheckInputs(const TensorShape& input_shape,
const TensorShape& weights_shape,
const TensorShape& bias_shape,
@ -21,6 +19,17 @@ class LongformerAttentionBase {
const TensorShape& global_bias_shape,
const TensorShape& global_shape) const;
protected:
LongformerAttentionBase(const OpKernelInfo& info) {
int64_t num_heads = 0;
ORT_ENFORCE(info.GetAttr("num_heads", &num_heads).IsOK() && num_heads > 0);
num_heads_ = static_cast<int>(num_heads);
int64_t window = 0;
ORT_ENFORCE(info.GetAttr("window", &window).IsOK() && window > 0);
window_ = static_cast<int>(window);
}
int num_heads_; // Number of attention heads
int window_; // Attention windows length (W). It is half (one-sided) of total window size.
};

1
onnxruntime/contrib_ops/cuda/bert/attention.cc
View file

@ -1,7 +1,6 @@
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
//#include "core/framework/tensorprotoutils.h"
#include "core/providers/cuda/cuda_common.h"
#include "core/providers/cuda/shared_inc/fpgeneric.h"
#include "attention.h"

1
onnxruntime/contrib_ops/cuda/bert/attention.h
View file

@ -3,7 +3,6 @@
#pragma once
//#include "core/common/common.h"
#include "core/providers/cuda/cuda_kernel.h"
#include "contrib_ops/cpu/bert/attention_base.h"

3
onnxruntime/contrib_ops/cuda/bert/embed_layer_norm.cc
View file

@ -1,9 +1,6 @@
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
//#include "core/providers/common.h"
//#include "core/framework/tensorprotoutils.h"
//#include "onnx/defs/tensor_proto_util.h"
#include "core/providers/cuda/cuda_common.h"
#include "contrib_ops/cpu/bert/embed_layer_norm_helper.h"
#include "embed_layer_norm.h"

1
onnxruntime/contrib_ops/cuda/bert/embed_layer_norm.h
View file

@ -2,7 +2,6 @@
// Licensed under the MIT License.
#pragma once
//#include "core/common/common.h"
#include "core/providers/cuda/cuda_kernel.h"
namespace onnxruntime {

1
onnxruntime/contrib_ops/cuda/bert/fast_gelu.cc
View file

@ -3,7 +3,6 @@
#include "core/providers/cuda/cuda_common.h"
#include "core/providers/cuda/cudnn_common.h"
//#include "core/framework/tensorprotoutils.h"
#include "fast_gelu.h"
#include "fast_gelu_impl.h"
#include "contrib_ops/cpu/bert/bias_gelu_helper.h"

1
onnxruntime/contrib_ops/cuda/bert/longformer_attention.cc
View file

@ -1,7 +1,6 @@
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
//#include "core/framework/tensorprotoutils.h"
#include "core/providers/cuda/cuda_common.h"
#include "core/providers/cuda/shared_inc/fpgeneric.h"
#include "core/platform/env_var_utils.h"

3
onnxruntime/contrib_ops/cuda/bert/skip_layer_norm.cc
View file

@ -1,10 +1,7 @@
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
//#include "core/providers/common.h"
#include "core/providers/cuda/cuda_common.h"
//#include "core/framework/tensorprotoutils.h"
//#include "onnx/defs/tensor_proto_util.h"
#include "skip_layer_norm.h"
#include "skip_layer_norm_impl.h"

161
onnxruntime/contrib_ops/cuda/cuda_contrib_kernels.cc
View file

@ -3,7 +3,6 @@
#include "core/providers/shared_library/provider_api.h"
#include "core/providers/cuda/cuda_common.h"
//#include "core/framework/kernel_registry.h"
using namespace onnxruntime::common;
@ -18,9 +17,9 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, BiasGelu);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, double, BiasGelu);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, BiasGelu);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, TransposeMatMul); // backward compatibility
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, double, TransposeMatMul); // backward compatibility
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, TransposeMatMul); // backward compatibility
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, TransposeMatMul); // backward compatibility
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, double, TransposeMatMul); // backward compatibility
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, TransposeMatMul); // backward compatibility
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, FusedMatMul);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, double, FusedMatMul);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, FusedMatMul);
@ -84,7 +83,7 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1
#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, BFloat16, FastGelu);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, BFloat16, TransposeMatMul); // backward compatibility
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, BFloat16, TransposeMatMul); // backward compatibility
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, BFloat16, FusedMatMul);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, BFloat16_float, LayerNormalization);
#endif
@ -97,87 +96,87 @@ KernelCreateInfo BuildKernelCreateInfo<void>() {
Status RegisterCudaContribKernels(KernelRegistry& kernel_registry) {
static const BuildKernelCreateInfoFn function_table[] = {
BuildKernelCreateInfo<void>, //default entry to avoid the list become empty after ops-reducing
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, FastGelu)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, FastGelu)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, Gelu)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, double, Gelu)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, Gelu)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, BiasGelu)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, double, BiasGelu)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, BiasGelu)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, TransposeMatMul)>, // backward compatibility
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, double, TransposeMatMul)>, // backward compatibility
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, TransposeMatMul)>, // backward compatibility
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, FusedMatMul)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, double, FusedMatMul)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, FusedMatMul)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, Rfft)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, double, Rfft)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, Rfft)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, Irfft)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, double, Irfft)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, Irfft)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, ComplexMul)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, ComplexMul)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, ComplexMulConj)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, ComplexMulConj)>,
BuildKernelCreateInfo<void>, //default entry to avoid the list become empty after ops-reducing
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, FastGelu)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, FastGelu)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, Gelu)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, double, Gelu)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, Gelu)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, BiasGelu)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, double, BiasGelu)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, BiasGelu)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, TransposeMatMul)>, // backward compatibility
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, double, TransposeMatMul)>, // backward compatibility
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, TransposeMatMul)>, // backward compatibility
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, FusedMatMul)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, double, FusedMatMul)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, FusedMatMul)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, Rfft)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, double, Rfft)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, Rfft)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, Irfft)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, double, Irfft)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, Irfft)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, ComplexMul)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, ComplexMul)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, ComplexMulConj)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, ComplexMulConj)>,
// These ops were experimental ops in onnx domain which have been removed now. We add them here as
// contrib ops to maintain backward compatibility
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, float, Affine)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, double, Affine)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, MLFloat16, Affine)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, Attention)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, Attention)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, ConvTransposeWithDynamicPads)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, float, Crop)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, double, Crop)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, MLFloat16, Crop)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, int32_t, DynamicSlice)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, int64_t, DynamicSlice)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, EmbedLayerNormalization)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, EmbedLayerNormalization)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, float, ImageScaler)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, double, ImageScaler)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, MLFloat16, ImageScaler)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, LongformerAttention)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, LongformerAttention)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, float, ParametricSoftplus)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, double, ParametricSoftplus)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, MLFloat16, ParametricSoftplus)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, float, ScaledTanh)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, double, ScaledTanh)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, MLFloat16, ScaledTanh)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, SkipLayerNormalization)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, SkipLayerNormalization)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, float, ThresholdedRelu)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, double, ThresholdedRelu)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, MLFloat16, ThresholdedRelu)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, float_float, LayerNormalization)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, double_double, LayerNormalization)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, MLFloat16_float, LayerNormalization)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, float_float, SimplifiedLayerNormalization)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, double_double, SimplifiedLayerNormalization)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, MLFloat16_float, SimplifiedLayerNormalization)>,
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, Inverse)>,
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, BiasSoftmax)>,
// These ops were experimental ops in onnx domain which have been removed now. We add them here as
// contrib ops to maintain backward compatibility
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, float, Affine)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, double, Affine)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, MLFloat16, Affine)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, Attention)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, Attention)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, ConvTransposeWithDynamicPads)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, float, Crop)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, double, Crop)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, MLFloat16, Crop)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, int32_t, DynamicSlice)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, int64_t, DynamicSlice)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, EmbedLayerNormalization)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, EmbedLayerNormalization)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, float, ImageScaler)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, double, ImageScaler)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, MLFloat16, ImageScaler)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, LongformerAttention)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, LongformerAttention)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, float, ParametricSoftplus)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, double, ParametricSoftplus)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, MLFloat16, ParametricSoftplus)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, float, ScaledTanh)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, double, ScaledTanh)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, MLFloat16, ScaledTanh)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, SkipLayerNormalization)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16, SkipLayerNormalization)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, float, ThresholdedRelu)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, double, ThresholdedRelu)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, MLFloat16, ThresholdedRelu)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, float_float, LayerNormalization)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, double_double, LayerNormalization)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, MLFloat16_float, LayerNormalization)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, float_float, SimplifiedLayerNormalization)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, double_double, SimplifiedLayerNormalization)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, MLFloat16_float, SimplifiedLayerNormalization)>,
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, Inverse)>,
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, BiasSoftmax)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, int8_t_MLFloat16, QuantizeLinear)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, uint8_t_MLFloat16, QuantizeLinear)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, int8_t_MLFloat16, DequantizeLinear)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, uint8_t_MLFloat16, DequantizeLinear)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float_int8_t, QAttention)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16_int8_t, QAttention)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, int8_t_MLFloat16, QuantizeLinear)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, uint8_t_MLFloat16, QuantizeLinear)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, int8_t_MLFloat16, DequantizeLinear)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, uint8_t_MLFloat16, DequantizeLinear)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float_int8_t, QAttention)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, MLFloat16_int8_t, QAttention)>,
#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, BFloat16, FastGelu)>,
// TransposedMatMul is still here for backward compatibility
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, BFloat16, TransposeMatMul)>, // backward compatibility
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, BFloat16, FusedMatMul)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, BFloat16_float, LayerNormalization)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, BFloat16, FastGelu)>,
// TransposedMatMul is still here for backward compatibility
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, BFloat16, TransposeMatMul)>, // backward compatibility
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, BFloat16, FusedMatMul)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, BFloat16_float, LayerNormalization)>,
#endif
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, FusedConv)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kMSDomain, 1, float, FusedConv)>,
};
for (auto& function_table_entry : function_table) {

1
onnxruntime/contrib_ops/cuda/cuda_contrib_kernels.h
View file

@ -2,7 +2,6 @@
// Licensed under the MIT License.
#pragma once
//#include "core/framework/kernel_registry.h"
namespace onnxruntime {
namespace contrib {

4
onnxruntime/contrib_ops/cuda/fused_conv.cc
View file

@ -119,9 +119,9 @@ ONNX_OPERATOR_TYPED_KERNEL_EX(
1,
float,
kCudaExecutionProvider,
KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType<float>()),
(*KernelDefBuilder::Create()).TypeConstraint("T", DataTypeImpl::GetTensorType<float>()),
FusedConv<float>);
} // namespace cuda
} // namespace contrib
} // namespace onnxruntime
} // namespace onnxruntime

2
onnxruntime/contrib_ops/cuda/inverse.cc
View file

@ -28,7 +28,7 @@ ONNX_OPERATOR_KERNEL_EX(
kMSDomain,
1,
kCudaExecutionProvider,
KernelDefBuilder()
(*KernelDefBuilder::Create())
.TypeConstraint("T", BuildKernelDefConstraints<float, double, MLFloat16>()),
Inverse);

8
onnxruntime/contrib_ops/cuda/layer_norm.cc
View file

@ -3,8 +3,6 @@
#include "layer_norm.h"
#include "layer_norm_impl.h"
#include "core/providers/common.h"
#include "core/providers/cuda/cuda_common.h"
namespace onnxruntime {
@ -18,7 +16,7 @@ namespace cuda {
1, \
T##_##U, \
kCudaExecutionProvider, \
KernelDefBuilder() \
(*KernelDefBuilder::Create()) \
.TypeConstraint("T", DataTypeImpl::GetTensorType<T>()) \
.TypeConstraint("U", DataTypeImpl::GetTensorType<U>()), \
LayerNorm<T, U, false>); \
@ -28,7 +26,7 @@ namespace cuda {
1, \
T##_##U, \
kCudaExecutionProvider, \
KernelDefBuilder() \
(*KernelDefBuilder::Create()) \
.TypeConstraint("T", DataTypeImpl::GetTensorType<T>()) \
.TypeConstraint("U", DataTypeImpl::GetTensorType<U>()), \
LayerNorm<T, U, true>);
@ -59,7 +57,7 @@ Status LayerNorm<T, U, simplified>::ComputeInternal(OpKernelContext* ctx) const
auto X_data = reinterpret_cast<const CudaT*>(X->template Data<T>());
auto scale_data = reinterpret_cast<const CudaT*>(scale->template Data<T>());
auto bias_data = (simplified || (nullptr == bias)) ? nullptr: reinterpret_cast<const CudaT*>(bias->template Data<T>());
auto bias_data = (simplified || (nullptr == bias)) ? nullptr : reinterpret_cast<const CudaT*>(bias->template Data<T>());
const TensorShape& x_shape = X->Shape();
const int64_t axis = HandleNegativeAxis(axis_, x_shape.NumDimensions());

1
onnxruntime/contrib_ops/cuda/layer_norm.h
View file

@ -2,7 +2,6 @@
// Licensed under the MIT License.
#pragma once
#include "core/common/common.h"
#include "core/providers/cuda/cuda_kernel.h"
namespace onnxruntime {

1
onnxruntime/contrib_ops/cuda/math/bias_softmax.cc
View file

@ -3,7 +3,6 @@
#include "core/providers/cuda/cuda_common.h"
#include "contrib_ops/cuda/math/bias_softmax.h"
//#include "core/providers/common.h"
using namespace onnxruntime;
using namespace onnxruntime::cuda;

1
onnxruntime/contrib_ops/cuda/math/bias_softmax_impl.cu
View file

@ -6,7 +6,6 @@
#include <limits>
#include <algorithm>
//#include "core/providers/common.h"
#include "core/providers/cuda/cuda_common.h"
#include "core/providers/cuda/cudnn_common.h"
#include "core/providers/cuda/cu_inc/binary_elementwise_impl.cuh"

2
onnxruntime/contrib_ops/cuda/quantization/attention_quantization.cc
View file

@ -4,8 +4,6 @@
#include "attention_quantization.h"
#include "attention_quantization_impl.cuh"
#include "contrib_ops/cuda/bert/attention_impl.h"
//#include "core/framework/tensorprotoutils.h"
//#include "core/providers/common.h"
#include "core/providers/cuda/cuda_common.h"
#include "core/providers/cuda/shared_inc/fpgeneric.h"
#include "core/providers/cuda/shared_inc/integer_gemm.h"

2
onnxruntime/core/framework/feeds_fetches_manager.h
View file

@ -6,7 +6,9 @@
#include <string>
#include <vector>
#ifndef SHARED_PROVIDER
#include "core/framework/ml_value.h"
#endif
namespace onnxruntime {
class ExecutionProviders;

4
onnxruntime/core/framework/ort_value_tensor_slicer.cc
View file

@ -80,8 +80,12 @@ void OrtValueTensorSlicer<T>::Iterator::MaterializeMLValue() const {
//
// TODO: Ideally we could avoid the overhead of creating a new Tensor (mainly cost of copying type and shape info)
// and would simply update Tensor::p_data_ given all other info remains constant for each slice.
#ifndef SHARED_PROVIDER
auto sub_tensor = onnxruntime::make_unique<Tensor>(tensor_data_type_, per_iteration_shape_,
const_cast<void*>(tensor_slice_data_raw), *tensor_location_);
#else
auto sub_tensor = Tensor::Create(tensor_data_type_, per_iteration_shape_, const_cast<void*>(tensor_slice_data_raw), *tensor_location_);
#endif
auto ml_tensor = DataTypeImpl::GetType<Tensor>();
current_ = OrtValue{sub_tensor.release(), ml_tensor, ml_tensor->GetDeleteFunc()};
}

2
onnxruntime/core/framework/ort_value_tensor_slicer.h
View file

@ -7,9 +7,11 @@
#include <limits>
#include <type_traits>
#ifndef SHARED_PROVIDER
#include "core/common/common.h"
#include "core/framework/ml_value.h"
#include "core/framework/tensor.h"
#endif
namespace onnxruntime {

255
onnxruntime/core/framework/provider_bridge_ort.cc
View file

@ -19,12 +19,14 @@ void Foo() {}
#include "core/session/abi_session_options_impl.h"
#include "core/session/ort_apis.h"
#include "core/util/math.h"
#include "core/framework/tensorprotoutils.h"
#include "core/framework/fallback_cpu_capability.h"
#include "core/framework/random_generator.h"
#include "core/providers/cpu/tensor/gatherbase.h"
#include "core/providers/cpu/tensor/gather_elements.h"
#include "core/providers/cpu/tensor/unsqueeze.h"
#include "core/providers/cpu/tensor/slice.h"
#include "core/framework/fallback_cpu_capability.h"
#include "core/providers/cpu/tensor/split.h"
#include "core/providers/cpu/tensor/size.h"
#include "core/providers/cpu/tensor/scatter_nd.h"
@ -32,12 +34,27 @@ void Foo() {}
#include "core/providers/cpu/tensor/concatbase.h"
#include "core/providers/cpu/tensor/gatherbase.h"
#include "core/providers/cpu/tensor/onehot.h"
#include "core/providers/cpu/tensor/tile.h"
#include "core/providers/cpu/controlflow/if.h"
#include "core/providers/cpu/controlflow/loop.h"
#include "core/providers/cpu/controlflow/scan.h"
#include "core/providers/cpu/math/einsum.h"
#include "core/providers/cpu/math/cumsum.h"
#include "core/providers/cpu/object_detection/non_max_suppression.h"
#ifdef USE_TENSORRT
#include "core/providers/cuda/cuda_allocator.h"
#include "core/providers/cuda/gpu_data_transfer.h"
#include "core/providers/cuda/math/unary_elementwise_ops_impl.h"
#include "core/providers/cuda/cuda_common.h"
#ifndef DISABLE_CONTRIB_OPS
#include "contrib_ops/cpu/bert/bias_gelu_helper.h"
#include "contrib_ops/cpu/bert/embed_layer_norm_helper.h"
#include "contrib_ops/cpu/bert/longformer_attention_base.h"
#include "contrib_ops/cpu/bert/attention_base.h"
#endif
#ifdef ENABLE_TRAINING
#include "orttraining/training_ops/cpu/controlflow/group.h"
#include "orttraining/training_ops/cpu/controlflow/record.h"
#include "orttraining/training_ops/cpu/controlflow/wait.h"
#include "orttraining/training_ops/cpu/loss/softmax_cross_entropy_loss.h"
#include "orttraining/training_ops/cpu/tensor/split.h"
#endif
namespace ONNX_NAMESPACE {
@ -82,6 +99,8 @@ namespace onnxruntime {
ProviderHost* g_host{};
ProviderInfo_CUDA* GetProviderInfo_CUDA();
struct TensorShapeProto_Dimension_Iterator_Impl : TensorShapeProto_Dimension_Iterator {
TensorShapeProto_Dimension_Iterator_Impl(google::protobuf::internal::RepeatedPtrIterator<const onnx::TensorShapeProto_Dimension>&& v) : v_{std::move(v)} {}
@ -100,8 +119,9 @@ struct NodeAttributes_Iterator_Impl : NodeAttributes_Iterator {
void operator++() override { v_.operator++(); }
const std::pair<const std::string, ONNX_NAMESPACE::AttributeProto>& operator*() const override { return v_.operator*(); }
// const std::string& first() const override { return v_->first; }
// const Provider_AttributeProto& second() const override { return v_->second; }
const std::string& first() const override { return v_->first; }
const ONNX_NAMESPACE::AttributeProto& second() const override { return v_->second; }
NodeAttributes::const_iterator v_;
};
@ -145,6 +165,7 @@ struct ProviderHostImpl : ProviderHost {
DataTypeImpl__GetTensorType_uint64 = &DataTypeImpl::GetTensorType<uint64_t>;
DataTypeImpl__GetTensorType_float = &DataTypeImpl::GetTensorType<float>;
DataTypeImpl__GetTensorType_double = &DataTypeImpl::GetTensorType<double>;
DataTypeImpl__GetTensorType_BFloat16 = &DataTypeImpl::GetTensorType<BFloat16>;
DataTypeImpl__GetTensorType_MLFloat16 = &DataTypeImpl::GetTensorType<MLFloat16>;
}
@ -156,29 +177,19 @@ struct ProviderHostImpl : ProviderHost {
return onnxruntime::make_unique<CPUAllocator>(memory_info);
};
#ifdef USE_TENSORRT
std::unique_ptr<IAllocator> CreateCUDAAllocator(int16_t device_id, const char* name) override {
return onnxruntime::make_unique<CUDAAllocator>(device_id, name);
}
void* CPUAllocator__Alloc(CPUAllocator* p, size_t size) override { return p->Alloc(size); }
void CPUAllocator__Free(CPUAllocator* p, void* allocation) override { return p->Free(allocation); }
std::unique_ptr<IAllocator> CreateCUDAPinnedAllocator(int16_t device_id, const char* name) override {
return onnxruntime::make_unique<CUDAPinnedAllocator>(device_id, name);
}
#ifdef USE_CUDA
std::unique_ptr<IAllocator> CreateCUDAAllocator(int16_t device_id, const char* name) override { return GetProviderInfo_CUDA()->CreateCUDAAllocator(device_id, name); }
std::unique_ptr<IAllocator> CreateCUDAPinnedAllocator(int16_t device_id, const char* name) override { return GetProviderInfo_CUDA()->CreateCUDAPinnedAllocator(device_id, name); }
std::unique_ptr<IDataTransfer> CreateGPUDataTransfer(void* stream) override { return GetProviderInfo_CUDA()->CreateGPUDataTransfer(stream); }
std::unique_ptr<IDataTransfer> CreateGPUDataTransfer(void* stream) override {
return onnxruntime::make_unique<GPUDataTransfer>(static_cast<cudaStream_t>(stream));
}
void cuda__Impl_Cast(void* stream, const int64_t* input_data, int32_t* output_data, size_t count) override { return GetProviderInfo_CUDA()->cuda__Impl_Cast(stream, input_data, output_data, count); }
void cuda__Impl_Cast(void* stream, const int32_t* input_data, int64_t* output_data, size_t count) override { return GetProviderInfo_CUDA()->cuda__Impl_Cast(stream, input_data, output_data, count); }
void cuda__Impl_Cast(void* stream, const int64_t* input_data, int32_t* output_data, size_t count) override {
return cuda::Impl_Cast(static_cast<cudaStream_t>(stream), input_data, output_data, count);
}
void cuda__Impl_Cast(void* stream, const int32_t* input_data, int64_t* output_data, size_t count) override {
return cuda::Impl_Cast(static_cast<cudaStream_t>(stream), input_data, output_data, count);
}
bool CudaCall_false(int retCode, const char* exprString, const char* libName, int successCode, const char* msg) override { return CudaCall<cudaError, false>(cudaError(retCode), exprString, libName, cudaError(successCode), msg); }
bool CudaCall_true(int retCode, const char* exprString, const char* libName, int successCode, const char* msg) override { return CudaCall<cudaError, true>(cudaError(retCode), exprString, libName, cudaError(successCode), msg); }
bool CudaCall_false(int retCode, const char* exprString, const char* libName, int successCode, const char* msg) override { return GetProviderInfo_CUDA()->CudaCall_false(retCode, exprString, libName, successCode, msg); }
bool CudaCall_true(int retCode, const char* exprString, const char* libName, int successCode, const char* msg) override { return GetProviderInfo_CUDA()->CudaCall_true(retCode, exprString, libName, successCode, msg); }
#endif
std::string GetEnvironmentVar(const std::string& var_name) override {
@ -189,6 +200,10 @@ struct ProviderHostImpl : ProviderHost {
return const_cast<logging::Logger*>(&logging::LoggingManager::DefaultLogger());
}
OrtStatus* CreateStatus(OrtErrorCode code, _In_ const char* msg) noexcept override {
return OrtApis::CreateStatus(code, msg);
}
std::unordered_set<NodeIndex> GetCpuPreferredNodes(const onnxruntime::GraphViewer& graph,
const std::string& provider_type,
const std::vector<const KernelRegistry*>& kernel_registries,
@ -199,21 +214,35 @@ struct ProviderHostImpl : ProviderHost {
int32_t PrimitiveDataTypeBase__GetDataType(const PrimitiveDataTypeBase* p) override { return p->GetDataType(); }
const char* DataTypeImpl__ToString(MLDataType type) override { return DataTypeImpl::ToString(type); }
bool DataTypeImpl__IsTensorType(const DataTypeImpl* p) override { return p->IsTensorType(); }
bool DataTypeImpl__IsTensorSequenceType(const DataTypeImpl* p) override { return p->IsTensorSequenceType(); }
bool DataTypeImpl__IsSparseTensorType(const DataTypeImpl* p) override { return p->IsSparseTensorType(); }
DeleteFunc DataTypeImpl__GetDeleteFunc(const DataTypeImpl* p) override { return p->GetDeleteFunc(); }
const std::vector<MLDataType>& DataTypeImpl__AllFixedSizeTensorTypes() override { return DataTypeImpl::AllFixedSizeTensorTypes(); }
const std::vector<MLDataType>& DataTypeImpl__AllTensorTypes() override { return DataTypeImpl::AllTensorTypes(); }
const std::vector<MLDataType>& DataTypeImpl__AllIEEEFloatTensorTypes() override { return DataTypeImpl::AllIEEEFloatTensorTypes(); }
size_t DataTypeImpl__Size(const DataTypeImpl* p) override { return p->Size(); }
const PrimitiveDataTypeBase* DataTypeImpl__AsPrimitiveDataType(const DataTypeImpl* p) override { return p->AsPrimitiveDataType(); }
Status UnpackTensor(const ONNX_NAMESPACE::TensorProto& tensor, const void* raw_data, size_t raw_data_len, /*out*/ bool* p_data, size_t expected_size) override { return utils::UnpackTensor(tensor, raw_data, raw_data_len, p_data, expected_size); }
Status UnpackTensor(const ONNX_NAMESPACE::TensorProto& tensor, const void* raw_data, size_t raw_data_len, /*out*/ float* p_data, size_t expected_size) override { return utils::UnpackTensor(tensor, raw_data, raw_data_len, p_data, expected_size); }
Status UnpackTensor(const ONNX_NAMESPACE::TensorProto& tensor, const void* raw_data, size_t raw_data_len, /*out*/ double* p_data, size_t expected_size) override { return utils::UnpackTensor(tensor, raw_data, raw_data_len, p_data, expected_size); }
Status UnpackTensor(const ONNX_NAMESPACE::TensorProto& tensor, const void* raw_data, size_t raw_data_len, /*out*/ MLFloat16* p_data, size_t expected_size) override { return utils::UnpackTensor(tensor, raw_data, raw_data_len, p_data, expected_size); }
Status UnpackTensor(const ONNX_NAMESPACE::TensorProto& tensor, const void* raw_data, size_t raw_data_len, /*out*/ int8_t* p_data, size_t expected_size) override { return utils::UnpackTensor(tensor, raw_data, raw_data_len, p_data, expected_size); }
Status UnpackTensor(const ONNX_NAMESPACE::TensorProto& tensor, const void* raw_data, size_t raw_data_len, /*out*/ uint8_t* p_data, size_t expected_size) override { return utils::UnpackTensor(tensor, raw_data, raw_data_len, p_data, expected_size); }
Status UnpackTensor(const ONNX_NAMESPACE::TensorProto& tensor, const void* raw_data, size_t raw_data_len, /*out*/ int16_t* p_data, size_t expected_size) override { return utils::UnpackTensor(tensor, raw_data, raw_data_len, p_data, expected_size); }
Status UnpackTensor(const ONNX_NAMESPACE::TensorProto& tensor, const void* raw_data, size_t raw_data_len, /*out*/ uint16_t* p_data, size_t expected_size) override { return utils::UnpackTensor(tensor, raw_data, raw_data_len, p_data, expected_size); }
Status UnpackTensor(const ONNX_NAMESPACE::TensorProto& tensor, const void* raw_data, size_t raw_data_len, /*out*/ int32_t* p_data, size_t expected_size) override { return utils::UnpackTensor(tensor, raw_data, raw_data_len, p_data, expected_size); }
Status UnpackTensor(const ONNX_NAMESPACE::TensorProto& tensor, const void* raw_data, size_t raw_data_len, /*out*/ uint32_t* p_data, size_t expected_size) override { return utils::UnpackTensor(tensor, raw_data, raw_data_len, p_data, expected_size); }
Status UnpackTensor(const ONNX_NAMESPACE::TensorProto& tensor, const void* raw_data, size_t raw_data_len, /*out*/ int64_t* p_data, size_t expected_size) override { return utils::UnpackTensor(tensor, raw_data, raw_data_len, p_data, expected_size); }
Status UnpackTensor(const ONNX_NAMESPACE::TensorProto& tensor, const void* raw_data, size_t raw_data_len, /*out*/ uint64_t* p_data, size_t expected_size) override { return utils::UnpackTensor(tensor, raw_data, raw_data_len, p_data, expected_size); }
void* HeapAllocate(size_t size) override { return new uint8_t[size]; }
void HeapFree(void* p) override { delete[] reinterpret_cast<uint8_t*>(p); }
std::vector<std::string> GetStackTrace() override { return onnxruntime::GetStackTrace(); }
uint16_t math__floatToHalf(float f) override { return math::floatToHalf(f); }
AutoPadType StringToAutoPadType(const std::string& str) override { return onnxruntime::StringToAutoPadType(str); }
int64_t HandleNegativeAxis(int64_t axis, int64_t tensor_rank) override { return onnxruntime::HandleNegativeAxis(axis, tensor_rank); }
float math__halfToFloat(uint16_t h) override { return math::halfToFloat(h); }
void LogRuntimeError(uint32_t session_id, const common::Status& status, const char* file, const char* function, uint32_t line) override {
return ::onnxruntime::LogRuntimeError(session_id, status, file, function, line);
@ -284,6 +313,7 @@ struct ProviderHostImpl : ProviderHost {
const int64_t& int64s__Get(const ONNX_NAMESPACE::int64s* p, int index) override { return p->Get(index); }
// TypeProto_Tensor
bool TypeProto_Tensor__has_shape(const ONNX_NAMESPACE::TypeProto_Tensor* p) override { return p->has_shape(); }
const ONNX_NAMESPACE::TensorShapeProto& TypeProto_Tensor__shape(const ONNX_NAMESPACE::TypeProto_Tensor* p) override { return p->shape(); }
ONNX_NAMESPACE::TensorShapeProto* TypeProto_Tensor__mutable_shape(ONNX_NAMESPACE::TypeProto_Tensor* p) override { return p->mutable_shape(); }
int32_t TypeProto_Tensor__elem_type(const ONNX_NAMESPACE::TypeProto_Tensor* p) override { return p->elem_type(); }
@ -291,6 +321,7 @@ struct ProviderHostImpl : ProviderHost {
// TypeProto
const ONNX_NAMESPACE::TypeProto_Tensor& TypeProto__tensor_type(const ONNX_NAMESPACE::TypeProto* p) override { return p->tensor_type(); }
ONNX_NAMESPACE::TypeProto_Tensor* TypeProto__mutable_tensor_type(ONNX_NAMESPACE::TypeProto* p) override { return p->mutable_tensor_type(); }
int TypeProto__value_case(const ONNX_NAMESPACE::TypeProto* p) override { return p->value_case(); }
// AttributeProto
std::unique_ptr<ONNX_NAMESPACE::AttributeProto> AttributeProto__construct() override { return onnxruntime::make_unique<ONNX_NAMESPACE::AttributeProto>(); }
@ -346,10 +377,19 @@ struct ProviderHostImpl : ProviderHost {
void ModelProto__set_ir_version(ONNX_NAMESPACE::ModelProto* p, int64_t value) override { p->set_ir_version(value); }
// TensorProto
std::unique_ptr<ONNX_NAMESPACE::TensorProto> TensorProto__construct() override { return onnxruntime::make_unique<ONNX_NAMESPACE::TensorProto>(); }
void TensorProto__operator_delete(ONNX_NAMESPACE::TensorProto* p) override { delete p; }
void TensorProto__operator_assign(ONNX_NAMESPACE::TensorProto* p, const ONNX_NAMESPACE::TensorProto& v) override { *p = v; }
bool TensorProto__has_name(const ONNX_NAMESPACE::TensorProto* p) override { return p->has_name(); }
int TensorProto__dims_size(const ONNX_NAMESPACE::TensorProto* p) override { return p->dims_size(); }
const ONNX_NAMESPACE::int64s& TensorProto__dims(const ONNX_NAMESPACE::TensorProto* p) override { return p->dims(); }
bool TensorProto__has_data_location(const ONNX_NAMESPACE::TensorProto* p) override { return p->has_data_location(); }
int TensorProto__data_location(const ONNX_NAMESPACE::TensorProto* p) override { return p->data_location(); }
bool TensorProto__has_raw_data(const ONNX_NAMESPACE::TensorProto* p) override { return p->has_raw_data(); }
const std::string& TensorProto__raw_data(const ONNX_NAMESPACE::TensorProto* p) override { return p->raw_data(); }
int32_t TensorProto__data_type(const ONNX_NAMESPACE::TensorProto* p) override { return p->data_type(); }
bool TensorProto_DataType_IsValid(int value) override { return ONNX_NAMESPACE::TensorProto::DataType_IsValid(value); }
// TensorProtos
ONNX_NAMESPACE::TensorProto* TensorProtos__Add(ONNX_NAMESPACE::TensorProtos* p) override { return p->Add(); }
@ -444,10 +484,13 @@ struct ProviderHostImpl : ProviderHost {
void KernelDefBuilder__TypeConstraint(KernelDefBuilder* p, const char* arg_name, MLDataType supported_type) override { p->TypeConstraint(arg_name, supported_type); }
void KernelDefBuilder__TypeConstraint(KernelDefBuilder* p, const char* arg_name, const std::vector<MLDataType>& supported_types) override { p->TypeConstraint(arg_name, supported_types); }
void KernelDefBuilder__InputMemoryType(KernelDefBuilder* p, OrtMemType type, int input_index) override { p->InputMemoryType(type, input_index); }
void KernelDefBuilder__InputMemoryType(KernelDefBuilder* p, OrtMemType type, const std::vector<int>& input_indexes) override { p->InputMemoryType(type, input_indexes); }
void KernelDefBuilder__OutputMemoryType(KernelDefBuilder* p, OrtMemType type, int input_index) override { p->OutputMemoryType(type, input_index); }
void KernelDefBuilder__ExecQueueId(KernelDefBuilder* p, int queue_id) override { p->ExecQueueId(queue_id); }
void KernelDefBuilder__MayInplace(KernelDefBuilder* p, int input_index, int output_index) override { p->MayInplace(input_index, output_index); }
void KernelDefBuilder__Alias(KernelDefBuilder* p, int input_index, int output_index) override { p->Alias(input_index, output_index); }
void KernelDefBuilder__Alias(KernelDefBuilder* p, const std::vector<std::pair<int, int>>& aliases) override { p->Alias(aliases); }
void KernelDefBuilder__VariadicAlias(KernelDefBuilder* p, int input_offset, int output_offset) override { p->VariadicAlias(input_offset, output_offset); }
std::unique_ptr<KernelDef> KernelDefBuilder__Build(KernelDefBuilder* p) override { return p->Build(); }
@ -571,7 +614,7 @@ struct ProviderHostImpl : ProviderHost {
const std::vector<const NodeArg*>& GraphViewer__GetOutputs(const GraphViewer* p) noexcept override { return p->GetOutputs(); }
const std::vector<const NodeArg*>& GraphViewer__GetValueInfo(const GraphViewer* p) noexcept override { return p->GetValueInfo(); }
const Provider_InitializedTensorSet& GraphViewer__GetAllInitializedTensors(const GraphViewer* p) override { return p->GetAllInitializedTensors(); }
const InitializedTensorSet& GraphViewer__GetAllInitializedTensors(const GraphViewer* p) override { return p->GetAllInitializedTensors(); }
bool GraphViewer__GetInitializedTensor(const GraphViewer* p, const std::string& tensor_name, const ONNX_NAMESPACE::TensorProto*& value) override { return p->GetInitializedTensor(tensor_name, value); }
const std::unordered_map<std::string, int>& GraphViewer__DomainToVersionMap(const GraphViewer* p) override { return p->DomainToVersionMap(); }
@ -587,9 +630,13 @@ struct ProviderHostImpl : ProviderHost {
// OpKernelContext
const Tensor* OpKernelContext__Input_Tensor(const OpKernelContext* p, int index) override { return p->Input<Tensor>(index); }
const Tensor& OpKernelContext__RequiredInput_Tensor(const OpKernelContext* p, int index) override { return p->RequiredInput<Tensor>(index); }
Tensor* OpKernelContext__Output(OpKernelContext* p, int index, const TensorShape& shape) override { return p->Output(index, shape); }
Tensor& OpKernelContext__RequiredOutput(OpKernelContext* p, int index, const TensorShape& shape) override { return p->RequiredOutput(index, shape); }
int OpKernelContext__InputCount(const OpKernelContext* p) override { return p->InputCount(); }
int OpKernelContext__OutputCount(const OpKernelContext* p) override { return p->OutputCount(); }
Status OpKernelContext__GetTempSpaceAllocator(const OpKernelContext* p, AllocatorPtr* output) override { return p->GetTempSpaceAllocator(output); }
bool OpKernelContext__GetUseDeterministicCompute(const OpKernelContext* p) override { return p->GetUseDeterministicCompute(); }
// OpKernelInfo
std::unique_ptr<OpKernelInfo> CopyOpKernelInfo(const OpKernelInfo& info) override { return onnxruntime::CopyOpKernelInfo(info); }
@ -598,6 +645,7 @@ struct ProviderHostImpl : ProviderHost {
Status OpKernelInfo__GetAttr_int64(const OpKernelInfo* p, const std::string& name, int64_t* value) override { return p->GetAttr(name, value); }
Status OpKernelInfo__GetAttr_float(const OpKernelInfo* p, const std::string& name, float* value) override { return p->GetAttr(name, value); }
Status OpKernelInfo__GetAttr_string(const OpKernelInfo* p, const std::string& name, std::string* value) override { return p->GetAttr(name, value); }
Status OpKernelInfo__GetAttr_TensorProto(const OpKernelInfo* p, const std::string& name, ONNX_NAMESPACE::TensorProto* value) override { return p->GetAttr(name, value); }
Status OpKernelInfo__GetAttrs(const OpKernelInfo* p, const std::string& name, std::vector<int64_t>& values) override { return p->GetAttrs(name, values); }
Status OpKernelInfo__GetAttrs(const OpKernelInfo* p, const std::string& name, std::vector<float>& values) override { return p->GetAttrs(name, values); }
Status OpKernelInfo__GetAttrs(const OpKernelInfo* p, const std::string& name, std::vector<std::string>& values) override { return p->GetAttrs(name, values); }
@ -614,6 +662,7 @@ struct ProviderHostImpl : ProviderHost {
// Tensor
std::unique_ptr<Tensor> Tensor__construct(MLDataType p_type, const TensorShape& shape, std::shared_ptr<IAllocator> allocator) override { return onnxruntime::make_unique<Tensor>(p_type, shape, allocator); }
std::unique_ptr<Tensor> Tensor__construct(MLDataType p_type, const TensorShape& shape, void* p_data, const OrtMemoryInfo& alloc, ptrdiff_t offset) override { return onnxruntime::make_unique<Tensor>(p_type, shape, p_data, alloc, offset); }
void Tensor__operator_delete(Tensor* p) override { delete p; }
bool* Tensor__MutableData_bool(Tensor* p) override { return p->MutableData<bool>(); }
@ -627,6 +676,7 @@ struct ProviderHostImpl : ProviderHost {
uint64_t* Tensor__MutableData_uint64(Tensor* p) override { return p->MutableData<uint64_t>(); }
float* Tensor__MutableData_float(Tensor* p) override { return p->MutableData<float>(); }
double* Tensor__MutableData_double(Tensor* p) override { return p->MutableData<double>(); }
BFloat16* Tensor__MutableData_BFloat16(Tensor* p) override { return p->MutableData<BFloat16>(); }
MLFloat16* Tensor__MutableData_MLFloat16(Tensor* p) override { return p->MutableData<MLFloat16>(); }
const bool* Tensor__Data_bool(const Tensor* p) override { return p->Data<bool>(); }
@ -640,21 +690,16 @@ struct ProviderHostImpl : ProviderHost {
const uint64_t* Tensor__Data_uint64(const Tensor* p) override { return p->Data<uint64_t>(); }
const float* Tensor__Data_float(const Tensor* p) override { return p->Data<float>(); }
const double* Tensor__Data_double(const Tensor* p) override { return p->Data<double>(); }
const BFloat16* Tensor__Data_BFloat16(const Tensor* p) override { return p->Data<BFloat16>(); }
const MLFloat16* Tensor__Data_MLFloat16(const Tensor* p) override { return p->Data<MLFloat16>(); }
gsl::span<const int64_t> Tensor__DataAsSpan_int64(const Tensor* p) override { return p->DataAsSpan<int64_t>(); }
void* Tensor__MutableDataRaw(Tensor* p, MLDataType type) override { return p->MutableDataRaw(type); }
const void* Tensor__DataRaw(const Tensor* p, MLDataType type) override { return p->DataRaw(type); }
void* Tensor__MutableDataRaw(Tensor* p) noexcept override { return p->MutableDataRaw(); }
const void* Tensor__DataRaw(const Tensor* p) noexcept override { return p->DataRaw(); }
const TensorShape& Tensor__Shape(const Tensor* p) override { return p->Shape(); }
size_t Tensor__SizeInBytes(const Tensor* p) override { return p->SizeInBytes(); }
const OrtMemoryInfo& Tensor__Location(const Tensor* p) override { return p->Location(); }
int32_t Tensor__GetElementType(const Tensor* p) override { return p->GetElementType(); }
MLDataType Tensor__DataType(const Tensor* p) override { return p->DataType(); }
bool Tensor__IsDataType_bool(const Tensor* p) noexcept override { return p->IsDataType<bool>(); }
bool Tensor__IsDataType_int8(const Tensor* p) noexcept override { return p->IsDataType<int8_t>(); }
bool Tensor__IsDataType_uint8(const Tensor* p) noexcept override { return p->IsDataType<uint8_t>(); }
@ -669,16 +714,34 @@ struct ProviderHostImpl : ProviderHost {
bool Tensor__IsDataType_MLFloat16(const Tensor* p) noexcept override { return p->IsDataType<MLFloat16>(); }
bool Tensor__IsDataTypeString(const Tensor* p) noexcept override { return p->IsDataTypeString(); }
const TensorShape& Tensor__Shape(const Tensor* p) override { return p->Shape(); }
void Tensor__Reshape(Tensor* p, const TensorShape& new_shape) override { return p->Reshape(new_shape); }
void Tensor__SetByteOffset(Tensor* p, ptrdiff_t byte_offset) override { p->SetByteOffset(byte_offset); }
ptrdiff_t Tensor__ByteOffset(const Tensor* p) override { return p->ByteOffset(); }
size_t Tensor__SizeInBytes(const Tensor* p) override { return p->SizeInBytes(); }
const OrtMemoryInfo& Tensor__Location(const Tensor* p) override { return p->Location(); }
int32_t Tensor__GetElementType(const Tensor* p) override { return p->GetElementType(); }
MLDataType Tensor__DataType(const Tensor* p) override { return p->DataType(); }
// GatherElements
Status GatherElements__ValidateInputShapes(const TensorShape& input_data_shape, const TensorShape& indices_shape, int64_t axis) override { return GatherElements::ValidateInputShapes(input_data_shape, indices_shape, axis); }
// cumsum.cc
Status cumsum_op__GetAxis(const Tensor* axis_tensor, int64_t input_rank, int64_t& axis_out) override { return cumsum_op::GetAxis(axis_tensor, input_rank, axis_out); }
// TileOp
bool TileOp__IsTileMemcpy(const TensorShape& input_shape, const int64_t* repeats, size_t rank, bool& is_batched_memcpy, size_t& num_of_elements_per_batch, size_t& num_of_copies_per_batch, size_t& num_of_batch_copies) override { return TileOp::IsTileMemcpy(input_shape, repeats, rank, is_batched_memcpy, num_of_elements_per_batch, num_of_copies_per_batch, num_of_batch_copies); }
// ROI
Status CheckROIAlignValidInput(const Tensor* X_ptr, const Tensor* rois_ptr, const Tensor* batch_indices_ptr) override { return CheckROIAlignValidInput(X_ptr, rois_ptr, batch_indices_ptr); }
Status NonMaxSuppressionBase__PrepareCompute(OpKernelContext* ctx, PrepareContext& pc) override { return NonMaxSuppressionBase::PrepareCompute(ctx, pc); }
Status NonMaxSuppressionBase__GetThresholdsFromInputs(const PrepareContext& pc, int64_t& max_output_boxes_per_class, float& iou_threshold, float& score_threshold) override { return NonMaxSuppressionBase::GetThresholdsFromInputs(pc, max_output_boxes_per_class, iou_threshold, score_threshold); }
// From onehot.h
Status ValidateInputs(const Tensor* depth, const Tensor* values) override { return onnxruntime::ValidateInputs(depth, values); }
Status PrepareOutputShape(const Tensor* indices, const int64_t depth_val, const int64_t axis, int64_t& prefix_dim_size, int64_t& suffix_dim_size, std::vector<int64_t>& output_shape) override { return onnxruntime::PrepareOutputShape(indices, depth_val, axis, prefix_dim_size, suffix_dim_size, output_shape); }
// From Providers/Common.h
bool IsScalarOr1ElementVector(const Tensor* input) override { return onnxruntime::IsScalarOr1ElementVector(input); }
// From cpu/tensor/unsqueeze.h
Status UnsqueezeBase__PrepareCompute(const UnsqueezeBase* p, OpKernelContext* ctx, UnsqueezeBase__Prepare& prepare) override { return p->PrepareCompute(ctx, reinterpret_cast<UnsqueezeBase::Prepare&>(prepare)); }
// From cpu/tensor/slice.h
@ -718,9 +781,63 @@ struct ProviderHostImpl : ProviderHost {
// From cpu/tensor/gatherbase.h
Status GatherBase__PrepareForCompute(const GatherBase* p, OpKernelContext* context, GatherBase__Prepare& prepare) override { return p->PrepareForCompute(context, reinterpret_cast<GatherBase::Prepare&>(prepare)); }
PhiloxGenerator& PhiloxGenerator__Default() override { return PhiloxGenerator::Default(); }
Status Einsum__Compute(const Einsum* p, OpKernelContext* context) override { return p->Compute(context); }
// EinsumComputePreprocessor
void EinsumComputePreprocessor__operator_delete(EinsumComputePreprocessor* p) override { delete p; }
virtual std::unique_ptr<EinsumComputePreprocessor> EinsumComputePreprocessor__Create(EinsumEquationPreprocessor& equation_preprocessor,
const std::vector<const Tensor*>& inputs,
AllocatorPtr allocator,
void* einsum_cuda_assets) override { return onnxruntime::make_unique<EinsumComputePreprocessor>(equation_preprocessor, inputs, allocator, einsum_cuda_assets); }
virtual Status EinsumComputePreprocessor__Run(EinsumComputePreprocessor* p) override { return p->Run(); }
virtual void EinsumComputePreprocessor__SetDeviceHelpers(EinsumComputePreprocessor* p, const EinsumOp::DeviceHelpers::Diagonal& diagonal_func, const EinsumOp::DeviceHelpers::Transpose& transpose_func) override { return p->SetDeviceHelpers(diagonal_func, transpose_func); }
// EinsumTypedComputeProcessor
void EinsumTypedComputeProcessor__operator_delete(EinsumTypedComputeProcessor<float>* p) override { delete p; }
void EinsumTypedComputeProcessor__operator_delete(EinsumTypedComputeProcessor<double>* p) override { delete p; }
void EinsumTypedComputeProcessor__operator_delete(EinsumTypedComputeProcessor<MLFloat16>* p) override { delete p; }
std::unique_ptr<EinsumTypedComputeProcessor<float>> EinsumTypedComputeProcessor_float__Create(OpKernelContext* context, AllocatorPtr allocator, concurrency::ThreadPool* tp, EinsumComputePreprocessor& einsum_compute_preprocessor, void* einsum_cuda_assets) override { return onnxruntime::make_unique<EinsumTypedComputeProcessor<float>>(context, allocator, tp, einsum_compute_preprocessor, einsum_cuda_assets); }
std::unique_ptr<EinsumTypedComputeProcessor<double>> EinsumTypedComputeProcessor_double__Create(OpKernelContext* context, AllocatorPtr allocator, concurrency::ThreadPool* tp, EinsumComputePreprocessor& einsum_compute_preprocessor, void* einsum_cuda_assets) override { return onnxruntime::make_unique<EinsumTypedComputeProcessor<double>>(context, allocator, tp, einsum_compute_preprocessor, einsum_cuda_assets); }
std::unique_ptr<EinsumTypedComputeProcessor<MLFloat16>> EinsumTypedComputeProcessor_MLFloat16__Create(OpKernelContext* context, AllocatorPtr allocator, concurrency::ThreadPool* tp, EinsumComputePreprocessor& einsum_compute_preprocessor, void* einsum_cuda_assets) override { return onnxruntime::make_unique<EinsumTypedComputeProcessor<MLFloat16>>(context, allocator, tp, einsum_compute_preprocessor, einsum_cuda_assets); }
void EinsumTypedComputeProcessor__SetDeviceHelpers(EinsumTypedComputeProcessor<float>* p, const EinsumOp::DeviceHelpers::Transpose& device_transpose_func, const EinsumOp::DeviceHelpers::MatMul<float>& device_matmul_func, const EinsumOp::DeviceHelpers::ReduceSum<float>& device_reduce_sum_func, const EinsumOp::DeviceHelpers::DataCopy& device_data_copy_func) override { return p->SetDeviceHelpers(device_transpose_func, device_matmul_func, device_reduce_sum_func, device_data_copy_func); }
void EinsumTypedComputeProcessor__SetDeviceHelpers(EinsumTypedComputeProcessor<double>* p, const EinsumOp::DeviceHelpers::Transpose& device_transpose_func, const EinsumOp::DeviceHelpers::MatMul<double>& device_matmul_func, const EinsumOp::DeviceHelpers::ReduceSum<double>& device_reduce_sum_func, const EinsumOp::DeviceHelpers::DataCopy& device_data_copy_func) override { return p->SetDeviceHelpers(device_transpose_func, device_matmul_func, device_reduce_sum_func, device_data_copy_func); }
void EinsumTypedComputeProcessor__SetDeviceHelpers(EinsumTypedComputeProcessor<MLFloat16>* p, const EinsumOp::DeviceHelpers::Transpose& device_transpose_func, const EinsumOp::DeviceHelpers::MatMul<MLFloat16>& device_matmul_func, const EinsumOp::DeviceHelpers::ReduceSum<MLFloat16>& device_reduce_sum_func, const EinsumOp::DeviceHelpers::DataCopy& device_data_copy_func) override { return p->SetDeviceHelpers(device_transpose_func, device_matmul_func, device_reduce_sum_func, device_data_copy_func); }
Status EinsumTypedComputeProcessor__Run(EinsumTypedComputeProcessor<float>* p) override { return p->Run(); }
Status EinsumTypedComputeProcessor__Run(EinsumTypedComputeProcessor<double>* p) override { return p->Run(); }
Status EinsumTypedComputeProcessor__Run(EinsumTypedComputeProcessor<MLFloat16>* p) override { return p->Run(); }
// AllocatorManager
void AllocatorManager__InsertAllocator(AllocatorManager* p, AllocatorPtr allocator) override { p->InsertAllocator(allocator); }
AllocatorPtr AllocatorManager__GetAllocator(AllocatorManager* p, int id, OrtMemType mem_type) override { return p->GetAllocator(id, mem_type); };
AllocatorPtr AllocatorManager__GetAllocator(const AllocatorManager* p, int id, OrtMemType mem_type) override { return p->GetAllocator(id, mem_type); };
std::unique_ptr<OpKernel> CreateOpKernel_CPU_If(const OpKernelInfo& info) override { return onnxruntime::make_unique<If>(info); }
std::unique_ptr<OpKernel> CreateOpKernel_CPU_Loop(const OpKernelInfo& info, const void* concat_output_func, void* stream) override { return Loop::Create(info, *reinterpret_cast<const Loop::ConcatOutput*>(concat_output_func), stream); }
std::unique_ptr<OpKernel> CreateOpKernel_CPU_Scan_8(const OpKernelInfo& info) override { return onnxruntime::make_unique<Scan<8>>(info); }
std::unique_ptr<OpKernel> CreateOpKernel_CPU_Scan_9(const OpKernelInfo& info) override { return onnxruntime::make_unique<Scan<9>>(info); }
// ContribOps
#ifndef DISABLE_CONTRIB_OPS
Status embed_layer_norm__CheckInputs(const OpKernelContext* context) override { return contrib::embed_layer_norm::CheckInputs(context); }
Status bias_gelu_helper__CheckInputs(const OpKernelContext* context) override { return contrib::bias_gelu_helper::CheckInputs(context); }
Status LongformerAttentionBase__CheckInputs(const contrib::LongformerAttentionBase* p, const TensorShape& input_shape, const TensorShape& weights_shape, const TensorShape& bias_shape, const TensorShape& mask_shape, const TensorShape& global_weights_shape, const TensorShape& global_bias_shape, const TensorShape& global_shape) override { return p->CheckInputs(input_shape, weights_shape, bias_shape, mask_shape, global_weights_shape, global_bias_shape, global_shape); }
Status AttentionBase__CheckInputs(const contrib::AttentionBase* p, const TensorShape& input_shape, const TensorShape& weights_shape, const TensorShape& bias_shape, const Tensor*& mask_index, const Tensor* past, const int max_threads_per_block) override { return p->CheckInputs(input_shape, weights_shape, bias_shape, mask_index, past, max_threads_per_block); }
Tensor* AttentionBase__GetPresent(const contrib::AttentionBase* p, OpKernelContext* context, const Tensor* past, int batch_size, int head_size, int sequence_length, int& past_sequence_length) override { return p->GetPresent(context, past, batch_size, head_size, sequence_length, past_sequence_length); }
#endif
#ifdef ENABLE_TRAINING
void contrib__record_event_in_tensor(const Tensor& event_id_tensor) override { return contrib::record_event_in_tensor(event_id_tensor); }
void contrib__wait_event_in_tensor(const Tensor& event_id_tensor) override { return contrib::wait_event_in_tensor(event_id_tensor); }
Status contrib__Group__Compute(const contrib::Group* p, OpKernelContext* context) override { return p->Compute(context); }
Status contrib__PassThrough__Compute(const contrib::PassThrough* p, OpKernelContext* context) override { return p->Compute(context); }
void contrib__VerifyLogitWeightAndLabelShape(const TensorShape& logit_shape, const TensorShape& label_shape, const TensorShape* weight_shape) override { contrib::VerifyLogitWeightAndLabelShape(logit_shape, label_shape, weight_shape); }
void contrib__GetNDCFromLogitAndLabelShape(const TensorShape& logit_shape, const TensorShape& label_shape, int64_t& N_D, int64_t& C) override { contrib::GetNDCFromLogitAndLabelShape(logit_shape, label_shape, N_D, C); }
void contrib__GetPermutationAndShape(bool ncd_to_ndc, const TensorShape& tensor_shape, std::vector<int64_t>& new_shape, std::vector<size_t>& permutations) override { contrib::GetPermutationAndShape(ncd_to_ndc, tensor_shape, new_shape, permutations); }
Status contrib__PrepareForTrainingCompute(const TensorShape& input_shape, int num_outputs, int64_t& axis, int& before_dims, int& after_dims_including_split_axis, int& after_dims_excluding_split, std::vector<int64_t>& split_sizes) override { return contrib::PrepareForTrainingCompute(input_shape, num_outputs, axis, before_dims, after_dims_including_split_axis, after_dims_excluding_split, split_sizes); }
#endif
} provider_host_;
@ -820,6 +937,14 @@ void UnloadSharedProviders() {
s_library_shared.Unload();
}
// Used by test code
std::unique_ptr<IAllocator> CreateCUDAPinnedAllocator(int16_t device_id, const char* name) {
if (auto* info = onnxruntime::GetProviderInfo_CUDA())
return info->CreateCUDAPinnedAllocator(device_id, name);
return nullptr;
}
std::shared_ptr<IExecutionProviderFactory> CreateExecutionProviderFactory_Cuda(const OrtCUDAProviderOptions* provider_options) {
if (auto provider = s_library_cuda.Get())
return provider->CreateExecutionProviderFactory(provider_options);
@ -855,12 +980,35 @@ std::shared_ptr<IExecutionProviderFactory> CreateExecutionProviderFactory_OpenVI
return nullptr;
}
const ProviderInfo_OpenVINO* GetProviderInfo_OpenVINO() {
ProviderInfo_OpenVINO* GetProviderInfo_OpenVINO() {
if (auto provider = s_library_openvino.Get())
return reinterpret_cast<const ProviderInfo_OpenVINO*>(provider->GetInfo());
return reinterpret_cast<ProviderInfo_OpenVINO*>(provider->GetInfo());
return nullptr;
}
ProviderInfo_CUDA* GetProviderInfo_CUDA() {
if (auto provider = s_library_cuda.Get())
return reinterpret_cast<ProviderInfo_CUDA*>(provider->GetInfo());
return nullptr;
}
void CopyGpuToCpu(
void* dst_ptr,
const void* src_ptr,
const size_t size,
const OrtMemoryInfo& dst_location,
const OrtMemoryInfo& src_location) {
if (auto* info = onnxruntime::GetProviderInfo_CUDA())
return info->CopyGpuToCpu(dst_ptr, src_ptr, size, dst_location, src_location);
ORT_THROW("GPU-to-CPU copy is not implemented.");
}
void cudaMemcpy_HostToDevice(void* dst, const void* src, size_t count) {
if (auto* info = onnxruntime::GetProviderInfo_CUDA())
return info->cudaMemcpy_HostToDevice(dst, src, count);
ORT_THROW("cudaMemcpy_HostToDevice is not implemented.");
}
} // namespace onnxruntime
ORT_API_STATUS_IMPL(OrtSessionOptionsAppendExecutionProvider_Dnnl, _In_ OrtSessionOptions* options, int use_arena) {
@ -916,6 +1064,18 @@ ORT_API_STATUS_IMPL(OrtSessionOptionsAppendExecutionProvider_CUDA, _In_ OrtSessi
return OrtApis::SessionOptionsAppendExecutionProvider_CUDA(options, &provider_options);
}
ORT_API_STATUS_IMPL(OrtApis::SetCurrentGpuDeviceId, _In_ int device_id) {
if (auto* info = onnxruntime::GetProviderInfo_CUDA())
return info->SetCurrentGpuDeviceId(device_id);
return nullptr;
}
ORT_API_STATUS_IMPL(OrtApis::GetCurrentGpuDeviceId, _In_ int* device_id) {
if (auto* info = onnxruntime::GetProviderInfo_CUDA())
return info->GetCurrentGpuDeviceId(device_id);
return nullptr;
}
ORT_API_STATUS_IMPL(OrtApis::SessionOptionsAppendExecutionProvider_CUDA, _In_ OrtSessionOptions* options, _In_ const OrtCUDAProviderOptions* cuda_options) {
auto factory = onnxruntime::CreateExecutionProviderFactory_Cuda(cuda_options);
if (!factory) {
@ -925,4 +1085,3 @@ ORT_API_STATUS_IMPL(OrtApis::SessionOptionsAppendExecutionProvider_CUDA, _In_ Or
options->provider_factories.push_back(factory);
return nullptr;
}

21
onnxruntime/core/framework/tensorprotoutils.h
View file

@ -6,6 +6,7 @@
#include <vector>
#include <type_traits>
#ifndef SHARED_PROVIDER
#include "core/common/common.h"
#include "core/common/path.h"
#include "core/common/status.h"
@ -26,15 +27,17 @@ bool operator==(const TensorShapeProto_Dimension& l, const TensorShapeProto_Dime
bool operator!=(const TensorShapeProto_Dimension& l, const TensorShapeProto_Dimension& r);
} // namespace ONNX_NAMESPACE
#endif
namespace onnxruntime {
class Tensor;
namespace utils {
#ifndef SHARED_PROVIDER
TensorShape GetTensorShapeFromTensorShapeProto(const ONNX_NAMESPACE::TensorShapeProto& tensor_shape_proto);
std::vector<int64_t> GetTensorShapeFromTensorProto(const ONNX_NAMESPACE::TensorProto& tensor_proto);
/**
/**
* deserialize a TensorProto into a preallocated memory buffer.
* \param tensor_proto_path A local file path of where the 'input' was loaded from. Can be NULL if the tensor proto doesn't
* have any external data or it was loaded from current working dir. This path could be either a
@ -51,8 +54,8 @@ common::Status TensorProtoToMLValue(const Env& env, const ORTCHAR_T* tensor_prot
* @return
*/
common::Status TensorProtoToTensor(const Env& env, const ORTCHAR_T* model_path,
const ONNX_NAMESPACE::TensorProto& tensor_proto,
Tensor& tensor);
const ONNX_NAMESPACE::TensorProto& tensor_proto,
Tensor& tensor);
/** Creates a TensorProto from a Tensor.
@param[in] tensor the Tensor whose data and shape will be used to create the TensorProto.
@ -99,6 +102,7 @@ common::Status DenseTensorToSparseTensorProto(const ONNX_NAMESPACE::TensorProto&
const Path& model_path,
ONNX_NAMESPACE::SparseTensorProto& sparse);
#endif // !ORT_MINIMAL_BUILD
#endif
inline bool HasDimValue(const ONNX_NAMESPACE::TensorShapeProto_Dimension& dim) {
return dim.value_case() == ONNX_NAMESPACE::TensorShapeProto_Dimension::kDimValue;
@ -122,10 +126,12 @@ inline bool HasShape(const ONNX_NAMESPACE::TypeProto_Tensor& ten_proto) {
return ten_proto.has_shape();
}
#ifndef SHARED_PROVIDER
inline bool HasShape(const ONNX_NAMESPACE::TypeProto_SparseTensor& ten_proto) {
// XXX: Figure out how do in proto3
return ten_proto.has_shape();
}
#endif
inline bool HasRawData(const ONNX_NAMESPACE::TensorProto& ten_proto) {
// Can not be UNDEFINED and can not be STRING but test for STRING is usually performed separately
@ -145,6 +151,7 @@ inline bool HasDataType(const ONNX_NAMESPACE::TensorProto& ten_proto) {
return ten_proto.data_type() != ONNX_NAMESPACE::TensorProto::UNDEFINED;
}
#ifndef SHARED_PROVIDER
inline bool HasName(const ONNX_NAMESPACE::TensorProto& ten_proto) {
return ten_proto.has_name(); // XXX
}
@ -168,11 +175,13 @@ inline bool HasKeyType(const ONNX_NAMESPACE::TypeProto_Map& map_proto) {
inline bool HasValueType(const ONNX_NAMESPACE::TypeProto_Map& map_proto) {
return map_proto.value_type().value_case() != ONNX_NAMESPACE::TypeProto::VALUE_NOT_SET;
}
#endif
inline bool HasType(const ONNX_NAMESPACE::ValueInfoProto& vi_proto) {
return vi_proto.type().value_case() != ONNX_NAMESPACE::TypeProto::VALUE_NOT_SET;
}
#ifndef SHARED_PROVIDER
inline bool HasName(const ONNX_NAMESPACE::ValueInfoProto& vi_proto) {
return vi_proto.has_name(); // XXX: Figure out proto3 way
}
@ -184,6 +193,7 @@ inline bool HasDomain(const ONNX_NAMESPACE::TypeProto_Opaque& op_proto) {
inline bool HasName(const ONNX_NAMESPACE::TypeProto_Opaque& op_proto) {
return !op_proto.name().empty();
}
#endif
inline bool HasType(const ONNX_NAMESPACE::AttributeProto& at_proto) {
return at_proto.type() != ONNX_NAMESPACE::AttributeProto::AttributeType::AttributeProto_AttributeType_UNDEFINED;
@ -229,6 +239,7 @@ inline bool HasGraphs(const ONNX_NAMESPACE::AttributeProto& at_proto) {
return at_proto.type() == ONNX_NAMESPACE::AttributeProto::AttributeType::AttributeProto_AttributeType_GRAPHS;
}
#ifndef SHARED_PROVIDER
inline bool HasName(const ONNX_NAMESPACE::AttributeProto& at_proto) {
return at_proto.has_name(); // XXX: Fugure out proto3
}
@ -249,6 +260,7 @@ inline bool HasName(const ONNX_NAMESPACE::NodeProto& node_proto) {
//XXX: Figure out proto3 style
return node_proto.has_name();
}
#endif
// UnpackTensor from raw data or the type specific data field. Does not handle external data.
// If the tensor does not contain raw data then raw_data should be nullptr and raw_data_len should be 0.
@ -276,6 +288,5 @@ common::Status UnpackInitializerData(const ONNX_NAMESPACE::TensorProto& initiali
const Path& model_path,
std::unique_ptr<unsigned char[]>& unpacked_tensor,
size_t& tensor_byte_size) ORT_MUST_USE_RESULT;
} // namespace utils
} // namespace onnxruntime

4
onnxruntime/core/platform/env_var_utils.h
View file

@ -14,7 +14,11 @@ namespace onnxruntime {
*/
template <typename T>
optional<T> ParseEnvironmentVariable(const std::string& name) {
#ifndef SHARED_PROVIDER
const std::string value_str = Env::Default().GetEnvironmentVar(name);
#else
const std::string value_str = GetEnvironmentVar(name);
#endif
if (value_str.empty()) {
return {};
}

2
onnxruntime/core/providers/common.h
View file

@ -6,8 +6,10 @@
#include <cstdint>
#include <functional>
#ifndef SHARED_PROVIDER
#include "core/common/common.h"
#include "core/framework/tensor.h"
#endif
namespace onnxruntime {

3
onnxruntime/core/providers/cpu/controlflow/if.h
View file

@ -5,9 +5,6 @@
#include <functional>
#include "gsl/gsl"
//#include "core/common/common.h"
//#include "core/framework/feeds_fetches_manager.h"
//#include "core/framework/op_kernel.h"
#include "core/providers/cpu/controlflow/utils.h"
namespace onnxruntime {

7
onnxruntime/core/providers/cpu/controlflow/loop.cc
View file

@ -259,6 +259,13 @@ Loop::Loop(const OpKernelInfo& info) : IControlFlowKernel(info) {
stream_ = nullptr;
}
std::unique_ptr<OpKernel> Loop::Create(const OpKernelInfo& info, const ConcatOutput& concat_output_func, void* stream) {
auto result = make_unique<Loop>(info);
result->SetConcatOutputFunc(concat_output_func);
result->SetComputeStream(stream);
return result;
}
// we need this to be in the .cc so 'unique_ptr<Info> info_' can be handled
Loop::~Loop() = default;

2
onnxruntime/core/providers/cpu/controlflow/loop.h
View file

@ -32,6 +32,8 @@ class Loop : public controlflow::IControlFlowKernel {
using ConcatOutput = std::function<Status(void* stream, std::vector<OrtValue>& per_iteration_output,
void* output, size_t output_size_in_bytes)>;
static std::unique_ptr<OpKernel> Create(const OpKernelInfo& info, const ConcatOutput& concat_output_func, void* stream);
protected:
// derived class can provide implementation for handling concatenation of Loop output on a different device
void SetConcatOutputFunc(const ConcatOutput& concat_output_func) { concat_output_func_ = concat_output_func; }

4
onnxruntime/core/providers/cpu/controlflow/scan.h
View file

@ -5,8 +5,11 @@
#include <functional>
#include "gsl/gsl"
#ifndef SHARED_PROVIDER
#include "core/common/common.h"
#include "core/framework/op_kernel.h"
#endif
#include "core/framework/feeds_fetches_manager.h"
#include "core/providers/cpu/controlflow/utils.h"
#include "core/framework/ort_value_tensor_slicer.h"
@ -56,7 +59,6 @@ class Scan : public controlflow::IControlFlowKernel {
struct Info;
~Scan();
protected:
void SetDeviceHelpers(const scan::detail::DeviceHelpers& device_helpers) {
device_helpers_ = device_helpers; // copy
}

1
onnxruntime/core/providers/cpu/controlflow/utils.h
View file

@ -7,7 +7,6 @@
#include <string>
#include <vector>
//#include "core/common/common.h"
#include "core/framework/feeds_fetches_manager.h"
#ifndef SHARED_PROVIDER
#include "core/framework/op_kernel.h"

16
onnxruntime/core/providers/cpu/generator/constant_of_shape_base.h
View file

@ -3,6 +3,7 @@
#pragma once
#ifndef SHARED_PROVIDER
#include "core/common/common.h"
#include "core/common/type_list.h"
#include "core/framework/data_types.h"
@ -10,6 +11,7 @@
#include "core/framework/op_kernel.h"
#include "core/framework/tensorprotoutils.h"
#include "core/providers/op_kernel_type_control_utils.h"
#endif
namespace onnxruntime {
@ -25,11 +27,17 @@ template <typename EnabledOutputTypeList = ConstantOfShapeDefaultOutputTypes>
class ConstantOfShapeBase {
protected:
ConstantOfShapeBase(const OpKernelInfo& info) {
#ifndef SHARED_PROVIDER
ONNX_NAMESPACE::TensorProto t_proto;
if (info.GetAttr<ONNX_NAMESPACE::TensorProto>("value", &t_proto).IsOK()) {
ORT_ENFORCE(t_proto.dims_size() == 1, "Must have a single dimension");
ORT_ENFORCE(t_proto.dims()[0] == 1, "Must have a single dimension of 1");
SetValueFromTensorProto(t_proto);
auto* t_proto_p = &t_proto;
#else
auto t_proto = ONNX_NAMESPACE::TensorProto::Create();
auto* t_proto_p = t_proto.get();
#endif
if (info.GetAttr<ONNX_NAMESPACE::TensorProto>("value", t_proto_p).IsOK()) {
ORT_ENFORCE(t_proto_p->dims_size() == 1, "Must have a single dimension");
ORT_ENFORCE(t_proto_p->dims()[0] == 1, "Must have a single dimension of 1");
SetValueFromTensorProto(*t_proto_p);
} else {
float f_value = 0.f;
SetValue(sizeof(float), reinterpret_cast<void*>(&f_value));

4
onnxruntime/core/providers/cpu/math/einsum.h
View file

@ -3,10 +3,12 @@
#pragma once
#ifndef SHARED_PROVIDER
#include "core/common/common.h"
#include "core/framework/op_kernel.h"
#include "einsum_utils/einsum_compute_preprocessor.h"
#include "einsum_utils/einsum_typed_compute_processor.h"
#endif
#include "einsum_utils/einsum_compute_preprocessor.h"
namespace onnxruntime {

19
onnxruntime/core/providers/cpu/math/einsum_utils/einsum_auxiliary_ops.cc
View file

@ -49,10 +49,10 @@ Status MatMul(const T* input_1_data, const T* input_2_data, T* output_data,
// CPU specific ReduceSum helper
template <typename T>
Tensor ReduceSum(const Tensor& input, const std::vector<int64_t>& reduce_axes,
bool keep_dims, AllocatorPtr allocator,
const TensorShape* input_shape_override,
concurrency::ThreadPool* tp, void* /*einsum_cuda_assets*/) {
std::unique_ptr<Tensor> ReduceSum(const Tensor& input, const std::vector<int64_t>& reduce_axes,
bool keep_dims, AllocatorPtr allocator,
const TensorShape* input_shape_override,
concurrency::ThreadPool* tp, void* /*einsum_cuda_assets*/) {
return onnxruntime::ReduceSum<T>::Impl(input, reduce_axes,
allocator, tp, keep_dims,
input_shape_override);
@ -352,8 +352,7 @@ std::unique_ptr<Tensor> ReduceSum(const Tensor& input, const std::vector<int64_t
concurrency::ThreadPool* tp, void* einsum_cuda_assets,
const DeviceHelpers::ReduceSum<T>& device_reduce_sum_func) {
TensorShape overriden_shape(input_shape_override);
auto output = device_reduce_sum_func(input, reduce_axes, true, allocator, &overriden_shape, tp, einsum_cuda_assets);
return onnxruntime::make_unique<Tensor>(std::move(output));
return device_reduce_sum_func(input, reduce_axes, true, allocator, &overriden_shape, tp, einsum_cuda_assets);
}
// Explicit template instantiations of functions
@ -371,7 +370,7 @@ template std::unique_ptr<Tensor> MatMul<float>(
AllocatorPtr allocator, concurrency::ThreadPool* tp, void* einsum_cuda_assets,
const DeviceHelpers::MatMul<float>& device_matmul_func);
template Tensor DeviceHelpers::CpuDeviceHelpers::ReduceSum<float>(
template std::unique_ptr<Tensor> DeviceHelpers::CpuDeviceHelpers::ReduceSum<float>(
const Tensor& input, const std::vector<int64_t>& reduce_axes,
bool keep_dims, AllocatorPtr allocator,
const TensorShape* input_shape_override,
@ -395,7 +394,7 @@ template std::unique_ptr<Tensor> MatMul<int32_t>(
AllocatorPtr allocator, concurrency::ThreadPool* tp, void* einsum_cuda_assets,
const DeviceHelpers::MatMul<int32_t>& device_matmul_func);
template Tensor DeviceHelpers::CpuDeviceHelpers::ReduceSum<int32_t>(
template std::unique_ptr<Tensor> DeviceHelpers::CpuDeviceHelpers::ReduceSum<int32_t>(
const Tensor& input, const std::vector<int64_t>& reduce_axes,
bool keep_dims, AllocatorPtr allocator,
const TensorShape* input_shape_override,
@ -420,7 +419,7 @@ template std::unique_ptr<Tensor> MatMul<double>(
AllocatorPtr allocator, concurrency::ThreadPool* tp, void* einsum_cuda_assets,
const DeviceHelpers::MatMul<double>& device_matmul_func);
template Tensor DeviceHelpers::CpuDeviceHelpers::ReduceSum<double>(
template std::unique_ptr<Tensor> DeviceHelpers::CpuDeviceHelpers::ReduceSum<double>(
const Tensor& input, const std::vector<int64_t>& reduce_axes,
bool keep_dims, AllocatorPtr allocator,
const TensorShape* input_shape_override,
@ -439,7 +438,7 @@ template Status DeviceHelpers::CpuDeviceHelpers::MatMul<int64_t>(
size_t num_batches, size_t M, size_t K, size_t N, concurrency::ThreadPool* tp,
void* einsum_cuda_assets);
template Tensor DeviceHelpers::CpuDeviceHelpers::ReduceSum<int64_t>(
template std::unique_ptr<Tensor> DeviceHelpers::CpuDeviceHelpers::ReduceSum<int64_t>(
const Tensor& input, const std::vector<int64_t>& reduce_axes,
bool keep_dims, AllocatorPtr allocator,
const TensorShape* input_shape_override,

16
onnxruntime/core/providers/cpu/math/einsum_utils/einsum_auxiliary_ops.h
View file

@ -41,10 +41,10 @@ using MatMul = std::function<Status(const T* input_1_data, const T* input_2_data
// ReduceSum op - Reduces along `reduce_axes`
template <typename T>
using ReduceSum = std::function<Tensor(const Tensor& input, const std::vector<int64_t>& reduce_axes,
bool keep_dims, AllocatorPtr allocator,
const TensorShape* input_shape_override,
concurrency::ThreadPool* tp, void* einsum_cuda_assets)>;
using ReduceSum = std::function<std::unique_ptr<Tensor>(const Tensor& input, const std::vector<int64_t>& reduce_axes,
bool keep_dims, AllocatorPtr allocator,
const TensorShape* input_shape_override,
concurrency::ThreadPool* tp, void* einsum_cuda_assets)>;
// Diagonal op
// Diagonal - A specialized implementation somewhat similar to Torch's Diagonal op
@ -73,10 +73,10 @@ Status MatMul(const T* input_1_data, const T* input_2_data, T* output_data,
void* einsum_cuda_assets);
template <typename T>
Tensor ReduceSum(const Tensor& input, const std::vector<int64_t>& reduce_axes,
bool keep_dims, AllocatorPtr allocator,
const TensorShape* input_shape_override,
concurrency::ThreadPool* tp, void* einsum_cuda_assets);
std::unique_ptr<Tensor> ReduceSum(const Tensor& input, const std::vector<int64_t>& reduce_axes,
bool keep_dims, AllocatorPtr allocator,
const TensorShape* input_shape_override,
concurrency::ThreadPool* tp, void* einsum_cuda_assets);
std::unique_ptr<Tensor> Diagonal(const Tensor& input, int64_t dim_1, int64_t dim_2, AllocatorPtr allocator, void* einsum_cuda_assets);

3
onnxruntime/core/providers/cpu/math/einsum_utils/einsum_compute_preprocessor.h
View file

@ -104,6 +104,7 @@ struct EinsumEquationPreprocessor {
// Subscript labels (letter) and subcript indices (a unique id to the letter) are interchangeable
// This is a pre-processor class that maps subscript labels to a dimension value, etc.
#ifndef SHARED_PROVIDER
class EinsumComputePreprocessor final {
public:
explicit EinsumComputePreprocessor(EinsumEquationPreprocessor& equation_preprocessor,
@ -224,5 +225,5 @@ class EinsumComputePreprocessor final {
// Holds EP-specific assets required for (auxiliary) ops that need to be executed on non-CPU EPs
void* einsum_ep_assets_;
};
#endif
} // namespace onnxruntime

2
onnxruntime/core/providers/cpu/nn/batch_norm_helper.h
View file

@ -3,8 +3,10 @@
#pragma once
#ifndef SHARED_PROVIDER
#include "core/common/status.h"
#include "core/framework/tensor.h"
#endif
#include <sstream>
namespace onnxruntime {

2
onnxruntime/core/providers/cpu/nn/conv_attributes.h
View file

@ -14,7 +14,7 @@ namespace onnxruntime {
// A helper struct holding attributes for Conv-family ops
struct ConvAttributes {
explicit ConvAttributes(const OpNodeProtoHelper<ProtoHelperNodeContext>& info) {
explicit ConvAttributes(const OpKernelInfo& info) {
std::string auto_pad_str;
auto status = info.GetAttr<std::string>("auto_pad", &auto_pad_str);
if (status.IsOK()) {

2
onnxruntime/core/providers/cpu/nn/conv_transpose_attributes.h
View file

@ -22,7 +22,7 @@
namespace onnxruntime {
struct ConvTransposeAttributes : public ConvAttributes {
explicit ConvTransposeAttributes(const OpNodeProtoHelper<ProtoHelperNodeContext>& info)
explicit ConvTransposeAttributes(const OpKernelInfo& info)
: ConvAttributes(info),
output_padding(info.GetAttrsOrDefault<int64_t>("output_padding")),
output_shape(info.GetAttrsOrDefault<int64_t>("output_shape")) {

2
onnxruntime/core/providers/cpu/nn/instance_norm_helper.h
View file

@ -3,8 +3,10 @@
#pragma once
#ifndef SHARED_PROVIDER
#include "core/common/status.h"
#include "core/framework/tensor.h"
#endif
#include <sstream>
namespace onnxruntime {

4
onnxruntime/core/providers/cpu/nn/pool_attributes.h
View file

@ -4,10 +4,12 @@
#pragma once
#include <cmath>
#ifndef SHARED_PROVIDER
#include "core/common/common.h"
#include "core/framework/op_node_proto_helper.h"
#include "core/framework/tensor_shape.h"
#include "core/providers/common.h"
#endif
namespace onnxruntime {
@ -17,7 +19,7 @@ struct PoolAttributes {
return op_name == "GlobalAveragePool" || op_name == "GlobalMaxPool" || op_name == "GlobalLpPool";
}
PoolAttributes(const OpNodeProtoHelper<ProtoHelperNodeContext>& info,
PoolAttributes(const OpKernelInfo& info,
const std::string& op_name, int start_version)
: global_pooling(IsGlobalPooling(op_name)) {
if (global_pooling) {

7
onnxruntime/core/providers/cpu/nn/pool_base.h
View file

@ -4,10 +4,12 @@
#pragma once
#include <cmath>
#ifndef SHARED_PROVIDER
#include "core/common/common.h"
#include "core/framework/op_kernel.h"
#include "core/providers/cpu/nn/pool_attributes.h"
#include "core/util/math.h"
#endif
#include "core/providers/cpu/nn/pool_attributes.h"
#include "core/mlas/inc/mlas.h"
namespace onnxruntime {
@ -106,8 +108,7 @@ class PoolBase {
protected:
PoolBase(const OpKernelInfo& info)
: op_name_(info.GetKernelDef().OpName().rfind("QLinear", 0) != 0 ?
info.GetKernelDef().OpName() : info.GetKernelDef().OpName().substr(7)),
: op_name_(info.GetKernelDef().OpName().rfind("QLinear", 0) != 0 ? info.GetKernelDef().OpName() : info.GetKernelDef().OpName().substr(7)),
pool_attrs_(info, op_name_, GetStartVersion(info)) {
}

9
onnxruntime/core/providers/cpu/object_detection/non_max_suppression.h
View file

@ -17,16 +17,17 @@ class NonMaxSuppressionBase {
ORT_ENFORCE(0 == center_point_box_ || 1 == center_point_box_, "center_point_box only support 0 or 1");
}
int64_t GetCenterPointBox() const {
return center_point_box_;
}
public:
static Status PrepareCompute(OpKernelContext* ctx, PrepareContext& pc);
static Status GetThresholdsFromInputs(const PrepareContext& pc,
int64_t& max_output_boxes_per_class,
float& iou_threshold,
float& score_threshold);
int64_t GetCenterPointBox() const {
return center_point_box_;
}
private:
int64_t center_point_box_;
};

18
onnxruntime/core/providers/cpu/reduction/reduction_ops.cc
View file

@ -586,9 +586,9 @@ Status ReduceSum<T>::Compute(OpKernelContext* ctx) const {
}
template <typename T>
Tensor ReduceSum<T>::Impl(const Tensor& input, const std::vector<int64_t>& reduce_axes,
AllocatorPtr allocator, concurrency::ThreadPool* tp, bool keep_dims,
const TensorShape* input_shape_override) {
std::unique_ptr<Tensor> ReduceSum<T>::Impl(const Tensor& input, const std::vector<int64_t>& reduce_axes,
AllocatorPtr allocator, concurrency::ThreadPool* tp, bool keep_dims,
const TensorShape* input_shape_override) {
std::vector<int64_t> axes;
auto reduced_dims = input.Shape().GetDims();
std::vector<int64_t> output_shape;
@ -597,10 +597,10 @@ Tensor ReduceSum<T>::Impl(const Tensor& input, const std::vector<int64_t>& reduc
SetupForReduce(&input, reduce_axes, axes, new_input_shape, output_shape, empty_reduce, input_shape_override);
if (empty_reduce) {
Tensor output(input.DataType(), keep_dims ? output_shape : std::vector<int64_t>(), allocator);
auto output = make_unique<Tensor>(input.DataType(), keep_dims ? output_shape : std::vector<int64_t>(), allocator);
if (new_input_shape.Size() == 1) {
const T* from_data = input.template Data<T>();
T* to_data = output.template MutableData<T>();
T* to_data = output->template MutableData<T>();
*to_data = *from_data;
} else {
ORT_ENFORCE(keep_dims,
@ -613,15 +613,15 @@ Tensor ReduceSum<T>::Impl(const Tensor& input, const std::vector<int64_t>& reduc
if (keep_dims) {
ResultsNoTransposePrepareForReduce last_results;
Tensor output(input.DataType(), output_shape, allocator);
NoTransposeReduce<T, ReduceAggregatorSum<T>>(&output, new_input_shape, input, axes, tp, last_results);
auto output = make_unique<Tensor>(input.DataType(), output_shape, allocator);
NoTransposeReduce<T, ReduceAggregatorSum<T>>(output.get(), new_input_shape, input, axes, tp, last_results);
return output;
} else {
ResultsNoTransposePrepareForReduce last_results;
std::vector<int64_t> dropped_axes;
DropDimensions(output_shape, axes, dropped_axes);
Tensor output(input.DataType(), dropped_axes, allocator);
NoTransposeReduce<T, ReduceAggregatorSum<T>>(&output, new_input_shape, input, axes, tp, last_results);
auto output = make_unique<Tensor>(input.DataType(), dropped_axes, allocator);
NoTransposeReduce<T, ReduceAggregatorSum<T>>(output.get(), new_input_shape, input, axes, tp, last_results);
return output;
}
}

6
onnxruntime/core/providers/cpu/reduction/reduction_ops.h
View file

@ -472,9 +472,9 @@ class ReduceSum final : public ReduceKernel<true> {
// For external calls requiring ReduceSum implementation - will return the reduced output.
//`input_shape_override` overrides the shape of `input` for compute purposes.
static Tensor Impl(const Tensor& input, const std::vector<int64_t>& reduce_axes,
AllocatorPtr allocator, concurrency::ThreadPool* tp, bool keep_dims,
const TensorShape* input_shape_override = nullptr);
static std::unique_ptr<Tensor> Impl(const Tensor& input, const std::vector<int64_t>& reduce_axes,
AllocatorPtr allocator, concurrency::ThreadPool* tp, bool keep_dims,
const TensorShape* input_shape_override = nullptr);
};
template <typename T>

4
onnxruntime/core/providers/cpu/tensor/reshape_helper.h
View file

@ -3,10 +3,6 @@
#pragma once
//#include "core/common/common.h"
//#include "core/framework/tensor.h"
//#include "gsl/gsl"
namespace onnxruntime {
// Verify and convert unknown dim during reshape

2
onnxruntime/core/providers/cpu/tensor/tile.h
View file

@ -1,9 +1,11 @@
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
#ifndef SHARED_PROVIDER
#include "core/common/common.h"
#include "core/framework/op_kernel.h"
#include "core/util/math_cpuonly.h"
#endif
namespace onnxruntime {

20
onnxruntime/core/providers/cuda/controlflow/if.cc
View file

@ -1,7 +1,7 @@
#if 0
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
#include "core/providers/shared_library/provider_api.h"
#include "core/providers/cuda/controlflow/if.h"
#include "core/providers/cuda/cuda_fwd.h"
@ -16,8 +16,8 @@ ONNX_OPERATOR_VERSIONED_KERNEL_EX(If,
kOnnxDomain,
1, 10,
kCudaExecutionProvider,
KernelDefBuilder()
.InputMemoryType<OrtMemTypeCPUInput>(0) // 'cond' needs to be on CPU
(*KernelDefBuilder::Create())
.InputMemoryType(OrtMemTypeCPUInput, 0) // 'cond' needs to be on CPU
.TypeConstraint("B", DataTypeImpl::GetTensorType<bool>())
.TypeConstraint("V", DataTypeImpl::AllFixedSizeTensorTypes()),
If);
@ -28,8 +28,8 @@ ONNX_OPERATOR_VERSIONED_KERNEL_EX(If,
kOnnxDomain,
11, 12,
kCudaExecutionProvider,
KernelDefBuilder()
.InputMemoryType<OrtMemTypeCPUInput>(0) // 'cond' needs to be on CPU
(*KernelDefBuilder::Create())
.InputMemoryType(OrtMemTypeCPUInput, 0) // 'cond' needs to be on CPU
.TypeConstraint("B", DataTypeImpl::GetTensorType<bool>())
.TypeConstraint("V", DataTypeImpl::AllFixedSizeTensorTypes()),
If);
@ -41,22 +41,24 @@ ONNX_OPERATOR_KERNEL_EX(If,
kOnnxDomain,
13,
kCudaExecutionProvider,
KernelDefBuilder()
.InputMemoryType<OrtMemTypeCPUInput>(0) // 'cond' needs to be on CPU
(*KernelDefBuilder::Create())
.InputMemoryType(OrtMemTypeCPUInput, 0) // 'cond' needs to be on CPU
.TypeConstraint("B", DataTypeImpl::GetTensorType<bool>())
.TypeConstraint("V", DataTypeImpl::AllFixedSizeTensorTypes()),
If);
If::If(const OpKernelInfo& info) : OpKernel(info), if_cpu_{CreateOpKernel_CPU_If(info)} {
}
Status If::Compute(OpKernelContext* ctx) const {
// call the base CPU version.
// we have this CUDA implementation so the inputs/outputs stay on GPU where possible.
// the logic to run the subgraph must be on CPU either way.
// technically we don't need this override of Compute, but it will be optimized out and it's easier to debug
// that this implementation is being called with it.
auto status = onnxruntime::If::Compute(ctx);
auto status = if_cpu_->Compute(ctx);
return status;
}
} // namespace cuda
} // namespace onnxruntime
#endif

6
onnxruntime/core/providers/cuda/controlflow/if.h
View file

@ -14,11 +14,13 @@ class SessionState;
namespace cuda {
// Use the CPU implementation for the logic
class If final : public onnxruntime::If {
class If final : public OpKernel {
public:
If(const OpKernelInfo& info) : onnxruntime::If(info) {}
If(const OpKernelInfo& info);
Status Compute(OpKernelContext* ctx) const override;
std::unique_ptr<OpKernel> if_cpu_;
};
} // namespace cuda
} // namespace onnxruntime

31
onnxruntime/core/providers/cuda/controlflow/loop.cc
View file

@ -1,10 +1,11 @@
#if 0
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
#include "core/providers/shared_library/provider_api.h"
#include "core/providers/cuda/controlflow/loop.h"
#include "core/providers/cuda/cuda_common.h"
#include "core/providers/cuda/cuda_fwd.h"
#include "core/framework/ml_value.h"
using namespace ONNX_NAMESPACE;
using namespace onnxruntime::common;
@ -16,9 +17,9 @@ ONNX_OPERATOR_VERSIONED_KERNEL_EX(Loop,
kOnnxDomain,
1, 10,
kCudaExecutionProvider,
KernelDefBuilder()
.InputMemoryType<OrtMemTypeCPUInput>(0) // 'M' needs to be on CPU
.InputMemoryType<OrtMemTypeCPUInput>(1) // 'cond' needs to be on CPU
(*KernelDefBuilder::Create())
.InputMemoryType(OrtMemTypeCPUInput, 0) // 'M' needs to be on CPU
.InputMemoryType(OrtMemTypeCPUInput, 1) // 'cond' needs to be on CPU
.TypeConstraint("I", DataTypeImpl::GetTensorType<int64_t>())
.TypeConstraint("B", DataTypeImpl::GetTensorType<bool>())
.TypeConstraint("V", DataTypeImpl::AllFixedSizeTensorTypes()),
@ -29,9 +30,9 @@ ONNX_OPERATOR_VERSIONED_KERNEL_EX(Loop,
kOnnxDomain,
11, 12,
kCudaExecutionProvider,
KernelDefBuilder()
.InputMemoryType<OrtMemTypeCPUInput>(0) // 'M' needs to be on CPU
.InputMemoryType<OrtMemTypeCPUInput>(1) // 'cond' needs to be on CPU
(*KernelDefBuilder::Create())
.InputMemoryType(OrtMemTypeCPUInput, 0) // 'M' needs to be on CPU
.InputMemoryType(OrtMemTypeCPUInput, 1) // 'cond' needs to be on CPU
.TypeConstraint("I", DataTypeImpl::GetTensorType<int64_t>())
.TypeConstraint("B", DataTypeImpl::GetTensorType<bool>())
.TypeConstraint("V", DataTypeImpl::AllFixedSizeTensorTypes()),
@ -44,9 +45,9 @@ ONNX_OPERATOR_KERNEL_EX(Loop,
kOnnxDomain,
13,
kCudaExecutionProvider,
KernelDefBuilder()
.InputMemoryType<OrtMemTypeCPUInput>(0) // 'M' needs to be on CPU
.InputMemoryType<OrtMemTypeCPUInput>(1) // 'cond' needs to be on CPU
(*KernelDefBuilder::Create())
.InputMemoryType(OrtMemTypeCPUInput, 0) // 'M' needs to be on CPU
.InputMemoryType(OrtMemTypeCPUInput, 1) // 'cond' needs to be on CPU
.TypeConstraint("I", DataTypeImpl::GetTensorType<int64_t>())
.TypeConstraint("B", DataTypeImpl::GetTensorType<bool>())
.TypeConstraint("V", DataTypeImpl::AllFixedSizeTensorTypes()),
@ -70,7 +71,7 @@ static Status ConcatenateGpuOutput(void* stream, std::vector<OrtValue>& per_iter
}
CUDA_RETURN_IF_ERROR(cudaMemcpyAsync(cur_output, iteration_data.DataRaw(), bytes_per_iteration,
cudaMemcpyDeviceToDevice, static_cast<cudaStream_t>(stream)));
cudaMemcpyDeviceToDevice, static_cast<cudaStream_t>(stream)));
cur_output = static_cast<void*>((static_cast<gsl::byte*>(cur_output) + bytes_per_iteration));
}
@ -81,9 +82,8 @@ static Status ConcatenateGpuOutput(void* stream, std::vector<OrtValue>& per_iter
return Status::OK();
}
Loop::Loop(const OpKernelInfo& info) : onnxruntime::Loop(info) {
SetConcatOutputFunc(ConcatenateGpuOutput);
SetComputeStream(static_cast<void*>(info.GetExecutionProvider()->GetComputeStream()));
Loop::Loop(const OpKernelInfo& info) : OpKernel(info) {
cpu_loop_ = onnxruntime::Loop::Create(info, ConcatenateGpuOutput, static_cast<void*>(info.GetExecutionProvider()->GetComputeStream()));
}
Status Loop::Compute(OpKernelContext* ctx) const {
@ -92,10 +92,9 @@ Status Loop::Compute(OpKernelContext* ctx) const {
// the logic to run the subgraph must be on CPU either way.
// technically we don't need this override of Compute, but it will be optimized out and it's easier to debug
// that this implementation is being called with it.
auto status = onnxruntime::Loop::Compute(ctx);
auto status = cpu_loop_->Compute(ctx);
return status;
}
} // namespace cuda
} // namespace onnxruntime
#endif

5
onnxruntime/core/providers/cuda/controlflow/loop.h
View file

@ -10,11 +10,14 @@ namespace onnxruntime {
namespace cuda {
// Use the CPU implementation for the logic
class Loop final : public onnxruntime::Loop {
class Loop final : public OpKernel {
public:
Loop(const OpKernelInfo& info);
Status Compute(OpKernelContext* ctx) const override;
private:
std::unique_ptr<OpKernel> cpu_loop_;
};
} // namespace cuda
} // namespace onnxruntime

25
onnxruntime/core/providers/cuda/controlflow/scan.cc
View file

@ -1,11 +1,13 @@
#if 0
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
#include "core/providers/shared_library/provider_api.h"
#include "core/providers/cuda/controlflow/scan.h"
#include "core/providers/cuda/cuda_common.h"
#include "core/providers/cuda/tensor/transpose.h"
#include "core/framework/ml_value.h"
#include "core/framework/ort_value_tensor_slicer.cc"
using namespace ONNX_NAMESPACE;
using namespace onnxruntime::common;
@ -14,7 +16,7 @@ namespace onnxruntime {
namespace cuda {
template <>
Scan<8>::Scan(const OpKernelInfo& info) : onnxruntime::Scan<8>(info) {
Scan<8>::Scan(const OpKernelInfo& info) : OpKernel(info), scan_cpu_{static_cast<onnxruntime::Scan<8>*>(CreateOpKernel_CPU_Scan_8(info).release())} {
scan::detail::DeviceHelpers helpers;
helpers.set_data_to_zero_func = [](void* data, size_t size_in_bytes) -> Status {
@ -23,11 +25,11 @@ Scan<8>::Scan(const OpKernelInfo& info) : onnxruntime::Scan<8>(info) {
};
// copy into base class
SetDeviceHelpers(helpers);
scan_cpu_->SetDeviceHelpers(helpers);
}
template <>
Scan<9>::Scan(const OpKernelInfo& info) : onnxruntime::Scan<9>(info) {
Scan<9>::Scan(const OpKernelInfo& info) : OpKernel(info), scan_cpu_{static_cast<onnxruntime::Scan<9>*>(CreateOpKernel_CPU_Scan_9(info).release())} {
scan::detail::DeviceHelpers helpers;
helpers.transpose_func = [this](const std::vector<size_t>& permutations, const Tensor& input, Tensor& output) {
@ -38,7 +40,7 @@ Scan<9>::Scan(const OpKernelInfo& info) : onnxruntime::Scan<9>(info) {
};
// copy into base class
SetDeviceHelpers(helpers);
scan_cpu_->SetDeviceHelpers(helpers);
}
template <>
@ -48,7 +50,7 @@ Status Scan<8>::Compute(OpKernelContext* ctx) const {
// the logic to run the subgraph must be on CPU either way.
// technically we don't need this override of Compute, but it will be optimized out and it's easier to debug
// that this implementation is being called with it.
auto status = onnxruntime::Scan<8>::Compute(ctx);
auto status = scan_cpu_->Compute(ctx);
return status;
}
@ -59,7 +61,7 @@ Status Scan<9>::Compute(OpKernelContext* ctx) const {
// the logic to run the subgraph must be on CPU either way.
// technically we don't need this override of Compute, but it will be optimized out and it's easier to debug
// that this implementation is being called with it.
auto status = onnxruntime::Scan<9>::Compute(ctx);
auto status = scan_cpu_->Compute(ctx);
return status;
}
@ -67,8 +69,8 @@ ONNX_OPERATOR_VERSIONED_KERNEL_EX(Scan,
kOnnxDomain,
8, 8,
kCudaExecutionProvider,
KernelDefBuilder()
.InputMemoryType<OrtMemTypeCPUInput>(0) // 'sequence_lens' needs to be on CPU
(*KernelDefBuilder::Create())
.InputMemoryType(OrtMemTypeCPUInput, 0) // 'sequence_lens' needs to be on CPU
.TypeConstraint("I", DataTypeImpl::GetTensorType<int64_t>())
.TypeConstraint("V", DataTypeImpl::AllTensorTypes()),
Scan<8>);
@ -77,7 +79,7 @@ ONNX_OPERATOR_VERSIONED_KERNEL_EX(Scan,
kOnnxDomain,
9, 10,
kCudaExecutionProvider,
KernelDefBuilder()
(*KernelDefBuilder::Create())
.TypeConstraint("I", DataTypeImpl::GetTensorType<int64_t>())
.TypeConstraint("V", DataTypeImpl::AllFixedSizeTensorTypes()),
Scan<9>);
@ -87,11 +89,10 @@ ONNX_OPERATOR_KERNEL_EX(Scan,
kOnnxDomain,
11,
kCudaExecutionProvider,
KernelDefBuilder()
(*KernelDefBuilder::Create())
.TypeConstraint("I", DataTypeImpl::GetTensorType<int64_t>())
.TypeConstraint("V", DataTypeImpl::AllFixedSizeTensorTypes()),
Scan<9>);
} // namespace cuda
} // namespace onnxruntime
#endif

6
onnxruntime/core/providers/cuda/controlflow/scan.h
View file

@ -5,7 +5,6 @@
#include <functional>
#include "gsl/gsl"
#include "core/common/common.h"
#include "core/providers/cpu/controlflow/scan.h"
namespace onnxruntime {
@ -15,11 +14,14 @@ namespace cuda {
// Use the CPU implementation for the logic
template <int OpSet>
class Scan final : public onnxruntime::Scan<OpSet> {
class Scan final : public OpKernel {
public:
Scan(const OpKernelInfo& info);
Status Compute(OpKernelContext* ctx) const override;
private:
std::unique_ptr<onnxruntime::Scan<OpSet>> scan_cpu_;
};
} // namespace cuda
} // namespace onnxruntime

2
onnxruntime/core/providers/cuda/cuda_allocator.cc
View file

@ -4,7 +4,6 @@
#include "cuda_allocator.h"
#include "cuda_common.h"
#include "core/framework/allocatormgr.h"
//#include "core/framework/session_state.h"
#include "cuda_fence.h"
#include "gpu_data_transfer.h"
@ -71,7 +70,6 @@ void* CUDAExternalAllocator::Alloc(size_t size) {
// review(codemzs): ORT_ENFORCE does not seem appropiate.
ORT_ENFORCE(p != nullptr);
}
return p;

3
onnxruntime/core/providers/cuda/cuda_call.cc
View file

@ -3,9 +3,6 @@
#include "core/providers/shared_library/provider_api.h"
#include "shared_inc/cuda_call.h"
//#include "core/common/common.h"
//#include "core/common/status.h"
//#include "core/common/logging/logging.h"
#ifdef _WIN32
#else // POSIX

1
onnxruntime/core/providers/cuda/cuda_common.h
View file

@ -10,7 +10,6 @@
#include "core/providers/cuda/shared_inc/cuda_call.h"
#include "core/providers/cuda/shared_inc/fast_divmod.h"
#include "gsl/gsl"
//#include "core/util/math.h"
// Can't include "core/util/math.h" in a provider, so this is the part we need for cuda:
namespace onnxruntime {

6
onnxruntime/core/providers/cuda/cuda_execution_provider.cc
View file

@ -3,12 +3,6 @@
#include "core/providers/shared_library/provider_api.h"
#include "core/providers/cuda/cuda_execution_provider.h"
//#include "core/framework/compute_capability.h"
//#include "core/framework/fallback_cpu_capability.h"
//#include "core/framework/kernel_registry.h"
//#include "core/framework/memcpy.h"
//#include "core/framework/provider_options_utils.h"
//#include "core/graph/graph_utils.h"
#include "core/providers/cuda/cuda_common.h"
#include "core/providers/cuda/cuda_allocator.h"
#include "core/providers/cuda/cuda_fence.h"

1
onnxruntime/core/providers/cuda/cuda_execution_provider.h
View file

@ -6,7 +6,6 @@
#include <set>
#include <vector>
//#include "core/graph/constants.h"
#include "core/framework/allocatormgr.h"
#include "core/framework/arena_extend_strategy.h"
#include "core/framework/execution_provider.h"

1
onnxruntime/core/providers/cuda/cuda_fence.cc
View file

@ -1,7 +1,6 @@
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
//#include "core/graph/constants.h"
#include "core/providers/cuda/cuda_common.h"
#include "core/providers/cuda/gpu_data_transfer.h"
#include "core/providers/cuda/cuda_fence.h"

3
onnxruntime/core/providers/cuda/cuda_fence.h
View file

@ -3,9 +3,6 @@
#pragma once
//#include "core/framework/tensor.h"
//#include "core/graph/basic_types.h"
//#include "core/providers/cuda/cuda_common.h"
#include "core/framework/fence.h"
namespace onnxruntime {

2
onnxruntime/core/providers/cuda/cuda_fwd.h
View file

@ -3,8 +3,6 @@
#pragma once
//#include "core/framework/op_kernel.h"
namespace onnxruntime {
namespace cuda {
template <typename T>

3
onnxruntime/core/providers/cuda/cuda_kernel.h
View file

@ -3,9 +3,6 @@
#pragma once
//#include "core/common/status.h"
//#include "core/framework/data_transfer_manager.h"
//#include "core/framework/op_kernel.h"
#include "core/providers/cuda/cuda_common.h"
#include "core/providers/cuda/cuda_execution_provider.h"
#include "core/providers/cuda/cuda_fwd.h"

144
onnxruntime/core/providers/cuda/cuda_provider_factory.cc
View file

@ -11,9 +11,9 @@
#include "core/common/make_unique.h"
#include "core/providers/cuda/cuda_execution_provider.h"
//#include "core/session/abi_session_options_impl.h"
//#include "core/session/onnxruntime_c_api.h"
//#include "core/session/ort_apis.h"
#include "core/providers/cuda/cuda_allocator.h"
#include "core/providers/cuda/gpu_data_transfer.h"
#include "core/providers/cuda/math/unary_elementwise_ops_impl.h"
using namespace onnxruntime;
@ -40,72 +40,86 @@ std::shared_ptr<IExecutionProviderFactory> CreateExecutionProviderFactory_CUDA(c
} // namespace onnxruntime
#if 0
ORT_API_STATUS_IMPL(OrtSessionOptionsAppendExecutionProvider_CUDA, _In_ OrtSessionOptions* options, int device_id) {
CUDAExecutionProviderInfo info{};
info.device_id = gsl::narrow<OrtDevice::DeviceId>(device_id);
options->provider_factories.push_back(onnxruntime::CreateExecutionProviderFactory_CUDA(info));
return nullptr;
}
ORT_API_STATUS_IMPL(OrtApis::SessionOptionsAppendExecutionProvider_CUDA,
_In_ OrtSessionOptions* options, _In_ const OrtCUDAProviderOptions* cuda_options) {
CUDAExecutionProviderInfo info{};
info.device_id = gsl::narrow<OrtDevice::DeviceId>(cuda_options->device_id);
info.cuda_mem_limit = cuda_options->cuda_mem_limit;
info.arena_extend_strategy = static_cast<onnxruntime::ArenaExtendStrategy>(cuda_options->arena_extend_strategy);
info.cudnn_conv_algo_search = cuda_options->cudnn_conv_algo_search;
info.do_copy_in_default_stream = cuda_options->do_copy_in_default_stream;
info.has_user_compute_stream = cuda_options->has_user_compute_stream;
info.user_compute_stream = cuda_options->user_compute_stream;
options->provider_factories.push_back(onnxruntime::CreateExecutionProviderFactory_CUDA(info));
return nullptr;
}
ORT_API_STATUS_IMPL(OrtApis::SetCurrentGpuDeviceId, _In_ int device_id) {
int num_devices;
auto cuda_err = cudaGetDeviceCount(&num_devices);
if (cuda_err != cudaSuccess) {
return CreateStatus(ORT_FAIL, "Failed to set device id since cudaGetDeviceCount failed.");
}
if (device_id >= num_devices) {
std::ostringstream ostr;
ostr << "Invalid device id. Device id should be less than total number of devices (" << num_devices << ")";
return CreateStatus(ORT_INVALID_ARGUMENT, ostr.str().c_str());
}
cuda_err = cudaSetDevice(device_id);
if (cuda_err != cudaSuccess) {
return CreateStatus(ORT_FAIL, "Failed to set device id.");
}
return nullptr;
}
ORT_API_STATUS_IMPL(OrtApis::GetCurrentGpuDeviceId, _In_ int* device_id) {
auto cuda_err = cudaGetDevice(device_id);
if (cuda_err != cudaSuccess) {
return CreateStatus(ORT_FAIL, "Failed to get device id.");
}
return nullptr;
}
#endif
namespace onnxruntime {
#if 0
struct ProviderInfo_CUDA_Impl : ProviderInfo_CUDA {
std::vector<std::string> GetAvailableDevices() const override {
InferenceEngine::Core ie_core;
return ie_core.GetAvailableDevices();
OrtStatus* SetCurrentGpuDeviceId(_In_ int device_id) override {
int num_devices;
auto cuda_err = cudaGetDeviceCount(&num_devices);
if (cuda_err != cudaSuccess) {
return CreateStatus(ORT_FAIL, "Failed to set device id since cudaGetDeviceCount failed.");
}
if (device_id >= num_devices) {
std::ostringstream ostr;
ostr << "Invalid device id. Device id should be less than total number of devices (" << num_devices << ")";
return CreateStatus(ORT_INVALID_ARGUMENT, ostr.str().c_str());
}
cuda_err = cudaSetDevice(device_id);
if (cuda_err != cudaSuccess) {
return CreateStatus(ORT_FAIL, "Failed to set device id.");
}
return nullptr;
}
OrtStatus* GetCurrentGpuDeviceId(_In_ int* device_id) override {
auto cuda_err = cudaGetDevice(device_id);
if (cuda_err != cudaSuccess) {
return CreateStatus(ORT_FAIL, "Failed to get device id.");
}
return nullptr;
}
std::unique_ptr<IAllocator> CreateCUDAAllocator(int16_t device_id, const char* name) override {
return onnxruntime::make_unique<CUDAAllocator>(device_id, name);
}
std::unique_ptr<IAllocator> CreateCUDAPinnedAllocator(int16_t device_id, const char* name) override {
return onnxruntime::make_unique<CUDAPinnedAllocator>(device_id, name);
}
std::unique_ptr<IDataTransfer> CreateGPUDataTransfer(void* stream) override {
return onnxruntime::make_unique<GPUDataTransfer>(static_cast<cudaStream_t>(stream));
}
void cuda__Impl_Cast(void* stream, const int64_t* input_data, int32_t* output_data, size_t count) override {
return cuda::Impl_Cast(static_cast<cudaStream_t>(stream), input_data, output_data, count);
}
void cuda__Impl_Cast(void* stream, const int32_t* input_data, int64_t* output_data, size_t count) override {
return cuda::Impl_Cast(static_cast<cudaStream_t>(stream), input_data, output_data, count);
}
bool CudaCall_false(int retCode, const char* exprString, const char* libName, int successCode, const char* msg) override { return CudaCall<cudaError, false>(cudaError(retCode), exprString, libName, cudaError(successCode), msg); }
bool CudaCall_true(int retCode, const char* exprString, const char* libName, int successCode, const char* msg) override { return CudaCall<cudaError, true>(cudaError(retCode), exprString, libName, cudaError(successCode), msg); }
void CopyGpuToCpu(void* dst_ptr, const void* src_ptr, const size_t size, const OrtMemoryInfo& dst_location, const OrtMemoryInfo& src_location) override {
ORT_ENFORCE(dst_location.device.Type() == OrtDevice::CPU);
// Current CUDA device.
int device;
CUDA_CALL(cudaGetDevice(&device));
if (device != src_location.id) {
// Need to switch to the allocating device.
CUDA_CALL(cudaSetDevice(src_location.id));
// Copy from GPU to CPU.
CUDA_CALL(cudaMemcpy(dst_ptr, src_ptr, size, cudaMemcpyDeviceToHost));
// Switch back to current device.
CUDA_CALL(cudaSetDevice(device));
} else {
// Copy from GPU to CPU.
CUDA_CALL(cudaMemcpy(dst_ptr, src_ptr, size, cudaMemcpyDeviceToHost));
}
}
// Used only by slice_concatenate_test.cc
void cudaMemcpy_HostToDevice(void* dst, const void* src, size_t count) override { cudaMemcpy(dst, src, count, cudaMemcpyHostToDevice); }
} g_info;
#endif
struct CUDA_Provider : Provider {
// const void* GetInfo() override { return &g_info; }
void* GetInfo() override { return &g_info; }
std::shared_ptr<IExecutionProviderFactory> CreateExecutionProviderFactory(const void* void_params) override {
auto params = reinterpret_cast<const OrtCUDAProviderOptions*>(void_params);
@ -116,12 +130,13 @@ struct CUDA_Provider : Provider {
info.arena_extend_strategy = static_cast<onnxruntime::ArenaExtendStrategy>(params->arena_extend_strategy);
info.cudnn_conv_algo_search = params->cudnn_conv_algo_search;
info.do_copy_in_default_stream = params->do_copy_in_default_stream;
info.has_user_compute_stream = params->has_user_compute_stream;
info.user_compute_stream = params->user_compute_stream;
return std::make_shared<CUDAProviderFactory>(info);
}
void Shutdown() override {
// openvino_ep::BackendManager::ReleaseGlobalContext();
}
} g_provider;
@ -134,4 +149,3 @@ ORT_API(onnxruntime::Provider*, GetProvider) {
return &onnxruntime::g_provider;
}
}

4
onnxruntime/core/providers/cuda/cudnn_common.h
View file

@ -5,8 +5,6 @@
#include <cfloat>
//#include "core/common/logging/logging.h"
//#include "core/framework/tensor.h"
#include "core/providers/cuda/cuda_common.h"
namespace onnxruntime {
@ -127,7 +125,7 @@ struct Consts<half> {
};
#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
template<>
template <>
struct Consts<nv_bfloat16> {
static const float Zero;
static const float One;

19
onnxruntime/core/providers/cuda/generator/constant_of_shape.cc
View file

@ -1,10 +1,6 @@
#if 0
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
#include "constant_of_shape.h"
#include "core/providers/common.h"
#include "gsl/gsl"
using namespace ::onnxruntime::common;
using namespace ONNX_NAMESPACE;
@ -16,8 +12,8 @@ ONNX_OPERATOR_KERNEL_EX(
kOnnxDomain,
9,
kCudaExecutionProvider,
KernelDefBuilder()
.InputMemoryType<OrtMemTypeCPUInput>(0)
(*KernelDefBuilder::Create())
.InputMemoryType(OrtMemTypeCPUInput, 0)
.TypeConstraint("T1", DataTypeImpl::GetTensorType<int64_t>())
.TypeConstraint("T2", DataTypeImpl::AllFixedSizeTensorTypes()),
ConstantOfShape);
@ -30,11 +26,11 @@ Status ConstantOfShape::ComputeInternal(OpKernelContext* ctx) const {
const void* value_ptr = GetValuePtr();
const auto element_size = output_tensor->DataType()->Size();
#define CASE(TYPE) \
case sizeof(TYPE): \
if (size > 0) { \
cuda::Fill(Stream(), reinterpret_cast<TYPE*>(output_data), *(reinterpret_cast<const TYPE*>(value_ptr)), size); \
} \
#define CASE(TYPE) \
case sizeof(TYPE): \
if (size > 0) { \
cuda::Fill(Stream(), reinterpret_cast<TYPE*>(output_data), *(reinterpret_cast<const TYPE*>(value_ptr)), size); \
} \
break;
switch (element_size) {
@ -52,4 +48,3 @@ Status ConstantOfShape::ComputeInternal(OpKernelContext* ctx) const {
} // namespace cuda
} // namespace onnxruntime
#endif

2
onnxruntime/core/providers/cuda/generator/constant_of_shape.h
View file

@ -3,8 +3,6 @@
#pragma once
#include "core/common/common.h"
#include "core/framework/data_types.h"
#include "core/providers/cuda/cuda_kernel.h"
#include "core/providers/cpu/generator/constant_of_shape_base.h"
#include "core/providers/cuda/shared_inc/cuda_utils.h"

14
onnxruntime/core/providers/cuda/generator/range.cc
View file

@ -1,8 +1,7 @@
#if 0
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
#include "core/framework/tensorprotoutils.h"
#include "core/providers/shared_library/provider_api.h"
#include "core/providers/cuda/cuda_common.h"
#include "range.h"
#include "range_impl.h"
@ -19,10 +18,10 @@ ONNX_OPERATOR_KERNEL_EX(
kOnnxDomain,
11,
kCudaExecutionProvider,
KernelDefBuilder()
.InputMemoryType<OrtMemTypeCPUInput>(0) // start
.InputMemoryType<OrtMemTypeCPUInput>(1) // limit
.InputMemoryType<OrtMemTypeCPUInput>(2) // delta
(*KernelDefBuilder::Create())
.InputMemoryType(OrtMemTypeCPUInput, 0) // start
.InputMemoryType(OrtMemTypeCPUInput, 1) // limit
.InputMemoryType(OrtMemTypeCPUInput, 2) // delta
.TypeConstraint("T", {DataTypeImpl::GetTensorType<float>(),
DataTypeImpl::GetTensorType<double>(),
DataTypeImpl::GetTensorType<int16_t>(),
@ -52,7 +51,7 @@ static Status ComputeRange(cudaStream_t stream, OpKernelContext* ctx) {
delta_tensor_ptr->Shape());
}
// Start, Limit and Delta are stored in CPU.
// Start, Limit and Delta are stored in CPU.
T start = *(start_tensor.template Data<T>());
T limit = *(limit_tensor.template Data<T>());
@ -105,4 +104,3 @@ Status Range::ComputeInternal(OpKernelContext* ctx) const {
} // namespace cuda
} // namespace onnxruntime
#endif

24
onnxruntime/core/providers/cuda/math/clip.cc
View file

@ -1,8 +1,6 @@
#if 0
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
#include "core/providers/common.h"
#include "core/providers/cuda/math/clip.h"
#include "core/providers/cuda/math/clip_impl.h"
@ -16,7 +14,7 @@ ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX(
10,
float,
kCudaExecutionProvider,
KernelDefBuilder()
(*KernelDefBuilder::Create())
.TypeConstraint("T", DataTypeImpl::GetTensorType<float>()),
Clip_6<float>);
@ -25,9 +23,9 @@ ONNX_OPERATOR_VERSIONED_KERNEL_EX(
kOnnxDomain,
11, 11,
kCudaExecutionProvider,
KernelDefBuilder()
.InputMemoryType<OrtMemTypeCPUInput>(1)
.InputMemoryType<OrtMemTypeCPUInput>(2)
(*KernelDefBuilder::Create())
.InputMemoryType(OrtMemTypeCPUInput, 1)
.InputMemoryType(OrtMemTypeCPUInput, 2)
.TypeConstraint("T", DataTypeImpl::GetTensorType<float>()),
Clip);
@ -36,9 +34,9 @@ ONNX_OPERATOR_VERSIONED_KERNEL_EX(
kOnnxDomain,
12, 12,
kCudaExecutionProvider,
KernelDefBuilder()
.InputMemoryType<OrtMemTypeCPUInput>(1)
.InputMemoryType<OrtMemTypeCPUInput>(2)
(*KernelDefBuilder::Create())
.InputMemoryType(OrtMemTypeCPUInput, 1)
.InputMemoryType(OrtMemTypeCPUInput, 2)
.TypeConstraint("T", BuildKernelDefConstraints<float, double, int8_t, uint8_t, int64_t, uint64_t>()),
Clip);
@ -47,13 +45,12 @@ ONNX_OPERATOR_KERNEL_EX(
kOnnxDomain,
13,
kCudaExecutionProvider,
KernelDefBuilder()
.InputMemoryType<OrtMemTypeCPUInput>(1)
.InputMemoryType<OrtMemTypeCPUInput>(2)
(*KernelDefBuilder::Create())
.InputMemoryType(OrtMemTypeCPUInput, 1)
.InputMemoryType(OrtMemTypeCPUInput, 2)
.TypeConstraint("T", BuildKernelDefConstraints<float, double, int8_t, uint8_t, int64_t, uint64_t>()),
Clip);
template <typename T>
Status Clip_6<T>::ComputeInternal(OpKernelContext* ctx) const {
const Tensor& X = *ctx->Input<Tensor>(0);
@ -132,4 +129,3 @@ Status Clip::ComputeInternal(OpKernelContext* ctx) const {
} // namespace cuda
} // namespace onnxruntime
#endif

12
onnxruntime/core/providers/cuda/math/cumsum.cc
View file

@ -1,4 +1,3 @@
#if 0
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
@ -15,8 +14,8 @@ ONNX_OPERATOR_VERSIONED_KERNEL_EX(
kOnnxDomain,
11, 13,
kCudaExecutionProvider,
KernelDefBuilder()
.InputMemoryType<OrtMemTypeCPUInput>(1) // 'axis' needs to be on CPU
(*KernelDefBuilder::Create())
.InputMemoryType(OrtMemTypeCPUInput, 1) // 'axis' needs to be on CPU
.TypeConstraint("T", std::vector<MLDataType>{
DataTypeImpl::GetTensorType<int32_t>(),
DataTypeImpl::GetTensorType<int64_t>(),
@ -32,8 +31,8 @@ ONNX_OPERATOR_KERNEL_EX(
kOnnxDomain,
14,
kCudaExecutionProvider,
KernelDefBuilder()
.InputMemoryType<OrtMemTypeCPUInput>(1) // 'axis' needs to be on CPU
(*KernelDefBuilder::Create())
.InputMemoryType(OrtMemTypeCPUInput, 1) // 'axis' needs to be on CPU
.TypeConstraint("T", std::vector<MLDataType>{
DataTypeImpl::GetTensorType<int32_t>(),
DataTypeImpl::GetTensorType<int64_t>(),
@ -41,7 +40,7 @@ ONNX_OPERATOR_KERNEL_EX(
DataTypeImpl::GetTensorType<uint64_t>(),
DataTypeImpl::GetTensorType<float>(),
DataTypeImpl::GetTensorType<double>(),
DataTypeImpl::GetTensorType<MLFloat16>()}) // MLFloat16 is added in opset 14
DataTypeImpl::GetTensorType<MLFloat16>()}) // MLFloat16 is added in opset 14
.TypeConstraint("T2", std::vector<MLDataType>{DataTypeImpl::GetTensorType<int32_t>(), DataTypeImpl::GetTensorType<int64_t>()}),
CumSum);
@ -143,4 +142,3 @@ Status CumSum::ComputeInternal(OpKernelContext* ctx) const {
} // namespace cuda
} // namespace onnxruntime
#endif

73
onnxruntime/core/providers/cuda/math/einsum.cc
View file

@ -1,4 +1,3 @@
#if 0
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
@ -6,6 +5,13 @@
namespace onnxruntime {
// This function must exist due to the C++ base class constructor needing this to be defined for the vtable, but it is never called.
Status Einsum::DeviceCompute(OpKernelContext* context, const std::vector<const Tensor*>& inputs,
AllocatorPtr allocator, concurrency::ThreadPool* tp) const {
assert(false);
return Status::OK();
}
namespace cuda {
ONNX_OPERATOR_KERNEL_EX(
@ -13,11 +19,7 @@ ONNX_OPERATOR_KERNEL_EX(
kOnnxDomain,
12,
kCudaExecutionProvider,
KernelDefBuilder().TypeConstraint("T",
std::vector<MLDataType>{
DataTypeImpl::GetTensorType<float>(),
DataTypeImpl::GetTensorType<double>(),
DataTypeImpl::GetTensorType<MLFloat16>()}),
(*KernelDefBuilder::Create()).TypeConstraint("T", std::vector<MLDataType>{DataTypeImpl::GetTensorType<float>(), DataTypeImpl::GetTensorType<double>(), DataTypeImpl::GetTensorType<MLFloat16>()}),
Einsum);
Status Einsum::Compute(OpKernelContext* context) const {
@ -31,46 +33,46 @@ Status Einsum::DeviceCompute(OpKernelContext* context, const std::vector<const T
EinsumOp::EinsumCudaAssets einsum_cuda_assets(cublas_handle, cuda_ep_);
// EinsumComputePreprocessor section -
auto einsum_compute_preprocessor = EinsumComputePreprocessor(*einsum_equation_preprocessor_, inputs, allocator,
&einsum_cuda_assets);
auto einsum_compute_preprocessor = EinsumComputePreprocessor::Create(*einsum_equation_preprocessor_, inputs, allocator,
&einsum_cuda_assets);
einsum_compute_preprocessor.SetDeviceHelpers(EinsumOp::DeviceHelpers::CudaDeviceHelpers::Diagonal,
EinsumOp::DeviceHelpers::CudaDeviceHelpers::Transpose);
einsum_compute_preprocessor->SetDeviceHelpers(EinsumOp::DeviceHelpers::CudaDeviceHelpers::Diagonal,
EinsumOp::DeviceHelpers::CudaDeviceHelpers::Transpose);
// Compute all required metadata to be used at Einsum compute time and return error status code if one was generated
ORT_RETURN_IF_ERROR(einsum_compute_preprocessor.Run());
ORT_RETURN_IF_ERROR(einsum_compute_preprocessor->Run());
// EinsumComputeProcessor section -
if (inputs[0]->IsDataType<float>()) {
auto einsum_compute_processor = EinsumTypedComputeProcessor<float>(context, allocator, tp,
einsum_compute_preprocessor,
&einsum_cuda_assets);
auto einsum_compute_processor = EinsumTypedComputeProcessor<float>::Create(context, allocator, tp,
*einsum_compute_preprocessor,
&einsum_cuda_assets);
einsum_compute_processor.SetDeviceHelpers(EinsumOp::DeviceHelpers::CudaDeviceHelpers::Transpose,
EinsumOp::DeviceHelpers::CudaDeviceHelpers::MatMul<float>,
EinsumOp::DeviceHelpers::CudaDeviceHelpers::ReduceSum<float>,
EinsumOp::DeviceHelpers::CudaDeviceHelpers::DataCopy);
return einsum_compute_processor.Run();
einsum_compute_processor->SetDeviceHelpers(EinsumOp::DeviceHelpers::CudaDeviceHelpers::Transpose,
EinsumOp::DeviceHelpers::CudaDeviceHelpers::MatMul<float>,
EinsumOp::DeviceHelpers::CudaDeviceHelpers::ReduceSum<float>,
EinsumOp::DeviceHelpers::CudaDeviceHelpers::DataCopy);
return einsum_compute_processor->Run();
} else if (inputs[0]->IsDataType<double>()) {
auto einsum_compute_processor = EinsumTypedComputeProcessor<double>(context, allocator, tp,
einsum_compute_preprocessor,
&einsum_cuda_assets);
auto einsum_compute_processor = EinsumTypedComputeProcessor<double>::Create(context, allocator, tp,
*einsum_compute_preprocessor,
&einsum_cuda_assets);
// Set device specific methods (CPU methods) to be used during processing
einsum_compute_processor.SetDeviceHelpers(EinsumOp::DeviceHelpers::CudaDeviceHelpers::Transpose,
EinsumOp::DeviceHelpers::CudaDeviceHelpers::MatMul<double>,
EinsumOp::DeviceHelpers::CudaDeviceHelpers::ReduceSum<double>,
EinsumOp::DeviceHelpers::CudaDeviceHelpers::DataCopy);
return einsum_compute_processor.Run();
einsum_compute_processor->SetDeviceHelpers(EinsumOp::DeviceHelpers::CudaDeviceHelpers::Transpose,
EinsumOp::DeviceHelpers::CudaDeviceHelpers::MatMul<double>,
EinsumOp::DeviceHelpers::CudaDeviceHelpers::ReduceSum<double>,
EinsumOp::DeviceHelpers::CudaDeviceHelpers::DataCopy);
return einsum_compute_processor->Run();
} else if (inputs[0]->IsDataType<MLFloat16>()) {
auto einsum_compute_processor = EinsumTypedComputeProcessor<MLFloat16>(context, allocator, tp,
einsum_compute_preprocessor,
&einsum_cuda_assets);
auto einsum_compute_processor = EinsumTypedComputeProcessor<MLFloat16>::Create(context, allocator, tp,
*einsum_compute_preprocessor,
&einsum_cuda_assets);
einsum_compute_processor.SetDeviceHelpers(EinsumOp::DeviceHelpers::CudaDeviceHelpers::Transpose,
EinsumOp::DeviceHelpers::CudaDeviceHelpers::MatMul<MLFloat16>,
EinsumOp::DeviceHelpers::CudaDeviceHelpers::ReduceSum<MLFloat16>,
EinsumOp::DeviceHelpers::CudaDeviceHelpers::DataCopy);
return einsum_compute_processor.Run();
einsum_compute_processor->SetDeviceHelpers(EinsumOp::DeviceHelpers::CudaDeviceHelpers::Transpose,
EinsumOp::DeviceHelpers::CudaDeviceHelpers::MatMul<MLFloat16>,
EinsumOp::DeviceHelpers::CudaDeviceHelpers::ReduceSum<MLFloat16>,
EinsumOp::DeviceHelpers::CudaDeviceHelpers::DataCopy);
return einsum_compute_processor->Run();
}
return ORT_MAKE_STATUS(ONNXRUNTIME, NOT_IMPLEMENTED,
@ -81,4 +83,3 @@ Status Einsum::DeviceCompute(OpKernelContext* context, const std::vector<const T
} // namespace cuda
} // namespace onnxruntime
#endif

5
onnxruntime/core/providers/cuda/math/einsum.h
View file

@ -3,10 +3,9 @@
#pragma once
#include "core/providers/cpu/math/einsum.h"
#include "core/common/common.h"
#include "core/framework/op_kernel.h"
#include "core/platform/threadpool.h"
#include "core/providers/cuda/cuda_common.h"
#include "core/providers/cpu/math/einsum.h"
#include "einsum_utils/einsum_auxiliary_ops.h"
#include "core/providers/cuda/cuda_execution_provider.h"

4
onnxruntime/core/providers/cuda/math/einsum_utils/einsum_auxiliary_ops.cc
View file

@ -1,7 +1,6 @@
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
#if 0
#include "core/providers/shared_library/provider_api.h"
namespace onnxruntime {
@ -181,7 +180,7 @@ template Status DeviceHelpers::CudaDeviceHelpers::MatMul<MLFloat16>(
size_t num_batches, size_t M, size_t K, size_t N, concurrency::ThreadPool* tp,
void* einsum_cuda_assets);
template Tensor DeviceHelpers::CudaDeviceHelpers::ReduceSum<MLFloat16>(
template std::unique_ptr<Tensor> DeviceHelpers::CudaDeviceHelpers::ReduceSum<MLFloat16>(
const Tensor& input, const std::vector<int64_t>& reduce_axes,
bool keep_dims, AllocatorPtr allocator,
const TensorShape* input_shape_override,
@ -190,4 +189,3 @@ template Tensor DeviceHelpers::CudaDeviceHelpers::ReduceSum<MLFloat16>(
} // namespace EinsumOp
} // namespace onnxruntime
#endif

8
onnxruntime/core/providers/cuda/math/einsum_utils/einsum_auxiliary_ops.h
View file

@ -47,10 +47,10 @@ Status MatMul(const T* input_1_data, const T* input_2_data, T* output_data,
void* einsum_cuda_assets);
template <typename T>
Tensor ReduceSum(const Tensor& input, const std::vector<int64_t>& reduce_axes,
bool keep_dims, AllocatorPtr allocator,
const TensorShape* input_shape_override,
concurrency::ThreadPool* /*tp*/, void* einsum_cuda_assets);
std::unique_ptr<Tensor> ReduceSum(const Tensor& input, const std::vector<int64_t>& reduce_axes,
bool keep_dims, AllocatorPtr allocator,
const TensorShape* input_shape_override,
concurrency::ThreadPool* /*tp*/, void* einsum_cuda_assets);
std::unique_ptr<Tensor> Diagonal(const Tensor& input, int64_t dim_1, int64_t dim_2, AllocatorPtr allocator, void* einsum_cuda_assets);

10
onnxruntime/core/providers/cuda/math/gemm.cc
View file

@ -1,4 +1,3 @@
#if 0
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
@ -18,7 +17,7 @@ namespace cuda {
8, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder() \
(*KernelDefBuilder::Create()) \
.TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
Gemm<T>); \
ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \
@ -28,7 +27,7 @@ namespace cuda {
10, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder() \
(*KernelDefBuilder::Create()) \
.TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
Gemm<T>); \
ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \
@ -38,7 +37,7 @@ namespace cuda {
12, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder() \
(*KernelDefBuilder::Create()) \
.TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
Gemm<T>); \
ONNX_OPERATOR_TYPED_KERNEL_EX( \
@ -47,7 +46,7 @@ namespace cuda {
13, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder() \
(*KernelDefBuilder::Create()) \
.TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
Gemm<T>);
@ -147,4 +146,3 @@ Status Gemm<T>::ComputeInternal(OpKernelContext* ctx) const {
} // namespace cuda
} // namespace onnxruntime
#endif

8
onnxruntime/core/providers/cuda/math/matmul.cc
View file

@ -1,4 +1,3 @@
#if 0
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
@ -17,7 +16,7 @@ namespace cuda {
1, 8, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder() \
(*KernelDefBuilder::Create()) \
.TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
MatMul<T>); \
ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \
@ -26,7 +25,7 @@ namespace cuda {
9, 12, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder() \
(*KernelDefBuilder::Create()) \
.TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
MatMul<T>); \
ONNX_OPERATOR_TYPED_KERNEL_EX( \
@ -35,7 +34,7 @@ namespace cuda {
13, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder() \
(*KernelDefBuilder::Create()) \
.TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
MatMul<T>);
@ -197,4 +196,3 @@ Status MatMul<T>::ComputeInternal(OpKernelContext* ctx) const {
} // namespace cuda
} // namespace onnxruntime
#endif

8
onnxruntime/core/providers/cuda/math/matmul_integer.cc
View file

@ -1,4 +1,3 @@
#if 0
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
@ -19,9 +18,9 @@ ONNX_OPERATOR_TYPED_KERNEL_EX(
10,
int8_t,
kCudaExecutionProvider,
KernelDefBuilder()
.InputMemoryType<OrtMemTypeCPUInput>(2)
.InputMemoryType<OrtMemTypeCPUInput>(3)
(*KernelDefBuilder::Create())
.InputMemoryType(OrtMemTypeCPUInput, 2)
.InputMemoryType(OrtMemTypeCPUInput, 3)
.TypeConstraint("T1", DataTypeImpl::GetTensorType<int8_t>())
.TypeConstraint("T2", DataTypeImpl::GetTensorType<int8_t>())
.TypeConstraint("T3", DataTypeImpl::GetTensorType<int32_t>()),
@ -113,4 +112,3 @@ Status MatMulInteger<int8_t, int8_t>::ComputeInternal(OpKernelContext* ctx) cons
} // namespace cuda
} // namespace onnxruntime
#endif

159
onnxruntime/core/providers/cuda/math/softmax.cc
View file

@ -1,4 +1,3 @@
#if 0
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
@ -51,7 +50,7 @@ Status SoftMaxComputeHelper(
return Status::OK();
}
#define SPECIALIZED_SOFTMAX_HELPER_IMPL(T) \
#define SPECIALIZED_SOFTMAX_HELPER_IMPL(T) \
template Status SoftMaxComputeHelper<T, false>(cudaStream_t stream, const T* input, const TensorShape& shape, T* Y, cudnnHandle_t handle, int64_t axis); \
template Status SoftMaxComputeHelper<T, true>(cudaStream_t stream, const T* input, const TensorShape& shape, T* Y, cudnnHandle_t handle, int64_t axis);
@ -61,81 +60,81 @@ SPECIALIZED_SOFTMAX_HELPER_IMPL(MLFloat16)
#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
// cudnnSoftmaxForward/Backward doesn't support BFloat16.
#define SPECIALIZED_SOFTMAX_HELPER_IMPL_BFloat16(is_log_softmax) \
template <> \
Status SoftMaxComputeHelper<BFloat16, is_log_softmax>( \
cudaStream_t stream, \
const BFloat16* X, \
const TensorShape& input_shape, \
BFloat16* Y, \
cudnnHandle_t, \
int64_t axis) { \
typedef typename ToCudaType<BFloat16>::MappedType CudaT; \
int64_t N = input_shape.SizeToDimension(axis); \
int64_t D = input_shape.SizeFromDimension(axis); \
auto Y_data = reinterpret_cast<CudaT*>(Y); \
auto X_data = reinterpret_cast<const CudaT*>(X); \
dispatch_softmax_forward<CudaT, CudaT, AccumulationType_t<CudaT>, is_log_softmax>( \
stream, Y_data, X_data, gsl::narrow_cast<int>(D), gsl::narrow_cast<int>(D), gsl::narrow_cast<int>(N)); \
return Status::OK(); \
#define SPECIALIZED_SOFTMAX_HELPER_IMPL_BFloat16(is_log_softmax) \
template <> \
Status SoftMaxComputeHelper<BFloat16, is_log_softmax>( \
cudaStream_t stream, \
const BFloat16* X, \
const TensorShape& input_shape, \
BFloat16* Y, \
cudnnHandle_t, \
int64_t axis) { \
typedef typename ToCudaType<BFloat16>::MappedType CudaT; \
int64_t N = input_shape.SizeToDimension(axis); \
int64_t D = input_shape.SizeFromDimension(axis); \
auto Y_data = reinterpret_cast<CudaT*>(Y); \
auto X_data = reinterpret_cast<const CudaT*>(X); \
dispatch_softmax_forward<CudaT, CudaT, AccumulationType_t<CudaT>, is_log_softmax>( \
stream, Y_data, X_data, gsl::narrow_cast<int>(D), gsl::narrow_cast<int>(D), gsl::narrow_cast<int>(N)); \
return Status::OK(); \
}
SPECIALIZED_SOFTMAX_HELPER_IMPL_BFloat16(true)
SPECIALIZED_SOFTMAX_HELPER_IMPL_BFloat16(false)
SPECIALIZED_SOFTMAX_HELPER_IMPL_BFloat16(false)
#endif
#define REGISTER_KERNEL_TYPED(T) \
ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \
Softmax, \
kOnnxDomain, \
1, 10, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
Softmax<T>); \
ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \
Softmax, \
kOnnxDomain, \
11, 12, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
Softmax<T>); \
ONNX_OPERATOR_TYPED_KERNEL_EX( \
Softmax, \
kOnnxDomain, \
13, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
Softmax<T>); \
ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \
LogSoftmax, \
kOnnxDomain, \
1, 10, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
Softmax<T>); \
ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \
LogSoftmax, \
kOnnxDomain, \
11, 12, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
Softmax<T>); \
ONNX_OPERATOR_TYPED_KERNEL_EX( \
LogSoftmax, \
kOnnxDomain, \
13, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
#define REGISTER_KERNEL_TYPED(T) \
ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \
Softmax, \
kOnnxDomain, \
1, 10, \
T, \
kCudaExecutionProvider, \
(*KernelDefBuilder::Create()).TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
Softmax<T>); \
ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \
Softmax, \
kOnnxDomain, \
11, 12, \
T, \
kCudaExecutionProvider, \
(*KernelDefBuilder::Create()).TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
Softmax<T>); \
ONNX_OPERATOR_TYPED_KERNEL_EX( \
Softmax, \
kOnnxDomain, \
13, \
T, \
kCudaExecutionProvider, \
(*KernelDefBuilder::Create()).TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
Softmax<T>); \
ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \
LogSoftmax, \
kOnnxDomain, \
1, 10, \
T, \
kCudaExecutionProvider, \
(*KernelDefBuilder::Create()).TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
Softmax<T>); \
ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \
LogSoftmax, \
kOnnxDomain, \
11, 12, \
T, \
kCudaExecutionProvider, \
(*KernelDefBuilder::Create()).TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
Softmax<T>); \
ONNX_OPERATOR_TYPED_KERNEL_EX( \
LogSoftmax, \
kOnnxDomain, \
13, \
T, \
kCudaExecutionProvider, \
(*KernelDefBuilder::Create()).TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
Softmax<T>);
template <typename T>
Status Softmax<T>::ComputeInternal(OpKernelContext* ctx) const {
template <typename T>
Status Softmax<T>::ComputeInternal(OpKernelContext* ctx) const {
const Tensor* X = ctx->Input<Tensor>(0);
const TensorShape& input_shape{X->Shape()};
size_t rank = input_shape.NumDimensions();
@ -149,9 +148,9 @@ Status Softmax<T>::ComputeInternal(OpKernelContext* ctx) const {
const size_t axis = static_cast<size_t>(HandleNegativeAxis(axis_, rank));
bool is_transpose_required = false;
Tensor transposed_input;
std::unique_ptr<Tensor> transposed_input;
std::vector<int64_t> transposed_input_dims;
Tensor intermediate_output; // output that the softmax implementation will write into while using transposed input
std::unique_ptr<Tensor> intermediate_output; // output that the softmax implementation will write into while using transposed input
std::vector<size_t> permutation(rank);
// The "semantic" meaning of axis has changed in opset-13.
@ -182,18 +181,17 @@ Status Softmax<T>::ComputeInternal(OpKernelContext* ctx) const {
}
// Allocate a temporary tensor to hold transposed input
Tensor temp_input(X->DataType(), TensorShape(transposed_input_dims), alloc);
auto temp_input = Tensor::Create(X->DataType(), TensorShape(transposed_input_dims), alloc);
// Perform the transpose
ORT_RETURN_IF_ERROR(Transpose::DoTranspose(cuda_ep_->GetDeviceProp(),
Stream(),
CublasHandle(),
permutation, *X, temp_input));
permutation, *X, *temp_input));
transposed_input = std::move(temp_input);
// Allocate memory for the intermediate output
Tensor temp_output(Y->DataType(), TensorShape(transposed_input_dims), alloc);
intermediate_output = std::move(temp_output);
intermediate_output = Tensor::Create(Y->DataType(), TensorShape(transposed_input_dims), alloc);
}
const T* X_data = nullptr;
@ -201,9 +199,9 @@ Status Softmax<T>::ComputeInternal(OpKernelContext* ctx) const {
const TensorShape* compute_input_shape = nullptr;
if (is_transpose_required) { // use intermediate buffers to compute the softmax values
X_data = transposed_input.template Data<T>();
Y_data = intermediate_output.template MutableData<T>();
compute_input_shape = &transposed_input.Shape();
X_data = transposed_input->template Data<T>();
Y_data = intermediate_output->template MutableData<T>();
compute_input_shape = &transposed_input->Shape();
} else { // use the node input/output directly
X_data = X->template Data<T>();
Y_data = Y->template MutableData<T>();
@ -233,7 +231,7 @@ Status Softmax<T>::ComputeInternal(OpKernelContext* ctx) const {
ORT_RETURN_IF_ERROR(Transpose::DoTranspose(cuda_ep_->GetDeviceProp(),
Stream(),
CublasHandle(),
reverse_permutation, intermediate_output, *Y));
reverse_permutation, *intermediate_output, *Y));
}
return Status::OK();
@ -252,4 +250,3 @@ SPECIALIZED_COMPUTE(BFloat16)
} // namespace cuda
} // namespace onnxruntime
#endif

8
onnxruntime/core/providers/cuda/math/topk.cc
View file

@ -1,4 +1,3 @@
#if 0
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
@ -13,7 +12,7 @@ ONNX_OPERATOR_VERSIONED_KERNEL_EX(
kOnnxDomain,
1, 9,
kCudaExecutionProvider,
KernelDefBuilder().TypeConstraint("T", DataTypeImpl::AllFixedSizeTensorTypes()),
(*KernelDefBuilder::Create()).TypeConstraint("T", DataTypeImpl::AllFixedSizeTensorTypes()),
TopK<false>);
ONNX_OPERATOR_VERSIONED_KERNEL_EX(
@ -21,7 +20,7 @@ ONNX_OPERATOR_VERSIONED_KERNEL_EX(
kOnnxDomain,
10, 10,
kCudaExecutionProvider,
KernelDefBuilder().InputMemoryType<OrtMemTypeCPUInput>(1).TypeConstraint("T", DataTypeImpl::AllFixedSizeTensorTypes()).TypeConstraint("I", DataTypeImpl::GetTensorType<int64_t>()),
(*KernelDefBuilder::Create()).InputMemoryType(OrtMemTypeCPUInput, 1).TypeConstraint("T", DataTypeImpl::AllFixedSizeTensorTypes()).TypeConstraint("I", DataTypeImpl::GetTensorType<int64_t>()),
TopK<true>);
ONNX_OPERATOR_KERNEL_EX(
@ -29,7 +28,7 @@ ONNX_OPERATOR_KERNEL_EX(
kOnnxDomain,
11,
kCudaExecutionProvider,
KernelDefBuilder().InputMemoryType<OrtMemTypeCPUInput>(1).TypeConstraint("T", DataTypeImpl::AllFixedSizeTensorTypes()).TypeConstraint("I", DataTypeImpl::GetTensorType<int64_t>()),
(*KernelDefBuilder::Create()).InputMemoryType(OrtMemTypeCPUInput, 1).TypeConstraint("T", DataTypeImpl::AllFixedSizeTensorTypes()).TypeConstraint("I", DataTypeImpl::GetTensorType<int64_t>()),
TopK<true>);
template <bool inputk>
@ -104,4 +103,3 @@ Status TopK<inputk>::ComputeInternal(OpKernelContext* ctx) const {
} // namespace cuda
} // namespace onnxruntime
#endif

1
onnxruntime/core/providers/cuda/math/topk_impl.cu
View file

@ -2,7 +2,6 @@
// Licensed under the MIT License.
#include "topk_impl.h"
//#include "core/framework/data_types.h"
#include "core/providers/cuda/cu_inc/common.cuh"
#include "device_atomic_functions.h"
#include "cub/cub.cuh"

18
onnxruntime/core/providers/cuda/math/variadic_elementwise_ops.cc
View file

@ -1,4 +1,3 @@
#if 0
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
@ -214,13 +213,13 @@ const auto k_hfd_datatypes =
// kernel registration
#define REGISTER_KERNEL(name, impl_class, version, datatypes) \
ONNX_OPERATOR_KERNEL_EX( \
name, \
kOnnxDomain, \
version, \
kCudaExecutionProvider, \
KernelDefBuilder().TypeConstraint("T", datatypes), \
#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_VERSIONED_KERNEL(name, impl_class, start_version, end_version, datatypes) \
@ -229,7 +228,7 @@ const auto k_hfd_datatypes =
kOnnxDomain, \
start_version, end_version, \
kCudaExecutionProvider, \
KernelDefBuilder().TypeConstraint("T", datatypes), \
(*KernelDefBuilder::Create()).TypeConstraint("T", datatypes), \
impl_class)
REGISTER_KERNEL(Sum, SumOp, 13, k_hfd_datatypes)
@ -249,4 +248,3 @@ REGISTER_VERSIONED_KERNEL(Max, MaxOp, 6, 11, k_hfd_datatypes)
} // namespace cuda
} // namespace onnxruntime
#endif

6
onnxruntime/core/providers/cuda/nn/batch_norm.cc
View file

@ -1,4 +1,3 @@
#if 0
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
@ -19,7 +18,7 @@ namespace cuda {
7, 8, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder() \
(*KernelDefBuilder::Create()) \
.TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
BatchNorm<T>); \
ONNX_OPERATOR_TYPED_KERNEL_EX( \
@ -28,7 +27,7 @@ namespace cuda {
9, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder() \
(*KernelDefBuilder::Create()) \
.TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
BatchNorm<T>);
@ -165,4 +164,3 @@ SPECIALIZED_COMPUTE(MLFloat16)
} // namespace cuda
} // namespace onnxruntime
#endif

1
onnxruntime/core/providers/cuda/nn/batch_norm.h
View file

@ -3,7 +3,6 @@
#pragma once
#include "gsl/gsl"
#include "core/providers/cuda/cuda_kernel.h"
#include "core/providers/cuda/cudnn_common.h"

38
onnxruntime/core/providers/cuda/nn/conv.cc
View file

@ -1,8 +1,6 @@
#if 0
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
#include "core/providers/common.h"
#include "core/providers/cuda/cuda_common.h"
#include "core/providers/cuda/nn/conv.h"
#include "core/providers/cuda/shared_inc/fpgeneric.h"
@ -13,22 +11,22 @@ namespace cuda {
// Op Set 11 for Conv only update document to clearify default dilations and strides value.
// which are already convered by op set 11 cpu versoin, so simply add declaration.
#define REGISTER_KERNEL_TYPED(T) \
ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \
Conv, \
kOnnxDomain, \
1, 10, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
Conv<T>); \
ONNX_OPERATOR_TYPED_KERNEL_EX( \
Conv, \
kOnnxDomain, \
11, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
#define REGISTER_KERNEL_TYPED(T) \
ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \
Conv, \
kOnnxDomain, \
1, 10, \
T, \
kCudaExecutionProvider, \
(*KernelDefBuilder::Create()).TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
Conv<T>); \
ONNX_OPERATOR_TYPED_KERNEL_EX( \
Conv, \
kOnnxDomain, \
11, \
T, \
kCudaExecutionProvider, \
(*KernelDefBuilder::Create()).TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
Conv<T>);
REGISTER_KERNEL_TYPED(float)
@ -36,8 +34,7 @@ REGISTER_KERNEL_TYPED(double)
REGISTER_KERNEL_TYPED(MLFloat16)
Status SliceOutUnwantedOutputSection(cudaStream_t stream,
const void* input_data,
const std::vector<int64_t>& input_dims,
const void* input_data, const std::vector<int64_t>& input_dims,
void* output_data,
const std::vector<int64_t>& output_dims,
std::vector<int64_t> starts,
@ -359,4 +356,3 @@ Status CudnnConvolutionDescriptor::Set(
} // namespace cuda
} // namespace onnxruntime
#endif

1
onnxruntime/core/providers/cuda/nn/conv.h
View file

@ -3,7 +3,6 @@
#pragma once
//#include "core/common/common.h"
#include "core/platform/ort_mutex.h"
#include "core/providers/cuda/cuda_kernel.h"
#include "core/providers/cuda/cudnn_common.h"

34
onnxruntime/core/providers/cuda/nn/conv_transpose.cc
View file

@ -1,4 +1,3 @@
#if 0
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
@ -9,22 +8,22 @@ namespace cuda {
// Op Set 11 for ConvTranspose only update document to clearify default dilations and strides value.
// which are already covered by op set 11 cpu version, so simply add declaration.
#define REGISTER_KERNEL_TYPED(T) \
ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \
ConvTranspose, \
kOnnxDomain, \
1, 10, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
ConvTranspose<T>); \
ONNX_OPERATOR_TYPED_KERNEL_EX( \
ConvTranspose, \
kOnnxDomain, \
11, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
#define REGISTER_KERNEL_TYPED(T) \
ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \
ConvTranspose, \
kOnnxDomain, \
1, 10, \
T, \
kCudaExecutionProvider, \
(*KernelDefBuilder::Create()).TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
ConvTranspose<T>); \
ONNX_OPERATOR_TYPED_KERNEL_EX( \
ConvTranspose, \
kOnnxDomain, \
11, \
T, \
kCudaExecutionProvider, \
(*KernelDefBuilder::Create()).TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
ConvTranspose<T>);
REGISTER_KERNEL_TYPED(float)
@ -208,4 +207,3 @@ Status ConvTranspose<T>::DoConvTranspose(OpKernelContext* context, bool dynamic_
} // namespace cuda
} // namespace onnxruntime
#endif

2
onnxruntime/core/providers/cuda/nn/conv_transpose.h
View file

@ -7,7 +7,7 @@
#include "core/providers/cuda/cuda_kernel.h"
#include "core/providers/cuda/cudnn_common.h"
#include "core/providers/cuda/nn/conv.h"
//#include "core/providers/cpu/nn/conv_transpose_attributes.h"
#include "core/providers/cpu/nn/conv_transpose_attributes.h"
namespace onnxruntime {
namespace cuda {

23
onnxruntime/core/providers/cuda/nn/dropout.cc
View file

@ -1,4 +1,3 @@
#if 0
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
@ -8,10 +7,11 @@ namespace onnxruntime {
namespace cuda {
#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
#define ALL_IEEE_FLOAT_TENSOR_TYPES {DataTypeImpl::GetTensorType<float>(), \
DataTypeImpl::GetTensorType<double>(), \
DataTypeImpl::GetTensorType<MLFloat16>(), \
DataTypeImpl::GetTensorType<BFloat16>()}
#define ALL_IEEE_FLOAT_TENSOR_TYPES \
{ DataTypeImpl::GetTensorType<float>(), \
DataTypeImpl::GetTensorType<double>(), \
DataTypeImpl::GetTensorType<MLFloat16>(), \
DataTypeImpl::GetTensorType<BFloat16>() }
#else
#define ALL_IEEE_FLOAT_TENSOR_TYPES DataTypeImpl::AllIEEEFloatTensorTypes()
#endif
@ -21,12 +21,12 @@ ONNX_OPERATOR_VERSIONED_KERNEL_EX(
kOnnxDomain,
12, 12,
kCudaExecutionProvider,
KernelDefBuilder()
(*KernelDefBuilder::Create())
.TypeConstraint("T", DataTypeImpl::AllIEEEFloatTensorTypes())
.TypeConstraint("T1", DataTypeImpl::AllIEEEFloatTensorTypes())
.TypeConstraint("T2", DataTypeImpl::GetTensorType<bool>())
.InputMemoryType<OrtMemTypeCPUInput>(1)
.InputMemoryType<OrtMemTypeCPUInput>(2),
.InputMemoryType(OrtMemTypeCPUInput, 1)
.InputMemoryType(OrtMemTypeCPUInput, 2),
Dropout);
ONNX_OPERATOR_KERNEL_EX(
@ -34,14 +34,13 @@ ONNX_OPERATOR_KERNEL_EX(
kOnnxDomain,
13,
kCudaExecutionProvider,
KernelDefBuilder()
(*KernelDefBuilder::Create())
.TypeConstraint("T", ALL_IEEE_FLOAT_TENSOR_TYPES)
.TypeConstraint("T1", ALL_IEEE_FLOAT_TENSOR_TYPES)
.TypeConstraint("T2", DataTypeImpl::GetTensorType<bool>())
.InputMemoryType<OrtMemTypeCPUInput>(1)
.InputMemoryType<OrtMemTypeCPUInput>(2),
.InputMemoryType(OrtMemTypeCPUInput, 1)
.InputMemoryType(OrtMemTypeCPUInput, 2),
Dropout);
} // namespace cuda
} // namespace onnxruntime
#endif

4
onnxruntime/core/providers/cuda/nn/instance_norm.cc
View file

@ -1,4 +1,3 @@
#if 0
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
@ -17,7 +16,7 @@ namespace cuda {
6, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder() \
(*KernelDefBuilder::Create()) \
.TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
InstanceNorm<T>);
@ -154,4 +153,3 @@ Status InstanceNorm<T>::ComputeInternal(OpKernelContext* p_op_kernel_context) co
} // namespace cuda
} // namespace onnxruntime
#endif

36
onnxruntime/core/providers/cuda/nn/lrn.cc
View file

@ -1,4 +1,3 @@
#if 0
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
@ -7,25 +6,25 @@
namespace onnxruntime {
namespace cuda {
#define REGISTER_KERNEL_VERSIONED_TYPED(START_VER, END_VER, T) \
ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \
LRN, \
kOnnxDomain, \
START_VER, \
END_VER, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
#define REGISTER_KERNEL_VERSIONED_TYPED(START_VER, END_VER, T) \
ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \
LRN, \
kOnnxDomain, \
START_VER, \
END_VER, \
T, \
kCudaExecutionProvider, \
(*KernelDefBuilder::Create()).TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
LRN<T>);
#define REGISTER_KERNEL_TYPED(VER, T) \
ONNX_OPERATOR_TYPED_KERNEL_EX( \
LRN, \
kOnnxDomain, \
VER, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
#define REGISTER_KERNEL_TYPED(VER, T) \
ONNX_OPERATOR_TYPED_KERNEL_EX( \
LRN, \
kOnnxDomain, \
VER, \
T, \
kCudaExecutionProvider, \
(*KernelDefBuilder::Create()).TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
LRN<T>);
REGISTER_KERNEL_VERSIONED_TYPED(1, 12, float)
@ -111,4 +110,3 @@ Status CudnnLRNDescriptor::Set(uint32_t N, double alpha, double beta, double K)
} // namespace cuda
} // namespace onnxruntime
#endif

44
onnxruntime/core/providers/cuda/nn/pool.cc
View file

@ -1,9 +1,8 @@
#if 0
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
#include "core/providers/shared_library/provider_api.h"
#include "core/providers/cuda/nn/pool.h"
#include "core/providers/common.h"
#include "core/providers/cuda/cudnn_common.h"
#include "core/providers/cuda/nn/max_pool_with_index.h"
#include "core/providers/cuda/math/unary_elementwise_ops_impl.h"
@ -12,28 +11,27 @@ using namespace onnxruntime::common;
namespace onnxruntime {
namespace cuda {
#define POOLING_KERNEL(op_name, data_type, pool_type, since_version) \
ONNX_OPERATOR_TYPED_KERNEL_EX( \
op_name, \
kOnnxDomain, \
since_version, \
data_type, \
kCudaExecutionProvider, \
KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType<data_type>()), \
#define POOLING_KERNEL(op_name, data_type, pool_type, since_version) \
ONNX_OPERATOR_TYPED_KERNEL_EX( \
op_name, \
kOnnxDomain, \
since_version, \
data_type, \
kCudaExecutionProvider, \
(*KernelDefBuilder::Create()).TypeConstraint("T", DataTypeImpl::GetTensorType<data_type>()), \
Pool<data_type, pool_type>);
#define POOLING_KERNEL_VERSIONED(op_name, data_type, pool_type, since_version, end_version) \
ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \
op_name, \
kOnnxDomain, \
since_version, \
end_version, \
data_type, \
kCudaExecutionProvider, \
KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType<data_type>()).TypeConstraint("I", DataTypeImpl::GetTensorType<int64_t>()), \
#define POOLING_KERNEL_VERSIONED(op_name, data_type, pool_type, since_version, end_version) \
ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \
op_name, \
kOnnxDomain, \
since_version, \
end_version, \
data_type, \
kCudaExecutionProvider, \
(*KernelDefBuilder::Create()).TypeConstraint("T", DataTypeImpl::GetTensorType<data_type>()).TypeConstraint("I", DataTypeImpl::GetTensorType<int64_t>()), \
Pool<data_type, pool_type>);
POOLING_KERNEL_VERSIONED(AveragePool, float, AveragePool, 7, 9)
POOLING_KERNEL_VERSIONED(AveragePool, double, AveragePool, 7, 9)
POOLING_KERNEL_VERSIONED(AveragePool, MLFloat16, AveragePool, 7, 9)
@ -65,7 +63,6 @@ POOLING_KERNEL(MaxPool, MLFloat16, MaxPool<8>, 12)
POOLING_KERNEL(MaxPool, int8_t, MaxPool<8>, 12)
POOLING_KERNEL(MaxPool, uint8_t, MaxPool<8>, 12)
POOLING_KERNEL(GlobalMaxPool, float, MaxPool<1>, 1)
POOLING_KERNEL(GlobalMaxPool, double, MaxPool<1>, 1)
POOLING_KERNEL(GlobalMaxPool, MLFloat16, MaxPool<1>, 1)
@ -168,8 +165,8 @@ Status Pool<T, PoolType>::ComputeInternal(OpKernelContext* context) const {
cudnnPoolingMode_t mode = CUDNN_POOLING_MAX;
if (PoolType::type == onnxruntime::PoolType::kAveragePool) {
mode = pool_attrs_.count_include_pad ? CUDNN_POOLING_AVERAGE_COUNT_INCLUDE_PADDING
: CUDNN_POOLING_AVERAGE_COUNT_EXCLUDE_PADDING;
mode = pool_attrs_.count_include_pad ? CUDNN_POOLING_AVERAGE_COUNT_INCLUDE_PADDING
: CUDNN_POOLING_AVERAGE_COUNT_EXCLUDE_PADDING;
}
CudnnPoolingDescriptor pooling_desc;
ORT_RETURN_IF_ERROR(pooling_desc.Set(mode, kernel_shape, pads, strides));
@ -259,4 +256,3 @@ Status Pool<T, MaxPool<8>>::ComputeInternal(OpKernelContext* context) const {
} // namespace cuda
} // namespace onnxruntime
#endif

5
onnxruntime/core/providers/cuda/nn/pool.h
View file

@ -3,7 +3,6 @@
#pragma once
#include "core/common/common.h"
#include "core/providers/cuda/cuda_kernel.h"
#include "core/providers/cuda/cudnn_common.h"
#include "core/providers/cpu/nn/pool_base.h"
@ -14,7 +13,7 @@ namespace cuda {
template <typename T, typename PoolType>
class Pool : public CudaKernel, public PoolBase {
public:
Pool(OpKernelInfo info) : CudaKernel(info), PoolBase(info) {}
Pool(const OpKernelInfo& info) : CudaKernel(info), PoolBase(info) {}
Status ComputeInternal(OpKernelContext* context) const override;
};
@ -22,7 +21,7 @@ class Pool : public CudaKernel, public PoolBase {
template <typename T>
class Pool<T, MaxPool<8>> final : public Pool<T, MaxPool<1>> {
public:
Pool(OpKernelInfo info) : Pool<T, MaxPool<1>>(info) {}
Pool(const OpKernelInfo& info) : Pool<T, MaxPool<1>>(info) {}
Status ComputeInternal(OpKernelContext* context) const override;
};

4
onnxruntime/core/providers/cuda/nn/shrink.cc
View file

@ -1,4 +1,3 @@
#if 0
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
@ -17,7 +16,7 @@ namespace cuda {
9, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder() \
(*KernelDefBuilder::Create()) \
.MayInplace(0, 0) \
.TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
Shrink<T>);
@ -53,4 +52,3 @@ SHRINK_REGISTER_KERNEL(int64_t)
} // namespace cuda
} // namespace onnxruntime
#endif

Some files were not shown because too many files have changed in this diff Show more