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
15285630
"...targets/git@developer.sourcefind.cn:gaoqiong/migraphx.git" did not exist on "d25493840c375a7cbd6d6e7bbb78b598c76a9f7a"
Commit
15285630
authored
Aug 11, 2022
by
Paul
Browse files
Format
parent
ba72ce42
Changes
2
Hide whitespace changes
Inline
Side-by-side
Showing
2 changed files
with
6 additions
and
6 deletions
+6
-6
src/targets/gpu/kernels/include/migraphx/kernels/layernorm.hpp
...argets/gpu/kernels/include/migraphx/kernels/layernorm.hpp
+2
-2
src/targets/gpu/kernels/include/migraphx/kernels/reduce.hpp
src/targets/gpu/kernels/include/migraphx/kernels/reduce.hpp
+4
-4
No files found.
src/targets/gpu/kernels/include/migraphx/kernels/layernorm.hpp
View file @
15285630
...
@@ -41,9 +41,9 @@ __device__ void generic_binary_layernorm(
...
@@ -41,9 +41,9 @@ __device__ void generic_binary_layernorm(
{
{
using
reduce_output
=
reduce
::
with_axis
<
Input1
,
Axis
>
;
using
reduce_output
=
reduce
::
with_axis
<
Input1
,
Axis
>
;
reduce
::
block
::
run
<
reduce_output
>
([
&
](
auto
,
auto
r
)
{
reduce
::
block
::
run
<
reduce_output
>
([
&
](
auto
,
auto
r
)
{
using
value_type
=
typename
Input1
::
type
;
using
value_type
=
typename
Input1
::
type
;
constexpr
auto
relements
=
r
.
template
elements
<
Input1
>();
constexpr
auto
relements
=
r
.
template
elements
<
Input1
>();
auto
mean
=
[
&
](
auto
f
)
{
auto
mean
=
[
&
](
auto
f
)
{
return
r
.
reduce
(
op
::
sum
{},
0
,
[
&
](
auto
x1
,
auto
x2
)
{
return
r
.
reduce
(
op
::
sum
{},
0
,
[
&
](
auto
x1
,
auto
x2
)
{
return
f
(
x1
,
x2
)
/
value_type
{
relements
};
return
f
(
x1
,
x2
)
/
value_type
{
relements
};
})(
input1
,
input2
);
})(
input1
,
input2
);
...
...
src/targets/gpu/kernels/include/migraphx/kernels/reduce.hpp
View file @
15285630
...
@@ -225,11 +225,11 @@ struct block
...
@@ -225,11 +225,11 @@ struct block
});
});
}
}
template
<
class
Input
>
template
<
class
Input
>
constexpr
auto
elements
()
const
constexpr
auto
elements
()
const
{
{
using
reduce_type
=
decltype
(
slicer
(
Input
{}));
using
reduce_type
=
decltype
(
slicer
(
Input
{}));
using
value_type
=
typename
Input
::
type
;
using
value_type
=
typename
Input
::
type
;
constexpr
auto
relements
=
get_shape_c
<
reduce_type
>
{}.
elements
();
constexpr
auto
relements
=
get_shape_c
<
reduce_type
>
{}.
elements
();
if
constexpr
(
vec_size
<
value_type
>
()
>
1
)
if
constexpr
(
vec_size
<
value_type
>
()
>
1
)
return
relements
*
vec_size
<
value_type
>
();
return
relements
*
vec_size
<
value_type
>
();
...
@@ -294,7 +294,7 @@ struct lane
...
@@ -294,7 +294,7 @@ struct lane
});
});
}
}
template
<
class
Input
>
template
<
class
Input
>
constexpr
auto
elements
()
const
constexpr
auto
elements
()
const
{
{
using
reduce_type
=
decltype
(
slicer
(
Input
{}));
using
reduce_type
=
decltype
(
slicer
(
Input
{}));
...
...
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