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
789f86fb
Commit
789f86fb
authored
Apr 04, 2022
by
Shucai Xiao
Browse files
clang format
parent
8e485cc8
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
6 additions
and
5 deletions
+6
-5
src/targets/gpu/device/layernorm.cpp
src/targets/gpu/device/layernorm.cpp
+6
-5
No files found.
src/targets/gpu/device/layernorm.cpp
View file @
789f86fb
...
@@ -358,16 +358,17 @@ void triadd_layernorm(hipStream_t stream,
...
@@ -358,16 +358,17 @@ void triadd_layernorm(hipStream_t stream,
auto
batch_item_num
=
in_s
.
lens
().
back
();
auto
batch_item_num
=
in_s
.
lens
().
back
();
if
(
type
==
shape
::
half_type
and
(
batch_item_num
%
2
)
==
0
)
if
(
type
==
shape
::
half_type
and
(
batch_item_num
%
2
)
==
0
)
{
{
auto
block_size
=
compute_block_size
(
batch_item_num
,
1024
);
auto
block_size
=
compute_block_size
(
batch_item_num
,
1024
);
int
block_num
=
in_s
.
elements
()
/
batch_item_num
;
int
block_num
=
in_s
.
elements
()
/
batch_item_num
;
int
shared_size
=
batch_item_num
*
2
*
in_s
.
type_size
();
int
shared_size
=
batch_item_num
*
2
*
in_s
.
type_size
();
auto
half2_block_size
=
block_size
/
4
;
auto
half2_block_size
=
block_size
/
4
;
triadd_layernorm_half2
<<<
block_num
,
half2_block_size
,
shared_size
,
stream
>>>
(
triadd_layernorm_half2
<<<
block_num
,
half2_block_size
,
shared_size
,
stream
>>>
(
arg1
.
data
(),
arg2
.
data
(),
arg3
.
data
(),
result
.
data
(),
batch_item_num
,
half2_block_size
);
arg1
.
data
(),
arg2
.
data
(),
arg3
.
data
(),
result
.
data
(),
batch_item_num
,
half2_block_size
);
// auto half_block_size = block_size / 2;
// auto half_block_size = block_size / 2;
// triadd_layernorm_half2<<<block_num, half_block_size, shared_size, stream>>>(
// triadd_layernorm_half2<<<block_num, half_block_size, shared_size, stream>>>(
// arg1.data(), arg2.data(), arg3.data(), result.data(), batch_item_num, half_block_size);
// arg1.data(), arg2.data(), arg3.data(), result.data(), batch_item_num,
// half_block_size);
}
}
else
else
{
{
...
@@ -427,10 +428,10 @@ void layernorm(hipStream_t stream, const argument& result, const argument& arg1)
...
@@ -427,10 +428,10 @@ void layernorm(hipStream_t stream, const argument& result, const argument& arg1)
auto
batch_item_num
=
in_s
.
lens
().
back
();
auto
batch_item_num
=
in_s
.
lens
().
back
();
if
(
type
==
shape
::
half_type
and
(
batch_item_num
%
2
)
==
0
)
if
(
type
==
shape
::
half_type
and
(
batch_item_num
%
2
)
==
0
)
{
{
auto
block_size
=
compute_block_size
(
batch_item_num
,
1024
);
auto
block_size
=
compute_block_size
(
batch_item_num
,
1024
);
int
block_num
=
in_s
.
elements
()
/
batch_item_num
;
int
block_num
=
in_s
.
elements
()
/
batch_item_num
;
int
shared_size
=
batch_item_num
*
2
*
in_s
.
type_size
();
int
shared_size
=
batch_item_num
*
2
*
in_s
.
type_size
();
auto
half2_block_size
=
block_size
/
4
;
auto
half2_block_size
=
block_size
/
4
;
layernorm_half2
<<<
block_num
,
half2_block_size
,
shared_size
,
stream
>>>
(
layernorm_half2
<<<
block_num
,
half2_block_size
,
shared_size
,
stream
>>>
(
arg1
.
data
(),
result
.
data
(),
batch_item_num
,
half2_block_size
);
arg1
.
data
(),
result
.
data
(),
batch_item_num
,
half2_block_size
);
...
...
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