Skip to content
GitLab
Menu
Projects
Groups
Snippets
Loading...
Help
Help
Support
Community forum
Keyboard shortcuts
?
Submit feedback
Contribute to GitLab
Sign in / Register
Toggle navigation
Menu
Open sidebar
gaoqiong
MIGraphX
Commits
cfbd91a9
Commit
cfbd91a9
authored
Jun 28, 2022
by
Paul
Browse files
Format
parent
c37b1636
Changes
2
Hide whitespace changes
Inline
Side-by-side
Showing
2 changed files
with
31 additions
and
37 deletions
+31
-37
src/targets/gpu/mlir.cpp
src/targets/gpu/mlir.cpp
+2
-5
src/targets/gpu/perfdb.cpp
src/targets/gpu/perfdb.cpp
+29
-32
No files found.
src/targets/gpu/mlir.cpp
View file @
cfbd91a9
...
...
@@ -465,7 +465,7 @@ struct mlir_program
ops
.
add_attribute_value
(
get_operator_value
(
ins
->
get_operator
()));
if
(
ins
->
name
()
!=
"@return"
)
ops
.
add_results
({
get_shape
(
ins
)});
if
(
ins
->
name
())
if
(
ins
->
name
())
pp
=
{
ins
->
get_operator
(),
ins
->
inputs
(),
ins
->
get_shape
()};
std
::
vector
<
MlirValue
>
inputs
;
...
...
@@ -525,10 +525,7 @@ struct mlir_program
MIGRAPHX_THROW
(
"Failed to compile mlir program"
);
}
std
::
string
get_tune_params
()
{
return
get_mlir_perf_for_conv
(
pp
);
}
std
::
string
get_tune_params
()
{
return
get_mlir_perf_for_conv
(
pp
);
}
mlir_context
ctx
;
MlirLocation
location
;
...
...
src/targets/gpu/perfdb.cpp
View file @
cfbd91a9
...
...
@@ -4,7 +4,6 @@
#include <migraphx/stringutils.hpp>
#include <migraphx/permutation.hpp>
namespace
migraphx
{
inline
namespace
MIGRAPHX_INLINE_NS
{
namespace
gpu
{
...
...
@@ -39,51 +38,50 @@ std::string get_type(const shape& s)
{
shape
::
int32_type
,
"INT32"
},
};
auto
it
=
m
.
find
(
s
.
type
());
if
(
it
==
m
.
end
())
if
(
it
==
m
.
end
())
return
"UNKNOWN"
;
return
it
->
second
;
}
std
::
string
generate_miopen_config
(
const
problem_params
&
pp
)
{
value
v
=
pp
.
op
.
to_value
();
auto
input
=
pp
.
inputs
[
0
].
lens
();
auto
weights
=
pp
.
inputs
[
1
].
lens
();
auto
output
=
pp
.
output
.
lens
();
auto
padding
=
v
[
"padding"
].
to_vector
<
std
::
size_t
>
();
auto
stride
=
v
[
"stride"
].
to_vector
<
std
::
size_t
>
();
value
v
=
pp
.
op
.
to_value
();
auto
input
=
pp
.
inputs
[
0
].
lens
();
auto
weights
=
pp
.
inputs
[
1
].
lens
();
auto
output
=
pp
.
output
.
lens
();
auto
padding
=
v
[
"padding"
].
to_vector
<
std
::
size_t
>
();
auto
stride
=
v
[
"stride"
].
to_vector
<
std
::
size_t
>
();
auto
dilation
=
v
[
"dilation"
].
to_vector
<
std
::
size_t
>
();
if
(
padding
.
size
()
!=
stride
.
size
())
padding
.
erase
(
padding
.
begin
()
+
padding
.
size
()
/
2
,
padding
.
end
());
if
(
padding
.
size
()
!=
stride
.
size
())
padding
.
erase
(
padding
.
begin
()
+
padding
.
size
()
/
2
,
padding
.
end
());
return
to_string_range
({
to_string
(
input
[
1
]),
to_string_range
(
input
.
begin
()
+
2
,
input
.
end
(),
"x"
),
to_string_range
(
weights
.
begin
()
+
2
,
weights
.
end
(),
"x"
),
to_string
(
weights
[
0
]),
to_string_range
(
output
.
begin
()
+
2
,
output
.
end
(),
"x"
),
to_string
(
input
[
0
]),
to_string_range
(
padding
.
begin
()
+
2
,
padding
.
end
(),
"x"
),
to_string_range
(
stride
.
begin
()
+
2
,
stride
.
end
(),
"x"
),
to_string_range
(
dilation
.
begin
()
+
2
,
dilation
.
end
(),
"x"
),
std
::
string
{
"0"
},
get_layout
(
pp
.
inputs
[
0
],
"NCHW"
),
get_layout
(
pp
.
inputs
[
1
],
"NCHW"
),
get_layout
(
pp
.
output
,
"NCHW"
),
get_type
(
pp
.
inputs
[
0
]),
std
::
string
{
"F"
}
},
"-"
);
return
to_string_range
({
to_string
(
input
[
1
]),
to_string_range
(
input
.
begin
()
+
2
,
input
.
end
(),
"x"
),
to_string_range
(
weights
.
begin
()
+
2
,
weights
.
end
(),
"x"
),
to_string
(
weights
[
0
]),
to_string_range
(
output
.
begin
()
+
2
,
output
.
end
(),
"x"
),
to_string
(
input
[
0
]),
to_string_range
(
padding
.
begin
()
+
2
,
padding
.
end
(),
"x"
),
to_string_range
(
stride
.
begin
()
+
2
,
stride
.
end
(),
"x"
),
to_string_range
(
dilation
.
begin
()
+
2
,
dilation
.
end
(),
"x"
),
std
::
string
{
"0"
},
get_layout
(
pp
.
inputs
[
0
],
"NCHW"
),
get_layout
(
pp
.
inputs
[
1
],
"NCHW"
),
get_layout
(
pp
.
output
,
"NCHW"
),
get_type
(
pp
.
inputs
[
0
]),
std
::
string
{
"F"
}},
"-"
);
}
}
}
// namespace
std
::
string
get_mlir_perf_for_conv
(
const
problem_params
&
pp
)
{
const
auto
dbpath
=
fs
::
path
{
"opt"
}
/
"rocm"
/
"share"
/
"miopen"
/
"db"
/
"miopen.db"
;
auto
db
=
sqlite
::
read
(
dbpath
);
const
auto
dbpath
=
fs
::
path
{
"opt"
}
/
"rocm"
/
"share"
/
"miopen"
/
"db"
/
"miopen.db"
;
auto
db
=
sqlite
::
read
(
dbpath
);
std
::
string
query
=
"select * from perf_db where config=${config}"
;
auto
results
=
db
.
execute
(
interpolate_string
(
query
,
{{
"config"
,
generate_miopen_config
(
pp
)}}));
if
(
results
.
empty
())
if
(
results
.
empty
())
return
""
;
return
results
.
front
().
at
(
"params"
);
}
...
...
@@ -91,4 +89,3 @@ std::string get_mlir_perf_for_conv(const problem_params& pp)
}
// namespace gpu
}
// namespace MIGRAPHX_INLINE_NS
}
// namespace migraphx
Write
Preview
Markdown
is supported
0%
Try again
or
attach a new file
.
Attach a file
Cancel
You are about to add
0
people
to the discussion. Proceed with caution.
Finish editing this message first!
Cancel
Please
register
or
sign in
to comment