Skip to content

Commit

Permalink
fix cuda/rocm provider info hash
Browse files Browse the repository at this point in the history
  • Loading branch information
@tianleiwu
tianleiwu committed Feb 4, 2024
1 parent 6d34a93 commit c003730
Show file tree
Hide file tree
Showing 4 changed files with 67 additions and 30 deletions.
39 changes: 32 additions & 7 deletions onnxruntime/core/providers/cuda/cuda_execution_provider_info.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -86,12 +86,37 @@ struct CUDAExecutionProviderInfo {
} // namespace onnxruntime

template <>
struct std::hash<::onnxruntime::cuda::TunableOpInfo> {
size_t operator()(const ::onnxruntime::cuda::TunableOpInfo& info) const {
size_t seed_and_value{0xbc9f1d34};
onnxruntime::HashCombine(info.enable, seed_and_value);
onnxruntime::HashCombine(info.tuning_enable, seed_and_value);
onnxruntime::HashCombine(info.max_tuning_duration_ms, seed_and_value);
return seed_and_value;
struct std::hash<::onnxruntime::CUDAExecutionProviderInfo> {
size_t operator()(const ::onnxruntime::CUDAExecutionProviderInfo& info) const {
size_t value{0xbc9f1d34}; // seed

// Bits: device_id (16), arena_extend_strategy/cudnn_conv_algo_search (reserved 2), boolean options (1 each)
size_t data = static_cast<size_t>(info.device_id) ^
(static_cast<size_t>(info.arena_extend_strategy) << 16) ^
(static_cast<size_t>(info.cudnn_conv_algo_search) << 18) ^
(static_cast<size_t>(info.do_copy_in_default_stream) << 20) ^
(static_cast<size_t>(info.has_user_compute_stream) << 21) ^
(static_cast<size_t>(info.cudnn_conv_use_max_workspace) << 22) ^
(static_cast<size_t>(info.enable_cuda_graph) << 23) ^
(static_cast<size_t>(info.tunable_op.enable) << 24) ^
(static_cast<size_t>(info.tunable_op.tuning_enable) << 25) ^
(static_cast<size_t>(info.cudnn_conv1d_pad_to_nc1d) << 26) ^
(static_cast<size_t>(info.enable_skip_layer_norm_strict_mode) << 27) ^
(static_cast<size_t>(info.prefer_nhwc) << 28) ^
(static_cast<size_t>(info.use_ep_level_unified_stream) << 29) ^
(static_cast<size_t>(info.use_tf32) << 30);
onnxruntime::HashCombine(data, value);

onnxruntime::HashCombine(info.gpu_mem_limit, value);
onnxruntime::HashCombine(info.tunable_op.max_tuning_duration_ms, value);

// Memory pointers
onnxruntime::HashCombine(reinterpret_cast<size_t>(info.user_compute_stream), value);
onnxruntime::HashCombine(reinterpret_cast<size_t>(info.external_allocator_info.alloc), value);
onnxruntime::HashCombine(reinterpret_cast<size_t>(info.external_allocator_info.free), value);
onnxruntime::HashCombine(reinterpret_cast<size_t>(info.external_allocator_info.empty_cache), value);

// The default memory arena cfg is not used in hashing right now.
return value;
}
};
34 changes: 27 additions & 7 deletions onnxruntime/core/providers/rocm/rocm_execution_provider_info.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -74,12 +74,32 @@ struct ROCMExecutionProviderInfo {
} // namespace onnxruntime

template <>
struct std::hash<::onnxruntime::rocm::TunableOpInfo> {
size_t operator()(const ::onnxruntime::rocm::TunableOpInfo& info) const {
size_t seed_and_value{0xbc9f1d34};
onnxruntime::HashCombine(info.enable, seed_and_value);
onnxruntime::HashCombine(info.tuning_enable, seed_and_value);
onnxruntime::HashCombine(info.max_tuning_duration_ms, seed_and_value);
return seed_and_value;
struct std::hash<::onnxruntime::ROCMExecutionProviderInfo> {
size_t operator()(const ::onnxruntime::ROCMExecutionProviderInfo& info) const {
size_t value{0xbc9f1d34}; // seed

// Bits: device_id (16), arena_extend_strategy/miopen_conv_exhaustive_search (reserved 2), boolean options (1 each)
size_t data = static_cast<size_t>(info.device_id) ^
(static_cast<size_t>(info.arena_extend_strategy) << 16) ^
(static_cast<size_t>(info.miopen_conv_exhaustive_search) << 18) ^
(static_cast<size_t>(info.do_copy_in_default_stream) << 20) ^
(static_cast<size_t>(info.has_user_compute_stream) << 21) ^
(static_cast<size_t>(info.miopen_conv_use_max_workspace) << 22) ^
(static_cast<size_t>(info.enable_hip_graph) << 23) ^
(static_cast<size_t>(info.tunable_op.enable) << 24) ^
(static_cast<size_t>(info.tunable_op.tuning_enable) << 25);
onnxruntime::HashCombine(data, value);

onnxruntime::HashCombine(info.gpu_mem_limit, value);
onnxruntime::HashCombine(info.tunable_op.max_tuning_duration_ms, value);

// Memory pointers
onnxruntime::HashCombine(reinterpret_cast<size_t>(info.user_compute_stream), value);
onnxruntime::HashCombine(reinterpret_cast<size_t>(info.external_allocator_info.alloc), value);
onnxruntime::HashCombine(reinterpret_cast<size_t>(info.external_allocator_info.free), value);
onnxruntime::HashCombine(reinterpret_cast<size_t>(info.external_allocator_info.empty_cache), value);

// The default memory arena cfg is not used in hashing right now.
return value;
}
};
4 changes: 4 additions & 0 deletions onnxruntime/test/python/onnxruntime_test_python.py
View file
Original file line number Diff line number Diff line change
Expand Up @@ -414,6 +414,8 @@ def test_get_and_set_option_with_values(option_name, option_values):
str(option_value),
)

test_get_and_set_option_with_values("enable_cuda_graph", ["1", "0"])

test_get_and_set_option_with_values("arena_extend_strategy", ["kNextPowerOfTwo", "kSameAsRequested"])

test_get_and_set_option_with_values("cudnn_conv_algo_search", ["DEFAULT", "EXHAUSTIVE", "HEURISTIC"])
Expand Down Expand Up @@ -555,6 +557,8 @@ def test_get_and_set_option_with_values(option_name, option_values):

test_get_and_set_option_with_values("tunable_op_max_tuning_duration_ms", ["-1", "1"])

test_get_and_set_option_with_values("enable_hip_graph", ["1", "0"])

run_rocm_options_test()

def test_invalid_set_providers(self):
Expand Down
20 changes: 4 additions & 16 deletions orttraining/orttraining/python/orttraining_python_module.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -47,7 +47,7 @@ void addObjectMethodsForLazyTensor(py::module& m);
#endif
bool InitArray();

bool GetDyanmicExecutionProviderHash(
bool GetDynamicExecutionProviderHash(
const std::string& ep_shared_lib_path,
const ProviderOptions& provider_options,
size_t& hash,
Expand Down Expand Up @@ -87,13 +87,7 @@ bool GetProviderInstanceHash(const std::string& type,
if (auto* cuda_provider_info = TryGetProviderInfo_CUDA()) {
const CUDAExecutionProviderInfo info = GetCudaExecutionProviderInfo(cuda_provider_info,
provider_options_map);
hash = static_cast<size_t>(info.device_id) ^
info.gpu_mem_limit ^
(static_cast<size_t>(info.arena_extend_strategy) << 16) ^
(static_cast<size_t>(info.cudnn_conv_algo_search) << 18) ^
(static_cast<size_t>(info.do_copy_in_default_stream) << 20) ^
(static_cast<size_t>(info.has_user_compute_stream) << 22) ^
std::hash<cuda::TunableOpInfo>{}(info.tunable_op);
hash = std::hash<CUDAExecutionProviderInfo>{}(info);
return true;
}
#endif
Expand All @@ -102,13 +96,7 @@ bool GetProviderInstanceHash(const std::string& type,
if (auto* rocm_provider_info = TryGetProviderInfo_ROCM()) {
const ROCMExecutionProviderInfo info = GetRocmExecutionProviderInfo(rocm_provider_info,
provider_options_map);
hash = static_cast<size_t>(info.device_id) ^
info.gpu_mem_limit ^
(static_cast<size_t>(info.arena_extend_strategy) << 16) ^
(static_cast<size_t>(info.miopen_conv_exhaustive_search) << 18) ^
(static_cast<size_t>(info.do_copy_in_default_stream) << 20) ^
(static_cast<size_t>(info.has_user_compute_stream) << 22) ^
std::hash<rocm::TunableOpInfo>{}(info.tunable_op);
hash = std::hash<ROCMExecutionProviderInfo>{}(info);
return true;
}
#endif
Expand All @@ -128,7 +116,7 @@ bool GetProviderInstanceHash(const std::string& type,
provider_options.insert(option);
}
}
return GetDyanmicExecutionProviderHash(shared_lib_path_it->second, provider_options, hash);
return GetDynamicExecutionProviderHash(shared_lib_path_it->second, provider_options, hash);
}
}
}
Expand Down

0 comments on commit c003730

Please sign in to comment.