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
b092d017
Commit
b092d017
authored
Mar 27, 2019
by
Paul
Browse files
Formatting
parent
eed607a3
Changes
2
Hide whitespace changes
Inline
Side-by-side
Showing
2 changed files
with
18 additions
and
16 deletions
+18
-16
src/targets/gpu/device/concat.cpp
src/targets/gpu/device/concat.cpp
+11
-6
src/targets/gpu/device/include/migraphx/gpu/device/tensor.hpp
...targets/gpu/device/include/migraphx/gpu/device/tensor.hpp
+7
-10
No files found.
src/targets/gpu/device/concat.cpp
View file @
b092d017
...
@@ -16,18 +16,23 @@ argument concat(hipStream_t stream,
...
@@ -16,18 +16,23 @@ argument concat(hipStream_t stream,
std
::
vector
<
std
::
size_t
>
offsets_vec
)
std
::
vector
<
std
::
size_t
>
offsets_vec
)
{
{
static
constexpr
const
std
::
size_t
limit
=
6
;
static
constexpr
const
std
::
size_t
limit
=
6
;
if
(
offsets_vec
.
size
()
>
limit
)
if
(
offsets_vec
.
size
()
>
limit
)
MIGRAPHX_THROW
(
"Too many arguments to concat"
);
MIGRAPHX_THROW
(
"Too many arguments to concat"
);
std
::
size_t
nelements
=
std
::
max_element
(
args_vec
.
begin
(),
std
::
prev
(
args_vec
.
end
()),
by
(
std
::
less
<>
{},
[
&
](
auto
&&
x
)
{
return
x
.
get_shape
().
elements
();
}))
->
get_shape
().
elements
();
std
::
size_t
nelements
=
std
::
max_element
(
args_vec
.
begin
(),
std
::
prev
(
args_vec
.
end
()),
by
(
std
::
less
<>
{},
[
&
](
auto
&&
x
)
{
return
x
.
get_shape
().
elements
();
}))
->
get_shape
()
.
elements
();
auto
offsets
=
to_hip_vector
<
limit
>
(
offsets_vec
);
auto
offsets
=
to_hip_vector
<
limit
>
(
offsets_vec
);
hip_visit_all
<
limit
+
1
>
(
args_vec
)([
&
](
auto
args
)
{
hip_visit_all
<
limit
+
1
>
(
args_vec
)([
&
](
auto
args
)
{
auto
output
=
args
.
back
();
auto
output
=
args
.
back
();
auto
ninputs
=
args
.
size
()
-
1
;
auto
ninputs
=
args
.
size
()
-
1
;
gs_launch
(
stream
,
nelements
)([
=
](
auto
i
)
{
gs_launch
(
stream
,
nelements
)([
=
](
auto
i
)
{
for
(
std
::
size_t
j
=
0
;
j
<
ninputs
;
j
++
)
for
(
std
::
size_t
j
=
0
;
j
<
ninputs
;
j
++
)
{
{
auto
&&
arg
=
args
[
j
];
auto
&&
arg
=
args
[
j
];
if
(
i
>=
arg
.
size
())
if
(
i
>=
arg
.
size
())
continue
;
continue
;
auto
idx
=
output
.
get_shape
().
index
(
arg
.
get_shape
().
multi
(
i
));
auto
idx
=
output
.
get_shape
().
index
(
arg
.
get_shape
().
multi
(
i
));
output
.
data
()[
idx
+
offsets
[
j
]]
=
arg
.
data
()[
i
];
output
.
data
()[
idx
+
offsets
[
j
]]
=
arg
.
data
()[
i
];
...
...
src/targets/gpu/device/include/migraphx/gpu/device/tensor.hpp
View file @
b092d017
...
@@ -124,21 +124,18 @@ hip_vector<T, N> to_hip_vector(const std::vector<T>& x)
...
@@ -124,21 +124,18 @@ hip_vector<T, N> to_hip_vector(const std::vector<T>& x)
return
result
;
return
result
;
}
}
template
<
std
::
size_t
N
>
template
<
std
::
size_t
N
>
struct
hip_shape
struct
hip_shape
{
{
using
hip_index
=
hip_array
<
std
::
size_t
,
N
>
;
using
hip_index
=
hip_array
<
std
::
size_t
,
N
>
;
hip_array
<
std
::
size_t
,
N
>
lens
=
{};
hip_array
<
std
::
size_t
,
N
>
lens
=
{};
hip_array
<
std
::
size_t
,
N
>
strides
=
{};
hip_array
<
std
::
size_t
,
N
>
strides
=
{};
std
::
size_t
elements
=
0
;
std
::
size_t
elements
=
0
;
bool
standard
=
false
;
bool
standard
=
false
;
__device__
__host__
hip_shape
()
=
default
;
__device__
__host__
hip_shape
()
=
default
;
hip_shape
(
const
shape
&
s
)
hip_shape
(
const
shape
&
s
)
:
elements
(
s
.
elements
()),
standard
(
s
.
standard
())
:
elements
(
s
.
elements
()),
standard
(
s
.
standard
())
{
{
assert
(
s
.
lens
().
size
()
==
N
);
assert
(
s
.
lens
().
size
()
==
N
);
assert
(
s
.
strides
().
size
()
==
N
);
assert
(
s
.
strides
().
size
()
==
N
);
...
@@ -240,7 +237,7 @@ auto hip_visit_all(T&& x, Ts&&... xs)
...
@@ -240,7 +237,7 @@ auto hip_visit_all(T&& x, Ts&&... xs)
{
{
return
[
&
](
auto
f
)
{
return
[
&
](
auto
f
)
{
visit_tensor_size
(
x
.
get_shape
().
lens
().
size
(),
[
&
](
auto
dim
)
{
visit_tensor_size
(
x
.
get_shape
().
lens
().
size
(),
[
&
](
auto
dim
)
{
visit_all
(
x
,
xs
...)([
&
](
auto
...
vs
)
{
f
(
make_hip_tensor_view
<
dim
>
(
vs
)...);
});
visit_all
(
x
,
xs
...)([
&
](
auto
...
vs
)
{
f
(
make_hip_tensor_view
<
dim
>
(
vs
)...);
});
});
});
};
};
}
}
...
@@ -250,7 +247,7 @@ auto hip_visit_all(const std::vector<T>& x)
...
@@ -250,7 +247,7 @@ auto hip_visit_all(const std::vector<T>& x)
{
{
return
[
&
](
auto
f
)
{
return
[
&
](
auto
f
)
{
visit_tensor_size
(
x
.
front
().
get_shape
().
lens
().
size
(),
[
&
](
auto
dim
)
{
visit_tensor_size
(
x
.
front
().
get_shape
().
lens
().
size
(),
[
&
](
auto
dim
)
{
visit_all
(
x
)([
&
](
auto
&&
v
)
{
f
(
make_hip_tensor_views
<
dim
,
N
>
(
v
));
});
visit_all
(
x
)([
&
](
auto
&&
v
)
{
f
(
make_hip_tensor_views
<
dim
,
N
>
(
v
));
});
});
});
};
};
}
}
...
...
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