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
a693711d
Commit
a693711d
authored
Jun 12, 2019
by
Khalique
Browse files
formatting
parent
75d5c660
Changes
5
Hide whitespace changes
Inline
Side-by-side
Showing
5 changed files
with
33 additions
and
24 deletions
+33
-24
src/include/migraphx/pad_calc.hpp
src/include/migraphx/pad_calc.hpp
+10
-4
src/onnx/onnx.cpp
src/onnx/onnx.cpp
+2
-1
src/targets/gpu/device/pad.cpp
src/targets/gpu/device/pad.cpp
+3
-2
src/tf/tf.cpp
src/tf/tf.cpp
+17
-16
test/gpu/miopen.cpp
test/gpu/miopen.cpp
+1
-1
No files found.
src/include/migraphx/pad_calc.hpp
View file @
a693711d
...
...
@@ -11,12 +11,18 @@ inline std::size_t calculate_padding(std::size_t weight_dim, std::size_t dilatio
return
(
dilation
*
(
weight_dim
-
1
))
/
2
;
}
inline
void
calculate_padding
(
int64_t
idx
,
std
::
vector
<
int64_t
>&
pads
,
int64_t
input_dim
,
int64_t
stride
,
int64_t
dilation
,
int64_t
weight_dim
)
inline
void
calculate_padding
(
int64_t
idx
,
std
::
vector
<
int64_t
>&
pads
,
int64_t
input_dim
,
int64_t
stride
,
int64_t
dilation
,
int64_t
weight_dim
)
{
int64_t
output_dim
=
input_dim
/
stride
;
int64_t
pad
=
std
::
max
(
static_cast
<
int64_t
>
(
0
),
(
output_dim
-
1
)
*
stride
+
dilation
*
weight_dim
-
input_dim
);
pads
[
idx
]
=
pad
/
2
;
pads
[
idx
+
2
]
=
pad
-
pad
/
2
;
int64_t
pad
=
std
::
max
(
static_cast
<
int64_t
>
(
0
),
(
output_dim
-
1
)
*
stride
+
dilation
*
weight_dim
-
input_dim
);
pads
[
idx
]
=
pad
/
2
;
pads
[
idx
+
2
]
=
pad
-
pad
/
2
;
}
}
// namespace MIGRAPHX_INLINE_NS
...
...
src/onnx/onnx.cpp
View file @
a693711d
...
...
@@ -352,7 +352,8 @@ struct onnx_parser
{
// insert zeros for pad op (args[0] has 4 dims)
padding
=
{
0
,
0
,
padding
[
0
],
padding
[
1
],
0
,
0
,
padding
[
2
],
padding
[
3
]};
l0
=
prog
.
add_instruction
(
op
::
pad
{
padding
,
std
::
numeric_limits
<
float
>::
lowest
()},
l0
);
l0
=
prog
.
add_instruction
(
op
::
pad
{
padding
,
std
::
numeric_limits
<
float
>::
lowest
()},
l0
);
}
else
{
...
...
src/targets/gpu/device/pad.cpp
View file @
a693711d
...
...
@@ -19,11 +19,12 @@ pad(hipStream_t stream, argument result, argument arg1, float value, std::vector
// visit_all(result)([&](auto output) {
// auto* outptr = output.data();
// gs_launch(stream, nelements)([=](auto i) {
// outptr[i] = std::numeric_limits<typename decltype(output)::value_type>::lowest();
// outptr[i] = std::numeric_limits<typename
// decltype(output)::value_type>::lowest();
// });
// });
// }
// else
// {
// visit_all(result)([&](auto output) {
...
...
src/tf/tf.cpp
View file @
a693711d
...
...
@@ -329,8 +329,8 @@ struct tf_parser
size_t
weight_w
=
weight_dims
[
3
];
auto
input_dims
=
l0
->
get_shape
().
lens
();
size_t
input_h
=
input_dims
[
2
];
size_t
input_w
=
input_dims
[
3
];
size_t
input_h
=
input_dims
[
2
];
size_t
input_w
=
input_dims
[
3
];
std
::
vector
<
int64_t
>
pads
(
input_dims
.
size
());
calculate_padding
(
0
,
pads
,
input_h
,
op
.
stride
[
0
],
op
.
dilation
[
0
],
weight_h
);
calculate_padding
(
1
,
pads
,
input_w
,
op
.
stride
[
1
],
op
.
dilation
[
1
],
weight_w
);
...
...
@@ -342,8 +342,8 @@ struct tf_parser
}
else
{
op
.
padding
[
0
]
=
pads
[
0
];
op
.
padding
[
1
]
=
pads
[
1
];
op
.
padding
[
0
]
=
pads
[
0
];
op
.
padding
[
1
]
=
pads
[
1
];
}
}
else
if
(
pad_mode
.
find
(
"VALID"
)
!=
std
::
string
::
npos
)
...
...
@@ -420,8 +420,8 @@ struct tf_parser
auto
l0
=
args
[
0
];
if
(
contains
(
attributes
,
"padding"
))
{
const
std
::
string
&
pad_mode
=
attributes
.
at
(
"padding"
).
s
();
const
std
::
string
&
pad_mode
=
attributes
.
at
(
"padding"
).
s
();
if
(
pad_mode
.
find
(
"SAME"
)
!=
std
::
string
::
npos
)
{
// op.padding_mode = op::padding_mode_t::same;
...
...
@@ -430,8 +430,8 @@ struct tf_parser
size_t
weight_w
=
weight_dims
[
3
];
auto
input_dims
=
l0
->
get_shape
().
lens
();
size_t
input_h
=
input_dims
[
2
];
size_t
input_w
=
input_dims
[
3
];
size_t
input_h
=
input_dims
[
2
];
size_t
input_w
=
input_dims
[
3
];
std
::
vector
<
int64_t
>
pads
(
input_dims
.
size
());
calculate_padding
(
0
,
pads
,
input_h
,
op
.
stride
[
0
],
op
.
dilation
[
0
],
weight_h
);
calculate_padding
(
1
,
pads
,
input_w
,
op
.
stride
[
1
],
op
.
dilation
[
1
],
weight_w
);
...
...
@@ -443,8 +443,8 @@ struct tf_parser
}
else
{
op
.
padding
[
0
]
=
pads
[
0
];
op
.
padding
[
1
]
=
pads
[
1
];
op
.
padding
[
0
]
=
pads
[
0
];
op
.
padding
[
1
]
=
pads
[
1
];
}
}
else
if
(
pad_mode
.
find
(
"VALID"
)
!=
std
::
string
::
npos
)
...
...
@@ -609,10 +609,10 @@ struct tf_parser
const
std
::
string
&
pad_mode
=
attributes
.
at
(
"padding"
).
s
();
if
(
pad_mode
.
find
(
"SAME"
)
!=
std
::
string
::
npos
)
{
//op.padding_mode = op::padding_mode_t::same;
//
op.padding_mode = op::padding_mode_t::same;
auto
input_dims
=
l0
->
get_shape
().
lens
();
size_t
input_h
=
input_dims
[
2
];
size_t
input_w
=
input_dims
[
3
];
size_t
input_h
=
input_dims
[
2
];
size_t
input_w
=
input_dims
[
3
];
std
::
vector
<
int64_t
>
pads
(
input_dims
.
size
());
calculate_padding
(
0
,
pads
,
input_h
,
op
.
stride
[
0
],
1
,
op
.
lengths
[
0
]);
calculate_padding
(
1
,
pads
,
input_w
,
op
.
stride
[
1
],
1
,
op
.
lengths
[
1
]);
...
...
@@ -624,12 +624,13 @@ struct tf_parser
if
(
pads
[
0
]
!=
pads
[
2
]
||
pads
[
1
]
!=
pads
[
3
])
{
std
::
vector
<
int64_t
>
padding
=
{
0
,
0
,
pads
[
0
],
pads
[
1
],
0
,
0
,
pads
[
2
],
pads
[
3
]};
l0
=
prog
.
add_instruction
(
migraphx
::
op
::
pad
{
padding
,
std
::
numeric_limits
<
float
>::
lowest
()},
l0
);
l0
=
prog
.
add_instruction
(
migraphx
::
op
::
pad
{
padding
,
std
::
numeric_limits
<
float
>::
lowest
()},
l0
);
}
else
{
op
.
padding
[
0
]
=
pads
[
0
];
op
.
padding
[
1
]
=
pads
[
1
];
op
.
padding
[
0
]
=
pads
[
0
];
op
.
padding
[
1
]
=
pads
[
1
];
}
}
else
if
(
pad_mode
.
find
(
"VALID"
)
!=
std
::
string
::
npos
)
...
...
test/gpu/miopen.cpp
View file @
a693711d
...
...
@@ -1581,7 +1581,7 @@ void pad_test()
auto
l0
=
p
.
add_literal
(
migraphx
::
literal
{
s0
,
data0
});
migraphx
::
op
::
pad
op
{};
op
.
value
=
std
::
numeric_limits
<
int8_t
>::
lowest
();
op
.
pads
=
{
0
,
0
,
1
,
1
};
op
.
pads
=
{
0
,
0
,
1
,
1
};
p
.
add_instruction
(
op
,
l0
);
p
.
compile
(
migraphx
::
gpu
::
target
{});
migraphx
::
program
::
parameter_map
m
;
...
...
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