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
7fa93728
"docs/vscode:/vscode.git/clone" did not exist on "6f1d014463496e1f1066e6fa2ffbcf5a537d489c"
Commit
7fa93728
authored
Nov 10, 2023
by
Paul
Browse files
Add concat kernel for input fusions
parent
3c160a3f
Changes
1
Show whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
17 additions
and
0 deletions
+17
-0
src/targets/gpu/kernels/include/migraphx/kernels/concat.hpp
src/targets/gpu/kernels/include/migraphx/kernels/concat.hpp
+17
-0
No files found.
src/targets/gpu/kernels/include/migraphx/kernels/concat.hpp
View file @
7fa93728
...
@@ -74,5 +74,22 @@ __device__ auto concat(Inputs... inputs)
...
@@ -74,5 +74,22 @@ __device__ auto concat(Inputs... inputs)
};
};
}
}
template
<
index_int
Axis
,
class
...
InputPacks
>
__device__
auto
concat2
(
InputPacks
...
input_packs
)
{
return
[
=
](
auto
f
,
auto
...
ts
)
{
auto
idx
=
make_index
();
fold
([
&
](
auto
start
,
auto
input_pack
)
{
return
input_pack
([
&
](
auto
g
,
auto
x
,
auto
...
xs
)
{
concat_slices
<
Axis
>
(
x
,
start
,
ts
...)([
&
](
auto
z
,
auto
...
ys
)
{
idx
.
global_stride
(
x
.
get_shape
().
elements
(),
[
&
](
auto
i
)
{
z
[
i
]
=
f
(
g
(
x
[
i
],
xs
[
i
]...),
ys
[
i
]...);
});
return
start
+
concat_ends
<
Axis
>
(
x
);
});
});
})(
_c
<
0
>
,
input_packs
...);
}
}
// namespace migraphx
}
// namespace migraphx
#endif // MIGRAPHX_GUARD_KERNELS_CONCAT_HPP
#endif // MIGRAPHX_GUARD_KERNELS_CONCAT_HPP
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