Skip to content
GitLab
Menu
Projects
Groups
Snippets
Loading...
Help
Help
Support
Community forum
Keyboard shortcuts
?
Submit feedback
Contribute to GitLab
Sign in
Toggle navigation
Menu
Open sidebar
gaoqiong
MIGraphX
Commits
d24755de
Commit
d24755de
authored
Nov 14, 2018
by
Paul
Browse files
Formatting
parent
96358e41
Changes
13
Hide whitespace changes
Inline
Side-by-side
Showing
13 changed files
with
99 additions
and
80 deletions
+99
-80
src/include/migraphx/matcher.hpp
src/include/migraphx/matcher.hpp
+11
-11
src/include/migraphx/requires.hpp
src/include/migraphx/requires.hpp
+3
-3
src/onnx/onnx.cpp
src/onnx/onnx.cpp
+1
-1
src/onnx/verify_onnx.cpp
src/onnx/verify_onnx.cpp
+4
-3
test/cpu_ops_test.cpp
test/cpu_ops_test.cpp
+14
-14
test/eliminate_allocation_test.cpp
test/eliminate_allocation_test.cpp
+4
-3
test/eliminate_concat_test.cpp
test/eliminate_concat_test.cpp
+18
-16
test/fwd_conv_batchnorm_rewrite_test.cpp
test/fwd_conv_batchnorm_rewrite_test.cpp
+4
-3
test/gpu/miopen.cpp
test/gpu/miopen.cpp
+19
-11
test/include/basic_ops.hpp
test/include/basic_ops.hpp
+3
-2
test/memory_coloring_test.cpp
test/memory_coloring_test.cpp
+2
-2
test/op_shape_test.cpp
test/op_shape_test.cpp
+9
-6
test/operation.cpp
test/operation.cpp
+7
-5
No files found.
src/include/migraphx/matcher.hpp
View file @
d24755de
...
@@ -169,21 +169,21 @@ basic_matcher<predicate_matcher<P>> make_basic_pred_matcher(P p)
...
@@ -169,21 +169,21 @@ basic_matcher<predicate_matcher<P>> make_basic_pred_matcher(P p)
}
}
/// This macro takes care of the boilerplate for defining a matcher
/// This macro takes care of the boilerplate for defining a matcher
#define MIGRAPH_BASIC_MATCHER(name, ...) \
#define MIGRAPH_BASIC_MATCHER(name, ...)
\
struct name##_m \
struct name##_m
\
{ \
{
\
instruction_ref match(__VA_ARGS__) const; \
instruction_ref match(__VA_ARGS__) const;
\
}; \
};
\
const constexpr auto name = migraphx::match::basic_matcher<name##_m>{{}}; \
const constexpr auto name = migraphx::match::basic_matcher<name##_m>{{}}; \
inline instruction_ref name##_m::match(__VA_ARGS__) const
inline instruction_ref name##_m::match(__VA_ARGS__) const
/// This macro takes care of the boilerplate for defining a predicate matcher
/// This macro takes care of the boilerplate for defining a predicate matcher
#define MIGRAPH_PRED_MATCHER(name, ...) \
#define MIGRAPH_PRED_MATCHER(name, ...)
\
struct name##_m \
struct name##_m
\
{ \
{
\
bool operator()(__VA_ARGS__) const; \
bool operator()(__VA_ARGS__) const;
\
}; \
};
\
const constexpr auto name = \
const constexpr auto name =
\
migraphx::match::basic_matcher<migraphx::match::predicate_matcher<name##_m>>{{}}; \
migraphx::match::basic_matcher<migraphx::match::predicate_matcher<name##_m>>{{}}; \
inline bool name##_m::operator()(__VA_ARGS__) const
inline bool name##_m::operator()(__VA_ARGS__) const
...
...
src/include/migraphx/requires.hpp
View file @
d24755de
...
@@ -32,14 +32,14 @@ struct requires_enum
...
@@ -32,14 +32,14 @@ struct requires_enum
#if 0
#if 0
// TODO: This currently crashed on clang
// TODO: This currently crashed on clang
#define MIGRAPH_REQUIRES(...) \
#define MIGRAPH_REQUIRES(...) \
typename migraphx::requires_enum<__LINE__>::e MIGRAPH_REQUIRES_CAT(
\
typename migraphx::requires_enum<__LINE__>::e MIGRAPH_REQUIRES_CAT( \
PrivateRequires, \
PrivateRequires, \
__LINE__) = migraphx::requires_enum<__LINE__>::a,
\
__LINE__) = migraphx::requires_enum<__LINE__>::a, \
class = typename std::enable_if<and_<__VA_ARGS__, \
class = typename std::enable_if<and_<__VA_ARGS__, \
MIGRAPH_REQUIRES_CAT(PrivateRequires, __LINE__) == \
MIGRAPH_REQUIRES_CAT(PrivateRequires, __LINE__) == \
migraphx::requires_enum<__LINE__>::a>{}>::type
migraphx::requires_enum<__LINE__>::a>{}>::type
#else
#else
#define MIGRAPH_REQUIRES(...) \
#define MIGRAPH_REQUIRES(...)
\
typename migraphx::requires_enum<__LINE__>::e MIGRAPH_REQUIRES_CAT( \
typename migraphx::requires_enum<__LINE__>::e MIGRAPH_REQUIRES_CAT( \
PrivateRequires, __LINE__) = migraphx::requires_enum<__LINE__>::a, \
PrivateRequires, __LINE__) = migraphx::requires_enum<__LINE__>::a, \
class = typename std::enable_if<and_<__VA_ARGS__>{}>::type
class = typename std::enable_if<and_<__VA_ARGS__>{}>::type
...
...
src/onnx/onnx.cpp
View file @
d24755de
...
@@ -410,7 +410,7 @@ struct onnx_parser
...
@@ -410,7 +410,7 @@ struct onnx_parser
auto
scale_tensor
=
prog
.
add_instruction
(
migraphx
::
op
::
scalar
{
input_shape
},
scale_val
);
auto
scale_tensor
=
prog
.
add_instruction
(
migraphx
::
op
::
scalar
{
input_shape
},
scale_val
);
auto
img_scaled
=
prog
.
add_instruction
(
migraphx
::
op
::
mul
{},
args
.
front
(),
scale_tensor
);
auto
img_scaled
=
prog
.
add_instruction
(
migraphx
::
op
::
mul
{},
args
.
front
(),
scale_tensor
);
auto
bias_bcast
=
prog
.
add_instruction
(
migraphx
::
op
::
broadcast
{
1
,
input_shape
},
bias_vals
);
auto
bias_bcast
=
prog
.
add_instruction
(
migraphx
::
op
::
broadcast
{
1
,
input_shape
},
bias_vals
);
return
prog
.
add_instruction
(
migraphx
::
op
::
add
{},
img_scaled
,
bias_bcast
);
return
prog
.
add_instruction
(
migraphx
::
op
::
add
{},
img_scaled
,
bias_bcast
);
}
}
...
...
src/onnx/verify_onnx.cpp
View file @
d24755de
...
@@ -38,7 +38,8 @@ migraphx::argument run_gpu(F f)
...
@@ -38,7 +38,8 @@ migraphx::argument run_gpu(F f)
migraphx
::
program
::
parameter_map
m
;
migraphx
::
program
::
parameter_map
m
;
for
(
auto
&&
x
:
p
.
get_parameter_shapes
())
for
(
auto
&&
x
:
p
.
get_parameter_shapes
())
{
{
m
[
x
.
first
]
=
migraphx
::
gpu
::
to_gpu
(
migraphx
::
generate_argument
(
x
.
second
,
get_hash
(
x
.
first
)));
m
[
x
.
first
]
=
migraphx
::
gpu
::
to_gpu
(
migraphx
::
generate_argument
(
x
.
second
,
get_hash
(
x
.
first
)));
}
}
auto
out
=
migraphx
::
gpu
::
from_gpu
(
p
.
eval
(
m
));
auto
out
=
migraphx
::
gpu
::
from_gpu
(
p
.
eval
(
m
));
std
::
cout
<<
p
<<
std
::
endl
;
std
::
cout
<<
p
<<
std
::
endl
;
...
@@ -101,7 +102,7 @@ void verify_reduced(F f, int n, double tolerance = 80)
...
@@ -101,7 +102,7 @@ void verify_reduced(F f, int n, double tolerance = 80)
auto
create_program
=
[
&
]
{
auto
create_program
=
[
&
]
{
migraphx
::
program
p
=
f
();
migraphx
::
program
p
=
f
();
auto
last
=
std
::
prev
(
p
.
end
(),
n
+
1
);
auto
last
=
std
::
prev
(
p
.
end
(),
n
+
1
);
p
.
remove_instructions
(
last
,
p
.
end
());
p
.
remove_instructions
(
last
,
p
.
end
());
return
p
;
return
p
;
};
};
...
@@ -114,7 +115,7 @@ template <class F>
...
@@ -114,7 +115,7 @@ template <class F>
void
verify_reduced_program
(
F
f
,
double
tolerance
=
80
)
void
verify_reduced_program
(
F
f
,
double
tolerance
=
80
)
{
{
migraphx
::
program
p
=
f
();
migraphx
::
program
p
=
f
();
auto
n
=
std
::
distance
(
p
.
begin
(),
p
.
end
());
auto
n
=
std
::
distance
(
p
.
begin
(),
p
.
end
());
for
(
int
i
=
0
;
i
<
n
;
i
++
)
for
(
int
i
=
0
;
i
<
n
;
i
++
)
{
{
verify_reduced
(
f
,
i
,
tolerance
);
verify_reduced
(
f
,
i
,
tolerance
);
...
...
test/cpu_ops_test.cpp
View file @
d24755de
...
@@ -605,20 +605,20 @@ TEST_CASE(imagescaler_test)
...
@@ -605,20 +605,20 @@ TEST_CASE(imagescaler_test)
migraphx
::
program
p
;
migraphx
::
program
p
;
migraphx
::
shape
s
{
migraphx
::
shape
::
float_type
,
{
1
,
3
,
2
,
2
}};
migraphx
::
shape
s
{
migraphx
::
shape
::
float_type
,
{
1
,
3
,
2
,
2
}};
auto
img
=
p
.
add_literal
(
migraphx
::
literal
{
s
,
auto
img
=
p
.
add_literal
(
migraphx
::
literal
{
s
,
{
0.2
,
{
0.2
,
0.3
,
0.3
,
0.5
,
0.5
,
0.4
,
0.4
,
0.7
,
0.7
,
0.8
,
0.8
,
0.1
,
0.1
,
0.9
,
0.9
,
0.15
,
0.15
,
0.25
,
0.25
,
0.35
,
0.35
,
0.45
}});
0.45
}});
auto
scale_val
=
p
.
add_literal
(
2.
f
);
auto
scale_val
=
p
.
add_literal
(
2.
f
);
auto
scaled_tensor
=
p
.
add_instruction
(
migraphx
::
op
::
scalar
{
s
},
scale_val
);
auto
scaled_tensor
=
p
.
add_instruction
(
migraphx
::
op
::
scalar
{
s
},
scale_val
);
auto
img_scaled
=
p
.
add_instruction
(
migraphx
::
op
::
mul
{},
img
,
scaled_tensor
);
auto
img_scaled
=
p
.
add_instruction
(
migraphx
::
op
::
mul
{},
img
,
scaled_tensor
);
...
...
test/eliminate_allocation_test.cpp
View file @
d24755de
...
@@ -10,7 +10,8 @@ struct eliminate_allocation_target
...
@@ -10,7 +10,8 @@ struct eliminate_allocation_target
std
::
string
name
()
const
{
return
"eliminate_allocation"
;
}
std
::
string
name
()
const
{
return
"eliminate_allocation"
;
}
std
::
vector
<
migraphx
::
pass
>
get_passes
(
migraphx
::
context
&
)
const
std
::
vector
<
migraphx
::
pass
>
get_passes
(
migraphx
::
context
&
)
const
{
{
return
{
migraphx
::
eliminate_allocation
{
"allocate"
,
align
},
migraphx
::
dead_code_elimination
{}};
return
{
migraphx
::
eliminate_allocation
{
"allocate"
,
align
},
migraphx
::
dead_code_elimination
{}};
}
}
migraphx
::
context
get_context
()
const
{
return
{};
}
migraphx
::
context
get_context
()
const
{
return
{};
}
};
};
...
@@ -25,8 +26,8 @@ struct allocate
...
@@ -25,8 +26,8 @@ struct allocate
return
s
;
return
s
;
}
}
migraphx
::
argument
compute
(
migraphx
::
context
&
,
migraphx
::
argument
compute
(
migraphx
::
context
&
,
const
migraphx
::
shape
&
output_shape
,
const
migraphx
::
shape
&
output_shape
,
const
std
::
vector
<
migraphx
::
argument
>&
)
const
const
std
::
vector
<
migraphx
::
argument
>&
)
const
{
{
return
{
output_shape
};
return
{
output_shape
};
}
}
...
...
test/eliminate_concat_test.cpp
View file @
d24755de
...
@@ -14,8 +14,8 @@ struct concat
...
@@ -14,8 +14,8 @@ struct concat
return
op
.
compute_shape
(
std
::
move
(
inputs
));
return
op
.
compute_shape
(
std
::
move
(
inputs
));
}
}
migraphx
::
argument
compute
(
migraphx
::
context
&
,
migraphx
::
argument
compute
(
migraphx
::
context
&
,
const
migraphx
::
shape
&
output_shape
,
const
migraphx
::
shape
&
output_shape
,
const
std
::
vector
<
migraphx
::
argument
>&
)
const
const
std
::
vector
<
migraphx
::
argument
>&
)
const
{
{
return
{
output_shape
};
return
{
output_shape
};
}
}
...
@@ -56,8 +56,8 @@ struct allocate
...
@@ -56,8 +56,8 @@ struct allocate
return
s
;
return
s
;
}
}
migraphx
::
argument
compute
(
migraphx
::
context
&
,
migraphx
::
argument
compute
(
migraphx
::
context
&
,
const
migraphx
::
shape
&
output_shape
,
const
migraphx
::
shape
&
output_shape
,
const
std
::
vector
<
migraphx
::
argument
>&
)
const
const
std
::
vector
<
migraphx
::
argument
>&
)
const
{
{
return
{
output_shape
};
return
{
output_shape
};
}
}
...
@@ -72,8 +72,8 @@ struct fred_op
...
@@ -72,8 +72,8 @@ struct fred_op
return
inputs
.
at
(
0
);
return
inputs
.
at
(
0
);
}
}
migraphx
::
argument
compute
(
migraphx
::
context
&
,
migraphx
::
argument
compute
(
migraphx
::
context
&
,
const
migraphx
::
shape
&
,
const
migraphx
::
shape
&
,
const
std
::
vector
<
migraphx
::
argument
>&
args
)
const
const
std
::
vector
<
migraphx
::
argument
>&
args
)
const
{
{
return
args
.
at
(
0
);
return
args
.
at
(
0
);
}
}
...
@@ -93,20 +93,22 @@ TEST_CASE(basic)
...
@@ -93,20 +93,22 @@ TEST_CASE(basic)
p
.
add_instruction
(
allocate
{
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
1
,
5
,
8
,
8
}}});
p
.
add_instruction
(
allocate
{
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
1
,
5
,
8
,
8
}}});
auto
p3
=
p
.
add_instruction
(
fred_op
{},
a3
);
auto
p3
=
p
.
add_instruction
(
fred_op
{},
a3
);
std
::
size_t
axis
=
1
;
std
::
size_t
axis
=
1
;
auto
a4
=
auto
a4
=
p
.
add_instruction
(
p
.
add_instruction
(
allocate
{
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
1
,
10
,
8
,
8
}}});
allocate
{
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
1
,
10
,
8
,
8
}}});
p
.
add_instruction
(
concat
(
axis
),
p1
,
p2
,
p3
,
a4
);
p
.
add_instruction
(
concat
(
axis
),
p1
,
p2
,
p3
,
a4
);
return
p
;
return
p
;
};
};
auto
create_control_program
=
[]()
{
auto
create_control_program
=
[]()
{
migraphx
::
program
p
;
migraphx
::
program
p
;
auto
a1
=
auto
a1
=
p
.
add_instruction
(
p
.
add_instruction
(
allocate
{
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
1
,
10
,
8
,
8
}}});
allocate
{
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
1
,
10
,
8
,
8
}}});
auto
l1
=
p
.
add_instruction
(
auto
l1
=
p
.
add_instruction
(
migraphx
::
op
::
load
{
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
1
,
2
,
8
,
8
}},
0
},
{
a1
});
migraphx
::
op
::
load
{
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
1
,
2
,
8
,
8
}},
0
},
{
a1
});
auto
p1
=
p
.
add_instruction
(
fred_op
{},
l1
);
auto
p1
=
p
.
add_instruction
(
fred_op
{},
l1
);
auto
l2
=
p
.
add_instruction
(
auto
l2
=
p
.
add_instruction
(
migraphx
::
op
::
load
{
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
1
,
3
,
8
,
8
}},
512
},
{
a1
});
migraphx
::
op
::
load
{
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
1
,
3
,
8
,
8
}},
512
},
{
a1
});
auto
p2
=
p
.
add_instruction
(
fred_op
{},
l2
);
auto
p2
=
p
.
add_instruction
(
fred_op
{},
l2
);
auto
l3
=
p
.
add_instruction
(
auto
l3
=
p
.
add_instruction
(
migraphx
::
op
::
load
{
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
1
,
5
,
8
,
8
}},
1280
},
migraphx
::
op
::
load
{
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
1
,
5
,
8
,
8
}},
1280
},
...
@@ -137,8 +139,8 @@ TEST_CASE(wont_work)
...
@@ -137,8 +139,8 @@ TEST_CASE(wont_work)
p
.
add_instruction
(
allocate
{
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
2
,
5
,
8
,
8
}}});
p
.
add_instruction
(
allocate
{
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
2
,
5
,
8
,
8
}}});
auto
p3
=
p
.
add_instruction
(
fred_op
{},
a3
);
auto
p3
=
p
.
add_instruction
(
fred_op
{},
a3
);
std
::
size_t
axis
=
1
;
std
::
size_t
axis
=
1
;
auto
a4
=
auto
a4
=
p
.
add_instruction
(
p
.
add_instruction
(
allocate
{
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
2
,
10
,
8
,
8
}}});
allocate
{
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
2
,
10
,
8
,
8
}}});
p
.
add_instruction
(
concat
(
axis
),
p1
,
p2
,
p3
,
a4
);
p
.
add_instruction
(
concat
(
axis
),
p1
,
p2
,
p3
,
a4
);
return
p
;
return
p
;
};
};
...
@@ -154,8 +156,8 @@ TEST_CASE(wont_work)
...
@@ -154,8 +156,8 @@ TEST_CASE(wont_work)
p
.
add_instruction
(
allocate
{
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
2
,
5
,
8
,
8
}}});
p
.
add_instruction
(
allocate
{
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
2
,
5
,
8
,
8
}}});
auto
p3
=
p
.
add_instruction
(
fred_op
{},
a3
);
auto
p3
=
p
.
add_instruction
(
fred_op
{},
a3
);
std
::
size_t
axis
=
1
;
std
::
size_t
axis
=
1
;
auto
a4
=
auto
a4
=
p
.
add_instruction
(
p
.
add_instruction
(
allocate
{
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
2
,
10
,
8
,
8
}}});
allocate
{
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
2
,
10
,
8
,
8
}}});
p
.
add_instruction
(
concat
(
axis
),
p1
,
p2
,
p3
,
a4
);
p
.
add_instruction
(
concat
(
axis
),
p1
,
p2
,
p3
,
a4
);
return
p
;
return
p
;
};
};
...
...
test/fwd_conv_batchnorm_rewrite_test.cpp
View file @
d24755de
...
@@ -36,9 +36,10 @@ TEST_CASE(fwd_conv_batchnorm_rewrite_test)
...
@@ -36,9 +36,10 @@ TEST_CASE(fwd_conv_batchnorm_rewrite_test)
auto
create_program
=
[
&
]()
{
auto
create_program
=
[
&
]()
{
migraphx
::
program
p
;
migraphx
::
program
p
;
auto
x
=
p
.
add_literal
(
xs
,
xdata
);
auto
x
=
p
.
add_literal
(
xs
,
xdata
);
auto
w
=
p
.
add_literal
(
ws
,
wdata
);
auto
w
=
p
.
add_literal
(
ws
,
wdata
);
auto
conv
=
p
.
add_instruction
(
migraphx
::
op
::
convolution
{{{
0
,
0
}},
{{
1
,
1
}},
{{
1
,
1
}}},
x
,
w
);
auto
conv
=
p
.
add_instruction
(
migraphx
::
op
::
convolution
{{{
0
,
0
}},
{{
1
,
1
}},
{{
1
,
1
}}},
x
,
w
);
auto
scale
=
p
.
add_literal
(
migraphx
::
literal
{
vars
,
{
3.0
f
}});
auto
scale
=
p
.
add_literal
(
migraphx
::
literal
{
vars
,
{
3.0
f
}});
auto
bias
=
p
.
add_literal
(
migraphx
::
literal
{
vars
,
{
8.1
f
}});
auto
bias
=
p
.
add_literal
(
migraphx
::
literal
{
vars
,
{
8.1
f
}});
auto
mean
=
p
.
add_literal
(
migraphx
::
literal
{
vars
,
{
4.0
f
}});
auto
mean
=
p
.
add_literal
(
migraphx
::
literal
{
vars
,
{
4.0
f
}});
...
...
test/gpu/miopen.cpp
View file @
d24755de
...
@@ -119,7 +119,8 @@ migraphx::argument run_gpu(migraphx::program& p)
...
@@ -119,7 +119,8 @@ migraphx::argument run_gpu(migraphx::program& p)
migraphx
::
program
::
parameter_map
m
;
migraphx
::
program
::
parameter_map
m
;
for
(
auto
&&
x
:
p
.
get_parameter_shapes
())
for
(
auto
&&
x
:
p
.
get_parameter_shapes
())
{
{
m
[
x
.
first
]
=
migraphx
::
gpu
::
to_gpu
(
migraphx
::
generate_argument
(
x
.
second
,
get_hash
(
x
.
first
)));
m
[
x
.
first
]
=
migraphx
::
gpu
::
to_gpu
(
migraphx
::
generate_argument
(
x
.
second
,
get_hash
(
x
.
first
)));
}
}
EXPECT
(
bool
{
m
.
find
(
"output"
)
!=
m
.
end
()});
EXPECT
(
bool
{
m
.
find
(
"output"
)
!=
m
.
end
()});
return
migraphx
::
gpu
::
from_gpu
(
p
.
eval
(
m
));
return
migraphx
::
gpu
::
from_gpu
(
p
.
eval
(
m
));
...
@@ -365,7 +366,8 @@ struct test_softmax2
...
@@ -365,7 +366,8 @@ struct test_softmax2
migraphx
::
program
create_program
()
const
migraphx
::
program
create_program
()
const
{
{
migraphx
::
program
p
;
migraphx
::
program
p
;
auto
x
=
p
.
add_parameter
(
"x"
,
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
1
,
1000
,
1
,
1
}});
auto
x
=
p
.
add_parameter
(
"x"
,
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
1
,
1000
,
1
,
1
}});
p
.
add_instruction
(
migraphx
::
op
::
softmax
{},
x
);
p
.
add_instruction
(
migraphx
::
op
::
softmax
{},
x
);
return
p
;
return
p
;
}
}
...
@@ -376,7 +378,8 @@ struct test_conv
...
@@ -376,7 +378,8 @@ struct test_conv
migraphx
::
program
create_program
()
const
migraphx
::
program
create_program
()
const
{
{
migraphx
::
program
p
;
migraphx
::
program
p
;
auto
input
=
p
.
add_parameter
(
"x"
,
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
4
,
3
,
3
,
3
}});
auto
input
=
p
.
add_parameter
(
"x"
,
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
4
,
3
,
3
,
3
}});
auto
weights
=
auto
weights
=
p
.
add_parameter
(
"w"
,
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
4
,
3
,
3
,
3
}});
p
.
add_parameter
(
"w"
,
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
4
,
3
,
3
,
3
}});
p
.
add_instruction
(
migraphx
::
op
::
convolution
{},
input
,
weights
);
p
.
add_instruction
(
migraphx
::
op
::
convolution
{},
input
,
weights
);
...
@@ -403,7 +406,8 @@ struct test_conv_relu
...
@@ -403,7 +406,8 @@ struct test_conv_relu
migraphx
::
program
create_program
()
const
migraphx
::
program
create_program
()
const
{
{
migraphx
::
program
p
;
migraphx
::
program
p
;
auto
input
=
p
.
add_parameter
(
"x"
,
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
4
,
3
,
3
,
3
}});
auto
input
=
p
.
add_parameter
(
"x"
,
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
4
,
3
,
3
,
3
}});
auto
weights
=
auto
weights
=
p
.
add_parameter
(
"w"
,
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
4
,
3
,
3
,
3
}});
p
.
add_parameter
(
"w"
,
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
4
,
3
,
3
,
3
}});
auto
conv
=
p
.
add_instruction
(
migraphx
::
op
::
convolution
{},
input
,
weights
);
auto
conv
=
p
.
add_instruction
(
migraphx
::
op
::
convolution
{},
input
,
weights
);
...
@@ -417,7 +421,8 @@ struct test_conv_relu_half
...
@@ -417,7 +421,8 @@ struct test_conv_relu_half
migraphx
::
program
create_program
()
const
migraphx
::
program
create_program
()
const
{
{
migraphx
::
program
p
;
migraphx
::
program
p
;
auto
input
=
p
.
add_parameter
(
"x"
,
migraphx
::
shape
{
migraphx
::
shape
::
half_type
,
{
4
,
3
,
3
,
3
}});
auto
input
=
p
.
add_parameter
(
"x"
,
migraphx
::
shape
{
migraphx
::
shape
::
half_type
,
{
4
,
3
,
3
,
3
}});
auto
weights
=
auto
weights
=
p
.
add_parameter
(
"w"
,
migraphx
::
shape
{
migraphx
::
shape
::
half_type
,
{
4
,
3
,
3
,
3
}});
p
.
add_parameter
(
"w"
,
migraphx
::
shape
{
migraphx
::
shape
::
half_type
,
{
4
,
3
,
3
,
3
}});
auto
conv
=
p
.
add_instruction
(
migraphx
::
op
::
convolution
{},
input
,
weights
);
auto
conv
=
p
.
add_instruction
(
migraphx
::
op
::
convolution
{},
input
,
weights
);
...
@@ -525,8 +530,10 @@ struct test_gemm_ld
...
@@ -525,8 +530,10 @@ struct test_gemm_ld
migraphx
::
program
create_program
()
const
migraphx
::
program
create_program
()
const
{
{
migraphx
::
program
p
;
migraphx
::
program
p
;
auto
a
=
p
.
add_parameter
(
"a"
,
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
4
,
5
},
{
10
,
1
}});
auto
a
=
auto
b
=
p
.
add_parameter
(
"b"
,
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
5
,
3
},
{
20
,
1
}});
p
.
add_parameter
(
"a"
,
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
4
,
5
},
{
10
,
1
}});
auto
b
=
p
.
add_parameter
(
"b"
,
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
5
,
3
},
{
20
,
1
}});
p
.
add_instruction
(
migraphx
::
op
::
dot
{},
a
,
b
);
p
.
add_instruction
(
migraphx
::
op
::
dot
{},
a
,
b
);
return
p
;
return
p
;
}
}
...
@@ -797,10 +804,11 @@ struct test_conv_bn_relu_pooling2
...
@@ -797,10 +804,11 @@ struct test_conv_bn_relu_pooling2
add_bn
(
migraphx
::
program
&
p
,
migraphx
::
instruction_ref
x
,
std
::
size_t
channels
)
add_bn
(
migraphx
::
program
&
p
,
migraphx
::
instruction_ref
x
,
std
::
size_t
channels
)
{
{
migraphx
::
shape
vars
{
migraphx
::
shape
::
float_type
,
{
channels
}};
migraphx
::
shape
vars
{
migraphx
::
shape
::
float_type
,
{
channels
}};
auto
scale
=
p
.
add_literal
(
migraphx
::
abs
(
migraphx
::
generate_literal
(
vars
,
1
+
channels
)));
auto
scale
=
p
.
add_literal
(
migraphx
::
abs
(
migraphx
::
generate_literal
(
vars
,
1
+
channels
)));
auto
bias
=
p
.
add_literal
(
migraphx
::
abs
(
migraphx
::
generate_literal
(
vars
,
2
+
channels
)));
auto
bias
=
p
.
add_literal
(
migraphx
::
abs
(
migraphx
::
generate_literal
(
vars
,
2
+
channels
)));
auto
mean
=
p
.
add_literal
(
migraphx
::
abs
(
migraphx
::
generate_literal
(
vars
,
3
+
channels
)));
auto
mean
=
p
.
add_literal
(
migraphx
::
abs
(
migraphx
::
generate_literal
(
vars
,
3
+
channels
)));
auto
variance
=
p
.
add_literal
(
migraphx
::
abs
(
migraphx
::
generate_literal
(
vars
,
4
+
channels
)));
auto
variance
=
p
.
add_literal
(
migraphx
::
abs
(
migraphx
::
generate_literal
(
vars
,
4
+
channels
)));
return
p
.
add_instruction
(
return
p
.
add_instruction
(
migraphx
::
op
::
batch_norm_inference
{},
x
,
scale
,
bias
,
mean
,
variance
);
migraphx
::
op
::
batch_norm_inference
{},
x
,
scale
,
bias
,
mean
,
variance
);
}
}
...
...
test/include/basic_ops.hpp
View file @
d24755de
...
@@ -110,8 +110,9 @@ struct pass_standard_op
...
@@ -110,8 +110,9 @@ struct pass_standard_op
struct
nop
struct
nop
{
{
std
::
string
name
()
const
{
return
"nop"
;
}
std
::
string
name
()
const
{
return
"nop"
;
}
migraphx
::
argument
migraphx
::
argument
compute
(
migraphx
::
context
&
,
compute
(
migraphx
::
context
&
,
const
migraphx
::
shape
&
,
const
std
::
vector
<
migraphx
::
argument
>&
)
const
const
migraphx
::
shape
&
,
const
std
::
vector
<
migraphx
::
argument
>&
)
const
{
{
return
{};
return
{};
}
}
...
...
test/memory_coloring_test.cpp
View file @
d24755de
...
@@ -25,8 +25,8 @@ struct allocate
...
@@ -25,8 +25,8 @@ struct allocate
return
inputs
.
front
();
return
inputs
.
front
();
}
}
migraphx
::
argument
compute
(
migraphx
::
context
&
,
migraphx
::
argument
compute
(
migraphx
::
context
&
,
const
migraphx
::
shape
&
output_shape
,
const
migraphx
::
shape
&
output_shape
,
const
std
::
vector
<
migraphx
::
argument
>&
)
const
const
std
::
vector
<
migraphx
::
argument
>&
)
const
{
{
return
{
output_shape
};
return
{
output_shape
};
}
}
...
...
test/op_shape_test.cpp
View file @
d24755de
...
@@ -120,12 +120,15 @@ TEST_CASE(flatten_shape)
...
@@ -120,12 +120,15 @@ TEST_CASE(flatten_shape)
expect_shape
(
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
1
,
2
*
4
*
6
*
8
}},
expect_shape
(
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
1
,
2
*
4
*
6
*
8
}},
migraphx
::
op
::
flatten
{
0
},
migraphx
::
op
::
flatten
{
0
},
input
);
input
);
expect_shape
(
expect_shape
(
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
2
,
4
*
6
*
8
}},
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
2
,
4
*
6
*
8
}},
migraphx
::
op
::
flatten
{
1
},
input
);
migraphx
::
op
::
flatten
{
1
},
expect_shape
(
input
);
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
2
*
4
,
6
*
8
}},
migraphx
::
op
::
flatten
{
2
},
input
);
expect_shape
(
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
2
*
4
,
6
*
8
}},
expect_shape
(
migraphx
::
op
::
flatten
{
2
},
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
2
*
4
*
6
,
8
}},
migraphx
::
op
::
flatten
{
3
},
input
);
input
);
expect_shape
(
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
2
*
4
*
6
,
8
}},
migraphx
::
op
::
flatten
{
3
},
input
);
expect_shape
(
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
2
*
4
*
6
*
8
,
1
}},
expect_shape
(
migraphx
::
shape
{
migraphx
::
shape
::
float_type
,
{
2
*
4
*
6
*
8
,
1
}},
migraphx
::
op
::
flatten
{
4
},
migraphx
::
op
::
flatten
{
4
},
input
);
input
);
...
...
test/operation.cpp
View file @
d24755de
...
@@ -17,8 +17,9 @@ struct simple_operation
...
@@ -17,8 +17,9 @@ struct simple_operation
{
{
MIGRAPH_THROW
(
"not computable"
);
MIGRAPH_THROW
(
"not computable"
);
}
}
migraphx
::
argument
migraphx
::
argument
compute
(
migraphx
::
context
&
,
compute
(
migraphx
::
context
&
,
const
migraphx
::
shape
&
,
const
std
::
vector
<
migraphx
::
argument
>&
)
const
const
migraphx
::
shape
&
,
const
std
::
vector
<
migraphx
::
argument
>&
)
const
{
{
MIGRAPH_THROW
(
"not computable"
);
MIGRAPH_THROW
(
"not computable"
);
}
}
...
@@ -36,8 +37,9 @@ struct simple_operation_no_print
...
@@ -36,8 +37,9 @@ struct simple_operation_no_print
{
{
MIGRAPH_THROW
(
"not computable"
);
MIGRAPH_THROW
(
"not computable"
);
}
}
migraphx
::
argument
migraphx
::
argument
compute
(
migraphx
::
context
&
,
compute
(
migraphx
::
context
&
,
const
migraphx
::
shape
&
,
const
std
::
vector
<
migraphx
::
argument
>&
)
const
const
migraphx
::
shape
&
,
const
std
::
vector
<
migraphx
::
argument
>&
)
const
{
{
MIGRAPH_THROW
(
"not computable"
);
MIGRAPH_THROW
(
"not computable"
);
}
}
...
@@ -58,7 +60,7 @@ TEST_CASE(operation_equal_test)
...
@@ -58,7 +60,7 @@ TEST_CASE(operation_equal_test)
{
{
simple_operation
s
{};
simple_operation
s
{};
migraphx
::
operation
op1
=
s
;
migraphx
::
operation
op1
=
s
;
s
.
data
=
2
;
s
.
data
=
2
;
migraphx
::
operation
op2
=
op1
;
// NOLINT
migraphx
::
operation
op2
=
op1
;
// NOLINT
migraphx
::
operation
op3
=
s
;
// NOLINT
migraphx
::
operation
op3
=
s
;
// NOLINT
...
...
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