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
6570ef7a
Commit
6570ef7a
authored
May 10, 2023
by
Po-Yen, Chen
Browse files
Move check for karg.K into CheckValidity()
parent
b0e02b8a
Changes
1
Show whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
19 additions
and
8 deletions
+19
-8
include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v1.hpp
...nsor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v1.hpp
+19
-8
No files found.
include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v1.hpp
View file @
6570ef7a
...
@@ -137,14 +137,10 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdl_cshuffle_v1
...
@@ -137,14 +137,10 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdl_cshuffle_v1
GemmSpec
==
GemmSpecialization
::
KPadding
||
GemmSpec
==
GemmSpecialization
::
KPadding
||
GemmSpec
==
GemmSpecialization
::
NKPadding
)
GemmSpec
==
GemmSpecialization
::
NKPadding
)
{
{
assert
(
CalculateKPadded
(
K
)
%
AK1Value
==
0
);
return
CalculateKPadded
(
K
)
/
AK1Value
;
return
CalculateKPadded
(
K
)
/
AK1Value
;
}
}
else
else
{
{
assert
(
K
%
AK1Value
==
0
);
return
K
/
AK1Value
;
return
K
/
AK1Value
;
}
}
}
}
...
@@ -158,14 +154,10 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdl_cshuffle_v1
...
@@ -158,14 +154,10 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdl_cshuffle_v1
GemmSpec
==
GemmSpecialization
::
KPadding
||
GemmSpec
==
GemmSpecialization
::
KPadding
||
GemmSpec
==
GemmSpecialization
::
MKPadding
)
GemmSpec
==
GemmSpecialization
::
MKPadding
)
{
{
assert
(
CalculateKPadded
(
K
)
%
BK1Value
==
0
);
return
CalculateKPadded
(
K
)
/
BK1Value
;
return
CalculateKPadded
(
K
)
/
BK1Value
;
}
}
else
else
{
{
assert
(
K
%
BK1Value
==
0
);
return
K
/
BK1Value
;
return
K
/
BK1Value
;
}
}
}
}
...
@@ -530,6 +522,25 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdl_cshuffle_v1
...
@@ -530,6 +522,25 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdl_cshuffle_v1
}
}
}
}
if
constexpr
(
GemmSpec
==
tensor_operation
::
device
::
GemmSpecialization
::
MKPadding
||
GemmSpec
==
tensor_operation
::
device
::
GemmSpecialization
::
MNKPadding
||
GemmSpec
==
tensor_operation
::
device
::
GemmSpecialization
::
KPadding
||
GemmSpec
==
tensor_operation
::
device
::
GemmSpecialization
::
NKPadding
)
{
if
(
!
(
CalculateKPadded
(
karg
.
K
)
%
AK1Value
==
0
)
||
!
(
CalculateKPadded
(
karg
.
K
)
%
BK1Value
==
0
))
{
return
false
;
}
}
else
{
if
(
!
(
karg
.
K
%
AK1Value
==
0
)
||
!
(
karg
.
K
%
BK1Value
==
0
))
{
return
false
;
}
}
if
constexpr
(
is_same
<
tensor_layout
::
gemm
::
RowMajor
,
ALayout
>::
value
)
if
constexpr
(
is_same
<
tensor_layout
::
gemm
::
RowMajor
,
ALayout
>::
value
)
{
{
if
(
karg
.
K
%
ABlockTransferSrcScalarPerVector
!=
0
)
if
(
karg
.
K
%
ABlockTransferSrcScalarPerVector
!=
0
)
...
...
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