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
910a26b4
Commit
910a26b4
authored
Sep 12, 2022
by
Po-Yen, Chen
Browse files
Check scalar-per-vector with padded length
parent
f3b3a61c
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
17 additions
and
3 deletions
+17
-3
include/ck/tensor_operation/gpu/device/device_permute.hpp
include/ck/tensor_operation/gpu/device/device_permute.hpp
+17
-3
No files found.
include/ck/tensor_operation/gpu/device/device_permute.hpp
View file @
910a26b4
...
...
@@ -232,13 +232,17 @@ struct DevicePermute
static
bool
IsSupportedArgument
(
const
Argument
&
arg
)
{
constexpr
auto
GetPaddedLength
=
[](
index_t
length
,
index_t
tile_length
)
{
return
math
::
integer_divide_ceil
(
length
,
tile_length
)
*
tile_length
;
};
constexpr
auto
IsScalarPerVectorValid
=
[](
index_t
length
,
index_t
stride
,
index_t
scalar
PerV
ector
)
{
if
(
stride
==
1
&&
length
%
scalar
PerV
ector
==
0
)
[](
index_t
length
,
index_t
stride
,
index_t
scalar
_per_v
ector
)
{
if
(
stride
==
1
&&
length
%
scalar
_per_v
ector
==
0
)
{
return
true
;
}
else
if
(
stride
!=
1
&&
scalar
PerV
ector
==
1
)
else
if
(
stride
!=
1
&&
scalar
_per_v
ector
==
1
)
{
return
true
;
}
...
...
@@ -249,9 +253,19 @@ struct DevicePermute
return
IsScalarPerVectorValid
(
arg
.
inLengths_
[
SrcVectorDim
],
arg
.
inStrides_
[
SrcVectorDim
],
SrcScalarPerVector
)
&&
IsScalarPerVectorValid
(
GetPaddedLength
(
arg
.
inLengths_
[
SrcVectorDim
],
(
SrcVectorDim
==
NumDim
-
2
?
HPerBlock
:
WPerBlock
)),
arg
.
inStrides_
[
SrcVectorDim
],
SrcScalarPerVector
)
&&
IsScalarPerVectorValid
(
arg
.
outLengths_
[
DstVectorDim
],
arg
.
outStrides_
[
DstVectorDim
],
DstScalarPerVector
)
&&
IsScalarPerVectorValid
(
GetPaddedLength
(
arg
.
outLengths_
[
DstVectorDim
],
(
DstVectorDim
==
NumDim
-
2
?
HPerBlock
:
WPerBlock
)),
arg
.
inStrides_
[
DstVectorDim
],
DstScalarPerVector
)
&&
GridwisePermute
::
CheckValidity
(
arg
.
in_grid_desc_
,
arg
.
out_grid_desc_
);
};
};
...
...
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