Skip to content

Commit 03b34e9

Browse files
turneramcausten
authored andcommitted
Failing fusion plan workaround (#995)
* Add workaround for devices that do not support miopen conv fusions
1 parent c4b85b2 commit 03b34e9

File tree

1 file changed

+11
-0
lines changed

1 file changed

+11
-0
lines changed

src/targets/gpu/fuse_ops.cpp

100755100644
Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -5,6 +5,7 @@
55
#include <migraphx/gpu/miopen.hpp>
66
#include <migraphx/gpu/clip.hpp>
77
#include <migraphx/gpu/convolution.hpp>
8+
#include <migraphx/gpu/device_name.hpp>
89
#include <migraphx/gpu/oper.hpp>
910
#include <migraphx/gpu/add.hpp>
1011
#include <migraphx/gpu/mul.hpp>
@@ -26,6 +27,7 @@
2627
#include <migraphx/array.hpp>
2728
#include <migraphx/op/clip.hpp>
2829
#include <cmath>
30+
#include <set>
2931

3032
namespace migraphx {
3133
inline namespace MIGRAPHX_INLINE_NS {
@@ -152,6 +154,12 @@ struct fusion
152154
}
153155
};
154156

157+
const std::unordered_set<std::string>& get_supported_archs()
158+
{
159+
static std::unordered_set<std::string> supported_archs{"gfx900", "gfx906", "gfx908", "gfx1030"};
160+
return supported_archs;
161+
}
162+
155163
MIGRAPHX_PRED_MATCHER(bias_shape, instruction_ref ins)
156164
{
157165
auto&& s = ins->get_shape();
@@ -161,6 +169,9 @@ MIGRAPHX_PRED_MATCHER(bias_shape, instruction_ref ins)
161169

162170
MIGRAPHX_PRED_MATCHER(fusable_conv, instruction_ref ins)
163171
{
172+
const auto device_name = split_string(get_device_name(), ':').front();
173+
if(not contains(get_supported_archs(), device_name))
174+
return false;
164175
if(enabled(MIGRAPHX_DISABLE_MIOPEN_FUSION{}))
165176
return false;
166177
if(ins->name() != "gpu::convolution")

0 commit comments

Comments
 (0)