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
c6700632
"vscode:/vscode.git/clone" did not exist on "967133c11ef03207e758b8b3f88f6791a4dd6b6d"
Commit
c6700632
authored
Mar 28, 2022
by
Shucai Xiao
Browse files
clang format
parent
69c94135
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
2 additions
and
3 deletions
+2
-3
src/targets/gpu/device/layernorm.cpp
src/targets/gpu/device/layernorm.cpp
+2
-3
No files found.
src/targets/gpu/device/layernorm.cpp
View file @
c6700632
...
@@ -237,7 +237,7 @@ __global__ void triadd_layernorm_kernel_half2(
...
@@ -237,7 +237,7 @@ __global__ void triadd_layernorm_kernel_half2(
__half2
*
input2
=
reinterpret_cast
<
__half2
*>
(
in2
);
__half2
*
input2
=
reinterpret_cast
<
__half2
*>
(
in2
);
__half2
*
input3
=
reinterpret_cast
<
__half2
*>
(
in3
);
__half2
*
input3
=
reinterpret_cast
<
__half2
*>
(
in3
);
__half2
*
output
=
reinterpret_cast
<
__half2
*>
(
data_out
);
__half2
*
output
=
reinterpret_cast
<
__half2
*>
(
data_out
);
auto
rnum
=
__float2half2_rn
(
1.0
f
/
batch_item_num
);
auto
rnum
=
__float2half2_rn
(
1.0
f
/
batch_item_num
);
batch_item_num
/=
2
;
batch_item_num
/=
2
;
extern
MIGRAPHX_DEVICE_SHARED
__half2
buffer2
[];
extern
MIGRAPHX_DEVICE_SHARED
__half2
buffer2
[];
__half2
*
in_data_reduce
=
buffer2
;
__half2
*
in_data_reduce
=
buffer2
;
...
@@ -410,8 +410,7 @@ void triadd_layernorm(hipStream_t stream,
...
@@ -410,8 +410,7 @@ void triadd_layernorm(hipStream_t stream,
int
shared_size
=
batch_item_num
*
2
*
in_s
.
type_size
();
int
shared_size
=
batch_item_num
*
2
*
in_s
.
type_size
();
half2_block_size
=
half2_block_size
/
4
;
half2_block_size
=
half2_block_size
/
4
;
triadd_layernorm_kernel_half2
<<<
block_num
,
half2_block_size
,
shared_size
,
stream
>>>
(
triadd_layernorm_kernel_half2
<<<
block_num
,
half2_block_size
,
shared_size
,
stream
>>>
(
arg1
.
data
(),
arg2
.
data
(),
arg3
.
data
(),
result
.
data
(),
batch_item_num
,
arg1
.
data
(),
arg2
.
data
(),
arg3
.
data
(),
result
.
data
(),
batch_item_num
,
half2_block_size
);
half2_block_size
);
}
}
// if(type == shape::half_type)
// if(type == shape::half_type)
// {
// {
...
...
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