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

[ROCm] do not use failed miopen fusion compile (#19012)

Browse source 
The FusedConv operator for the ROCm EP could fail to compile the fused
operation, in which case it should not attempt to use the failed fusion
plan. In addition, the hash for the miopenConvolutionDescriptor_t for
newer ROCm versions was failing to use all components of the descriptor.
This commit is contained in:
Jeff Daily 2024-01-08 03:06:45 -08:00 committed by GitHub
parent 4190c29d22
commit db3c076081
No known key found for this signature in database
GPG key ID: 4AEE18F83AFDEB23
1 changed files with 10 additions and 2 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

12
onnxruntime/contrib_ops/rocm/fused_conv.cc
View file

@ -76,7 +76,12 @@ struct FNVHash {
void HashConvolutionDescriptor(miopenConvolutionDescriptor_t cdesc) {
int spatial_dim = 1;
#if ROCM_VERSION >= 50500
miopenGetConvolutionSpatialDim(cdesc, &spatial_dim);
MIOPEN_CALL(miopenGetConvolutionSpatialDim(cdesc, &spatial_dim));
std::vector<int> pads{spatial_dim};
std::vector<int> strides{spatial_dim};
std::vector<int> dilations{spatial_dim};
miopenConvolutionMode_t mode;
MIOPEN_CALL(miopenGetConvolutionNdDescriptor(cdesc, spatial_dim, &spatial_dim, pads.data(), strides.data(), dilations.data(), &mode));
#else
// Previous versions of MIOpen doesn't provide API to probe the dimension of a
// miopenConvolutionDescriptor_t, so we have to guess.
@ -100,11 +105,12 @@ struct FNVHash {
pads.resize(spatial_dim);
strides.resize(spatial_dim);
dilations.resize(spatial_dim);
#endif
(*this) << spatial_dim;
(*this) << pads;
(*this) << strides;
(*this) << dilations;
#endif
(*this) << mode;
}
private:
@ -313,6 +319,8 @@ class FusedConv : public onnxruntime::rocm::Conv<T, false> {
auto ret = miopenCompileFusionPlan(handle, fusion->plan);
if (miopenStatusSuccess == ret) {
fusion->compiled_on.insert(handle);
} else {
return ret;
}
return miopenStatusSuccess;
}