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
b3e2901e
Commit
b3e2901e
authored
Oct 26, 2018
by
Paul
Browse files
Formatting
parent
9979fada
Changes
9
Hide whitespace changes
Inline
Side-by-side
Showing
9 changed files
with
59 additions
and
26 deletions
+59
-26
src/targets/gpu/concat.cpp
src/targets/gpu/concat.cpp
+3
-2
src/targets/gpu/contiguous.cpp
src/targets/gpu/contiguous.cpp
+3
-2
src/targets/gpu/device/add.cpp
src/targets/gpu/device/add.cpp
+5
-1
src/targets/gpu/device/add_relu.cpp
src/targets/gpu/device/add_relu.cpp
+8
-3
src/targets/gpu/device/concat.cpp
src/targets/gpu/device/concat.cpp
+2
-1
src/targets/gpu/device/include/migraph/gpu/device/nary.hpp
src/targets/gpu/device/include/migraph/gpu/device/nary.hpp
+23
-12
src/targets/gpu/include/migraph/gpu/device/add.hpp
src/targets/gpu/include/migraph/gpu/device/add.hpp
+5
-1
src/targets/gpu/include/migraph/gpu/device/add_relu.hpp
src/targets/gpu/include/migraph/gpu/device/add_relu.hpp
+6
-2
src/targets/gpu/include/migraph/gpu/device/concat.hpp
src/targets/gpu/include/migraph/gpu/device/concat.hpp
+4
-2
No files found.
src/targets/gpu/concat.cpp
View file @
b3e2901e
...
...
@@ -14,8 +14,9 @@ shape hip_concat::compute_shape(std::vector<shape> inputs) const
return
op
.
compute_shape
(
inputs
);
}
argument
hip_concat
::
compute
(
context
&
ctx
,
const
shape
&
output_shape
,
const
std
::
vector
<
argument
>&
args
)
const
argument
hip_concat
::
compute
(
context
&
ctx
,
const
shape
&
output_shape
,
const
std
::
vector
<
argument
>&
args
)
const
{
std
::
vector
<
std
::
size_t
>
offsets
=
op
.
compute_offsets
(
output_shape
,
args
);
return
device
::
concat
(
ctx
.
get_stream
().
get
(),
output_shape
,
args
,
offsets
);
...
...
src/targets/gpu/contiguous.cpp
View file @
b3e2901e
...
...
@@ -12,8 +12,9 @@ shape miopen_contiguous::compute_shape(const std::vector<shape>& inputs) const
check_shapes
{
inputs
,
*
this
}.
has
(
2
);
return
op
.
compute_shape
({
inputs
.
at
(
0
)});
}
argument
miopen_contiguous
::
compute
(
context
&
ctx
,
shape
output_shape
,
const
std
::
vector
<
argument
>&
args
)
const
argument
miopen_contiguous
::
compute
(
context
&
ctx
,
shape
output_shape
,
const
std
::
vector
<
argument
>&
args
)
const
{
assert
(
output_shape
==
args
[
1
].
get_shape
());
assert
(
output_shape
.
standard
());
...
...
src/targets/gpu/device/add.cpp
View file @
b3e2901e
...
...
@@ -10,7 +10,11 @@ void add(hipStream_t stream, const argument& result, const argument& arg1, const
nary
(
stream
,
result
,
arg1
,
arg2
)([](
auto
x
,
auto
y
)
{
return
x
+
y
;
});
}
void
add
(
hipStream_t
stream
,
const
argument
&
result
,
const
argument
&
arg1
,
const
argument
&
arg2
,
const
argument
&
arg3
)
void
add
(
hipStream_t
stream
,
const
argument
&
result
,
const
argument
&
arg1
,
const
argument
&
arg2
,
const
argument
&
arg3
)
{
nary
(
stream
,
result
,
arg1
,
arg2
,
arg3
)([](
auto
x
,
auto
y
,
auto
z
)
{
return
x
+
y
+
z
;
});
}
...
...
src/targets/gpu/device/add_relu.cpp
View file @
b3e2901e
...
...
@@ -5,12 +5,17 @@ namespace migraph {
namespace
gpu
{
namespace
device
{
void
add_relu
(
hipStream_t
stream
,
const
argument
&
result
,
const
argument
&
arg1
,
const
argument
&
arg2
)
void
add_relu
(
hipStream_t
stream
,
const
argument
&
result
,
const
argument
&
arg1
,
const
argument
&
arg2
)
{
nary
(
stream
,
result
,
arg1
,
arg2
)([](
auto
x
,
auto
y
)
{
return
std
::
max
<
decltype
(
x
+
y
)
>
(
0
,
x
+
y
);
});
nary
(
stream
,
result
,
arg1
,
arg2
)(
[](
auto
x
,
auto
y
)
{
return
std
::
max
<
decltype
(
x
+
y
)
>
(
0
,
x
+
y
);
});
}
void
add_relu
(
hipStream_t
stream
,
const
argument
&
result
,
void
add_relu
(
hipStream_t
stream
,
const
argument
&
result
,
const
argument
&
arg1
,
const
argument
&
arg2
,
const
argument
&
arg3
)
...
...
src/targets/gpu/device/concat.cpp
View file @
b3e2901e
...
...
@@ -8,7 +8,8 @@ namespace migraph {
namespace
gpu
{
namespace
device
{
argument
concat
(
hipStream_t
stream
,
const
migraph
::
shape
&
output_shape
,
argument
concat
(
hipStream_t
stream
,
const
migraph
::
shape
&
output_shape
,
std
::
vector
<
migraph
::
argument
>
args
,
std
::
vector
<
std
::
size_t
>
offsets
)
{
...
...
src/targets/gpu/device/include/migraph/gpu/device/nary.hpp
View file @
b3e2901e
...
...
@@ -52,8 +52,12 @@ auto nary_nonstandard_impl(hipStream_t stream, F f, argument result, Arguments..
}
template
<
class
F
>
void
trinary_broadcast_vec_impl
(
hipStream_t
stream
,
F
f
,
const
argument
&
result
,
const
argument
&
arg1
,
const
argument
&
arg2
,
const
argument
&
arg3
)
void
trinary_broadcast_vec_impl
(
hipStream_t
stream
,
F
f
,
const
argument
&
result
,
const
argument
&
arg1
,
const
argument
&
arg2
,
const
argument
&
arg3
)
{
const
auto
&
output_shape
=
result
.
get_shape
();
const
auto
&
b_shape
=
arg3
.
get_shape
();
...
...
@@ -107,8 +111,12 @@ void trinary_broadcast_vec_impl(
}
template
<
class
F
>
void
trinary_broadcast_impl
(
hipStream_t
stream
,
F
f
,
const
argument
&
result
,
const
argument
&
arg1
,
const
argument
&
arg2
,
const
argument
&
arg3
)
void
trinary_broadcast_impl
(
hipStream_t
stream
,
F
f
,
const
argument
&
result
,
const
argument
&
arg1
,
const
argument
&
arg2
,
const
argument
&
arg3
)
{
const
auto
&
output_shape
=
result
.
get_shape
();
const
auto
&
b_shape
=
arg3
.
get_shape
();
...
...
@@ -154,10 +162,8 @@ void trinary_broadcast_impl(
}
template
<
class
F
>
void
binary_broadcast_vec_impl
(
hipStream_t
stream
,
F
f
,
const
argument
&
result
,
const
argument
&
arg1
,
const
argument
&
arg2
)
void
binary_broadcast_vec_impl
(
hipStream_t
stream
,
F
f
,
const
argument
&
result
,
const
argument
&
arg1
,
const
argument
&
arg2
)
{
const
auto
&
output_shape
=
result
.
get_shape
();
const
auto
&
b_shape
=
arg2
.
get_shape
();
...
...
@@ -209,7 +215,8 @@ void binary_broadcast_vec_impl(hipStream_t stream, F f,
}
template
<
class
F
>
void
binary_broadcast_impl
(
hipStream_t
stream
,
F
f
,
const
argument
&
result
,
const
argument
&
arg1
,
const
argument
&
arg2
)
void
binary_broadcast_impl
(
hipStream_t
stream
,
F
f
,
const
argument
&
result
,
const
argument
&
arg1
,
const
argument
&
arg2
)
{
const
auto
&
output_shape
=
result
.
get_shape
();
const
auto
&
b_shape
=
arg2
.
get_shape
();
...
...
@@ -321,7 +328,8 @@ auto nary(hipStream_t stream, argument result, Arguments... args)
return
[
=
](
auto
f
)
{
nary_impl
(
stream
,
f
,
result
,
args
...);
};
}
inline
auto
nary
(
hipStream_t
stream
,
const
argument
&
result
,
const
argument
&
arg1
,
const
argument
&
arg2
)
inline
auto
nary
(
hipStream_t
stream
,
const
argument
&
result
,
const
argument
&
arg1
,
const
argument
&
arg2
)
{
return
[
=
](
auto
f
)
{
// TODO: Check result and arg1 shape is the same
...
...
@@ -349,8 +357,11 @@ inline auto nary(hipStream_t stream, const argument& result, const argument& arg
};
}
inline
auto
nary
(
hipStream_t
stream
,
const
argument
&
result
,
const
argument
&
arg1
,
const
argument
&
arg2
,
const
argument
&
arg3
)
inline
auto
nary
(
hipStream_t
stream
,
const
argument
&
result
,
const
argument
&
arg1
,
const
argument
&
arg2
,
const
argument
&
arg3
)
{
return
[
=
](
auto
f
)
{
// TODO: Check result and arg1 shape is the same
...
...
src/targets/gpu/include/migraph/gpu/device/add.hpp
View file @
b3e2901e
...
...
@@ -11,7 +11,11 @@ namespace device {
void
add
(
hipStream_t
stream
,
const
argument
&
result
,
const
argument
&
arg1
,
const
argument
&
arg2
);
void
add
(
hipStream_t
stream
,
const
argument
&
result
,
const
argument
&
arg1
,
const
argument
&
arg2
,
const
argument
&
arg3
);
void
add
(
hipStream_t
stream
,
const
argument
&
result
,
const
argument
&
arg1
,
const
argument
&
arg2
,
const
argument
&
arg3
);
}
// namespace device
}
// namespace gpu
...
...
src/targets/gpu/include/migraph/gpu/device/add_relu.hpp
View file @
b3e2901e
...
...
@@ -9,9 +9,13 @@ namespace migraph {
namespace
gpu
{
namespace
device
{
void
add_relu
(
hipStream_t
stream
,
const
argument
&
result
,
const
argument
&
arg1
,
const
argument
&
arg2
);
void
add_relu
(
hipStream_t
stream
,
const
argument
&
result
,
const
argument
&
arg1
,
const
argument
&
arg2
);
void
add_relu
(
hipStream_t
stream
,
const
argument
&
result
,
void
add_relu
(
hipStream_t
stream
,
const
argument
&
result
,
const
argument
&
arg1
,
const
argument
&
arg2
,
const
argument
&
arg3
);
...
...
src/targets/gpu/include/migraph/gpu/device/concat.hpp
View file @
b3e2901e
...
...
@@ -8,8 +8,10 @@ namespace migraph {
namespace
gpu
{
namespace
device
{
argument
concat
(
hipStream_t
stream
,
const
shape
&
output_shape
,
std
::
vector
<
argument
>
args
,
std
::
vector
<
std
::
size_t
>
offsets
);
argument
concat
(
hipStream_t
stream
,
const
shape
&
output_shape
,
std
::
vector
<
argument
>
args
,
std
::
vector
<
std
::
size_t
>
offsets
);
}
// namespace device
}
// namespace gpu
...
...
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