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
composable_kernel
Commits
1af0a0a4
Commit
1af0a0a4
authored
Sep 26, 2023
by
letaoqin
Browse files
change the check for vector load
parent
b19fe02b
Changes
5
Show whitespace changes
Inline
Side-by-side
Showing
5 changed files
with
5 additions
and
11 deletions
+5
-11
example/52_flash_atten_bias/grouped_mutihead_attention_bias_infer.cpp
...lash_atten_bias/grouped_mutihead_attention_bias_infer.cpp
+1
-1
include/ck/tensor_operation/gpu/device/impl/device_batched_mha_fwd_xdl_cshuffle_v2.hpp
...pu/device/impl/device_batched_mha_fwd_xdl_cshuffle_v2.hpp
+1
-2
include/ck/tensor_operation/gpu/device/impl/device_batched_mha_infer_xdl_cshuffle.hpp
...gpu/device/impl/device_batched_mha_infer_xdl_cshuffle.hpp
+1
-2
include/ck/tensor_operation/gpu/device/impl/device_grouped_mha_fwd_xdl_cshuffle_v2.hpp
...pu/device/impl/device_grouped_mha_fwd_xdl_cshuffle_v2.hpp
+1
-3
include/ck/tensor_operation/gpu/device/impl/device_grouped_mha_infer_xdl_cshuffle.hpp
...gpu/device/impl/device_grouped_mha_infer_xdl_cshuffle.hpp
+1
-3
No files found.
example/52_flash_atten_bias/grouped_mutihead_attention_bias_infer.cpp
View file @
1af0a0a4
...
@@ -120,7 +120,7 @@ using DeviceGemmInstance =
...
@@ -120,7 +120,7 @@ using DeviceGemmInstance =
8
,
8
,
8
,
8
,
true
,
true
,
4
,
1
,
S
<
16
,
16
,
1
>
,
// B1BlockTransfer
S
<
16
,
16
,
1
>
,
// B1BlockTransfer
S
<
0
,
2
,
1
>
,
S
<
0
,
2
,
1
>
,
S
<
0
,
2
,
1
>
,
S
<
0
,
2
,
1
>
,
...
...
include/ck/tensor_operation/gpu/device/impl/device_batched_mha_fwd_xdl_cshuffle_v2.hpp
View file @
1af0a0a4
...
@@ -1028,8 +1028,7 @@ struct DeviceBatchedMultiheadAttentionForward_Xdl_CShuffle_V2
...
@@ -1028,8 +1028,7 @@ struct DeviceBatchedMultiheadAttentionForward_Xdl_CShuffle_V2
{
{
if
(
arg
.
d0_n_length_stride_
[
1
]
==
1
)
if
(
arg
.
d0_n_length_stride_
[
1
]
==
1
)
{
{
if
(
!
(
arg
.
d0_n_length_stride_
[
0
]
%
Acc0BiasTransferSrcScalarPerVector
==
0
||
if
(
arg
.
d0_n_length_stride_
[
0
]
%
Acc0BiasTransferSrcScalarPerVector
!=
0
)
Transform
::
matrix_padder
.
PadN
))
return
false
;
return
false
;
}
}
else
if
(
Acc0BiasTransferSrcScalarPerVector
!=
1
)
else
if
(
Acc0BiasTransferSrcScalarPerVector
!=
1
)
...
...
include/ck/tensor_operation/gpu/device/impl/device_batched_mha_infer_xdl_cshuffle.hpp
View file @
1af0a0a4
...
@@ -744,8 +744,7 @@ struct DeviceBatchedMultiheadAttentionInfer_Xdl_CShuffle
...
@@ -744,8 +744,7 @@ struct DeviceBatchedMultiheadAttentionInfer_Xdl_CShuffle
{
{
if
(
arg
.
d0_n_length_stride_
[
1
]
==
1
)
if
(
arg
.
d0_n_length_stride_
[
1
]
==
1
)
{
{
if
(
!
(
arg
.
d0_n_length_stride_
[
0
]
%
Acc0BiasTransferSrcScalarPerVector
==
0
||
if
(
arg
.
d0_n_length_stride_
[
0
]
%
Acc0BiasTransferSrcScalarPerVector
!=
0
)
Transform
::
matrix_padder
.
PadN
))
return
false
;
return
false
;
}
}
else
if
(
Acc0BiasTransferSrcScalarPerVector
!=
1
)
else
if
(
Acc0BiasTransferSrcScalarPerVector
!=
1
)
...
...
include/ck/tensor_operation/gpu/device/impl/device_grouped_mha_fwd_xdl_cshuffle_v2.hpp
View file @
1af0a0a4
...
@@ -1104,9 +1104,7 @@ struct DeviceGroupedMultiheadAttentionForward_Xdl_CShuffle_V2
...
@@ -1104,9 +1104,7 @@ struct DeviceGroupedMultiheadAttentionForward_Xdl_CShuffle_V2
{
{
if
(
device_arg
.
d0_n_length_stride_
[
1
]
==
1
)
if
(
device_arg
.
d0_n_length_stride_
[
1
]
==
1
)
{
{
if
(
!
(
device_arg
.
d0_n_length_stride_
[
0
]
%
Acc0BiasTransferSrcScalarPerVector
==
if
(
device_arg
.
d0_n_length_stride_
[
0
]
%
Acc0BiasTransferSrcScalarPerVector
!=
0
)
0
||
Transform
::
matrix_padder
.
PadN
))
return
false
;
return
false
;
}
}
else
if
(
Acc0BiasTransferSrcScalarPerVector
!=
1
)
else
if
(
Acc0BiasTransferSrcScalarPerVector
!=
1
)
...
...
include/ck/tensor_operation/gpu/device/impl/device_grouped_mha_infer_xdl_cshuffle.hpp
View file @
1af0a0a4
...
@@ -788,9 +788,7 @@ struct DeviceGroupedMultiheadAttentionInfer_Xdl_CShuffle
...
@@ -788,9 +788,7 @@ struct DeviceGroupedMultiheadAttentionInfer_Xdl_CShuffle
if
(
device_arg
.
d0_n_length_stride_
[
1
]
==
1
)
if
(
device_arg
.
d0_n_length_stride_
[
1
]
==
1
)
{
{
if
(
!
(
device_arg
.
d0_n_length_stride_
[
0
]
%
Acc0BiasTransferSrcScalarPerVector
==
if
(
device_arg
.
d0_n_length_stride_
[
0
]
%
Acc0BiasTransferSrcScalarPerVector
!=
0
)
0
||
Transform
::
matrix_padder
.
PadN
))
{
{
return
false
;
return
false
;
}
}
...
...
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