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
d356c871
Commit
d356c871
authored
Sep 06, 2022
by
Po-Yen, Chen
Browse files
Remove no-longer used type argument
parent
7a6dbadc
Changes
3
Show whitespace changes
Inline
Side-by-side
Showing
3 changed files
with
18 additions
and
32 deletions
+18
-32
example/36_permute/permute_fp16.cpp
example/36_permute/permute_fp16.cpp
+2
-2
include/ck/tensor_operation/gpu/device/device_permute.hpp
include/ck/tensor_operation/gpu/device/device_permute.hpp
+8
-15
include/ck/tensor_operation/gpu/grid/gridwise_permute.hpp
include/ck/tensor_operation/gpu/grid/gridwise_permute.hpp
+8
-15
No files found.
example/36_permute/permute_fp16.cpp
View file @
d356c871
...
...
@@ -6,8 +6,8 @@
using
ADataType
=
F16
;
using
BDataType
=
F16
;
using
DevicePermuteInstance
=
ck
::
tensor_operation
::
device
::
DevicePermute
<
ADataType
,
BDataType
,
PassThrough
,
4
,
8
,
S
<
8
>
,
S
<
1
>
>
;
using
DevicePermuteInstance
=
ck
::
tensor_operation
::
device
::
DevicePermute
<
ADataType
,
BDataType
,
PassThrough
,
4
,
8
,
8
,
1
>
;
#include "run_permute_example.inc"
...
...
include/ck/tensor_operation/gpu/device/device_permute.hpp
View file @
d356c871
...
...
@@ -73,23 +73,16 @@ template <typename InDataType,
typename
ElementwiseOperation
,
index_t
NumDim
,
index_t
MPerThread
,
typename
InScalarPerVector
Seq
,
typename
OutScalarPerVector
Seq
>
index_t
InScalarPerVector
,
index_t
OutScalarPerVector
>
struct
DevicePermute
:
detail
::
DevicePermuteBase
<
DevicePermute
<
InDataType
,
OutDataType
,
ElementwiseOperation
,
NumDim
,
MPerThread
,
InScalarPerVector
Seq
,
OutScalarPerVector
Seq
>>
InScalarPerVector
,
OutScalarPerVector
>>
{
static
constexpr
int
NumInput
=
1
;
static
constexpr
int
NumOutput
=
1
;
static_assert
(
NumInput
==
InScalarPerVectorSeq
::
Size
()
&&
NumOutput
==
OutScalarPerVectorSeq
::
Size
(),
"Tuple size is inconsistent with the number of in/out!"
);
using
InDataTypePointer
=
const
InDataType
*
;
using
OutDataTypePointer
=
OutDataType
*
;
...
...
@@ -156,8 +149,8 @@ struct DevicePermute : detail::DevicePermuteBase<DevicePermute<InDataType,
OutDataTypePointer
,
ElementwiseOperation
,
MPerThread
,
InScalarPerVector
Seq
,
OutScalarPerVector
Seq
>
;
InScalarPerVector
,
OutScalarPerVector
>
;
struct
Argument
:
public
BaseArgument
{
...
...
@@ -243,12 +236,12 @@ struct DevicePermute : detail::DevicePermuteBase<DevicePermute<InDataType,
};
bool
valid
=
true
;
if
(
!
IsScalarPerVectorValid
(
arg
.
inLengths_
,
arg
.
inStrides_
,
InScalarPerVector
Seq
::
At
(
0
)
))
if
(
!
IsScalarPerVectorValid
(
arg
.
inLengths_
,
arg
.
inStrides_
,
InScalarPerVector
))
{
valid
=
false
;
}
if
(
!
IsScalarPerVectorValid
(
arg
.
inLengths_
,
arg
.
outStrides_
,
OutScalarPerVector
Seq
::
At
(
0
)
))
if
(
!
IsScalarPerVectorValid
(
arg
.
inLengths_
,
arg
.
outStrides_
,
OutScalarPerVector
))
{
valid
=
false
;
}
...
...
include/ck/tensor_operation/gpu/grid/gridwise_permute.hpp
View file @
d356c871
...
...
@@ -32,17 +32,10 @@ template <typename InGrid1dDesc,
typename
OutDataTypePointer
,
typename
ElementwiseOperation
,
index_t
MPerThread
,
typename
InScalarPerVector
Seq
,
typename
OutScalarPerVector
Seq
>
index_t
InScalarPerVector
,
index_t
OutScalarPerVector
>
struct
GridwisePermute
{
static
constexpr
index_t
NumInput
=
1
;
static
constexpr
index_t
NumOutput
=
1
;
static_assert
(
NumInput
==
InScalarPerVectorSeq
::
Size
()
&&
NumOutput
==
OutScalarPerVectorSeq
::
Size
(),
"Tuple size is inconsistent with the number of in/out!"
);
static
constexpr
auto
I0
=
Number
<
0
>
{};
static
constexpr
auto
thread_buffer_desc_m
=
...
...
@@ -86,7 +79,7 @@ struct GridwisePermute
Sequence
<
MPerThread
>
,
// SliceLengths
Sequence
<
0
>
,
// DimAccessOrder
0
,
// SrcVectorDim
InScalarPerVector
Seq
::
At
(
0
),
// ScalarPerVector
InScalarPerVector
,
// ScalarPerVector
1
,
// SrcScalarStrideInVector
false
>
{
in_grid_1d_desc
,
thread_global_offset
};
...
...
@@ -99,7 +92,7 @@ struct GridwisePermute
Sequence
<
MPerThread
>
,
// SliceLengths
Sequence
<
0
>
,
// DimAccessOrder
0
,
// SrcVectorDim
OutScalarPerVector
Seq
::
At
(
0
)
,
OutScalarPerVector
,
InMemoryDataOperationEnum
::
Set
,
1
,
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