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
3de56715
Commit
3de56715
authored
Oct 18, 2018
by
wsttiger
Browse files
Formatting
parent
447eec93
Changes
5
Hide whitespace changes
Inline
Side-by-side
Showing
5 changed files
with
20 additions
and
19 deletions
+20
-19
src/targets/gpu/concat.cpp
src/targets/gpu/concat.cpp
+4
-4
src/targets/gpu/device/concat.cpp
src/targets/gpu/device/concat.cpp
+10
-10
src/targets/gpu/include/migraph/gpu/device/concat.hpp
src/targets/gpu/include/migraph/gpu/device/concat.hpp
+2
-1
src/targets/gpu/lowering.cpp
src/targets/gpu/lowering.cpp
+2
-2
test/gpu/miopen.cpp
test/gpu/miopen.cpp
+2
-2
No files found.
src/targets/gpu/concat.cpp
View file @
3de56715
...
...
@@ -8,14 +8,14 @@
namespace
migraph
{
namespace
gpu
{
shape
hip_concat
::
compute_shape
(
std
::
vector
<
shape
>
inputs
)
const
shape
hip_concat
::
compute_shape
(
std
::
vector
<
shape
>
inputs
)
const
{
inputs
.
pop_back
();
return
op
.
compute_shape
(
inputs
);
}
std
::
vector
<
std
::
size_t
>
hip_concat
::
compute_offsets
(
const
shape
&
output_shape
,
const
std
::
vector
<
argument
>
args
)
const
const
std
::
vector
<
argument
>
args
)
const
{
std
::
vector
<
std
::
size_t
>
offsets
;
std
::
vector
<
std
::
size_t
>
offset
(
args
[
0
].
get_shape
().
lens
().
size
(),
0
);
...
...
@@ -29,8 +29,8 @@ std::vector<std::size_t> hip_concat::compute_offsets(const shape& output_shape,
}
argument
hip_concat
::
compute
(
context
&
ctx
,
const
shape
&
output_shape
,
const
std
::
vector
<
argument
>&
args
)
const
const
shape
&
output_shape
,
const
std
::
vector
<
argument
>&
args
)
const
{
std
::
vector
<
std
::
size_t
>
offsets
=
compute_offsets
(
output_shape
,
args
);
return
device
::
concat
(
output_shape
,
args
,
offsets
);
...
...
src/targets/gpu/device/concat.cpp
View file @
3de56715
...
...
@@ -8,30 +8,30 @@ namespace migraph {
namespace
gpu
{
namespace
device
{
argument
concat
(
const
migraph
::
shape
&
output_shape
,
std
::
vector
<
migraph
::
argument
>
args
,
std
::
vector
<
std
::
size_t
>
offsets
)
argument
concat
(
const
migraph
::
shape
&
output_shape
,
std
::
vector
<
migraph
::
argument
>
args
,
std
::
vector
<
std
::
size_t
>
offsets
)
{
//migraph::argument& result = args.back();
for
(
std
::
size_t
l
=
0
;
l
<
args
.
size
()
-
1
;
l
++
)
//
migraph::argument& result = args.back();
for
(
std
::
size_t
l
=
0
;
l
<
args
.
size
()
-
1
;
l
++
)
{
auto
argl
=
args
[
l
];
auto
argl
=
args
[
l
];
std
::
size_t
nelements
=
argl
.
get_shape
().
elements
();
visit_all
(
args
.
back
(),
argl
)([
&
](
auto
output
,
auto
input
)
{
visit_tensor_size
(
output_shape
.
lens
().
size
(),
[
&
](
auto
ndim
)
{
visit_tensor_size
(
output_shape
.
lens
().
size
(),
[
&
](
auto
ndim
)
{
auto
*
outptr
=
output
.
data
()
+
offsets
[
l
];
const
auto
*
inptr
=
input
.
data
();
hip_tensor_descriptor
<
ndim
>
desc_input
(
input
.
get_shape
());
hip_tensor_descriptor
<
ndim
>
desc_output
(
output
.
get_shape
());
gs_launch
(
nelements
)([
=
](
auto
i
)
{
outptr
[
desc_output
.
linear
(
desc_input
.
multi
(
i
))]
=
inptr
[
i
];
});
gs_launch
(
nelements
)(
[
=
](
auto
i
)
{
outptr
[
desc_output
.
linear
(
desc_input
.
multi
(
i
))]
=
inptr
[
i
];
});
});
});
}
//return result;
//
return result;
return
args
.
back
();
}
}
// namespace device
}
// namespace gpu
}
// namespace migraph
src/targets/gpu/include/migraph/gpu/device/concat.hpp
View file @
3de56715
...
...
@@ -5,7 +5,8 @@ namespace migraph {
namespace
gpu
{
namespace
device
{
argument
concat
(
const
shape
&
output_shape
,
std
::
vector
<
argument
>
args
,
std
::
vector
<
std
::
size_t
>
offsets
);
argument
concat
(
const
shape
&
output_shape
,
std
::
vector
<
argument
>
args
,
std
::
vector
<
std
::
size_t
>
offsets
);
}
// namespace device
}
// namespace gpu
...
...
src/targets/gpu/lowering.cpp
View file @
3de56715
...
...
@@ -165,8 +165,8 @@ struct miopen_apply
instruction_ref
apply_concat
(
instruction_ref
ins
)
{
auto
&&
op
=
any_cast
<
op
::
concat
>
(
ins
->
get_operator
());
auto
output
=
insert_allocation
(
ins
,
ins
->
get_shape
());
auto
&&
op
=
any_cast
<
op
::
concat
>
(
ins
->
get_operator
());
auto
output
=
insert_allocation
(
ins
,
ins
->
get_shape
());
std
::
vector
<
instruction_ref
>
refs
=
ins
->
inputs
();
refs
.
push_back
(
output
);
return
prog
->
replace_instruction
(
ins
,
hip_concat
{
op
},
refs
);
...
...
test/gpu/miopen.cpp
View file @
3de56715
...
...
@@ -523,7 +523,7 @@ struct test_concat
migraph
::
program
create_program
()
const
{
migraph
::
program
p
;
std
::
size_t
axis
=
1
;
std
::
size_t
axis
=
1
;
migraph
::
shape
s0
{
migraph
::
shape
::
int32_type
,
{
2
,
2
}};
migraph
::
shape
s1
{
migraph
::
shape
::
int32_type
,
{
2
,
3
}};
migraph
::
shape
s2
{
migraph
::
shape
::
int32_type
,
{
2
,
1
}};
...
...
@@ -540,7 +540,7 @@ struct test_concat2
migraph
::
program
create_program
()
const
{
migraph
::
program
p
;
std
::
size_t
axis
=
0
;
std
::
size_t
axis
=
0
;
migraph
::
shape
s0
{
migraph
::
shape
::
int32_type
,
{
2
,
2
}};
migraph
::
shape
s1
{
migraph
::
shape
::
int32_type
,
{
3
,
2
}};
migraph
::
shape
s2
{
migraph
::
shape
::
int32_type
,
{
1
,
2
}};
...
...
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