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
f8301636
Commit
f8301636
authored
Mar 10, 2022
by
Chao Liu
Browse files
clean
parent
b7a6f810
Changes
1
Show whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
12 additions
and
12 deletions
+12
-12
include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer.hpp
...operation/gpu/thread/threadwise_tensor_slice_transfer.hpp
+12
-12
No files found.
include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer.hpp
View file @
f8301636
...
@@ -123,9 +123,9 @@ struct ThreadwiseTensorSliceTransfer_v1r3
...
@@ -123,9 +123,9 @@ struct ThreadwiseTensorSliceTransfer_v1r3
typename
vector_type_maker
<
DstData
,
DstScalarPerVector
>::
type
dst_vector
;
typename
vector_type_maker
<
DstData
,
DstScalarPerVector
>::
type
dst_vector
;
using
dst_vector_t
=
typename
vector_type_maker
<
DstData
,
DstScalarPerVector
>::
type
::
type
;
using
dst_vector_t
=
typename
vector_type_maker
<
DstData
,
DstScalarPerVector
>::
type
::
type
;
constexpr
auto
num_access
es
=
SpaceFillingCurve
::
GetNumOfAccess
();
constexpr
auto
num_access
=
SpaceFillingCurve
::
GetNumOfAccess
();
static_for
<
0
,
num_access
es
,
1
>
{}([
&
](
auto
idx_1d
)
{
static_for
<
0
,
num_access
,
1
>
{}([
&
](
auto
idx_1d
)
{
constexpr
auto
idx_md
=
SpaceFillingCurve
::
GetIndex
(
idx_1d
);
constexpr
auto
idx_md
=
SpaceFillingCurve
::
GetIndex
(
idx_1d
);
// copy data from src_buf into dst_vector
// copy data from src_buf into dst_vector
...
@@ -152,7 +152,7 @@ struct ThreadwiseTensorSliceTransfer_v1r3
...
@@ -152,7 +152,7 @@ struct ThreadwiseTensorSliceTransfer_v1r3
is_dst_valid
,
is_dst_valid
,
dst_vector
.
template
AsType
<
dst_vector_t
>()[
Number
<
0
>
{}]);
dst_vector
.
template
AsType
<
dst_vector_t
>()[
Number
<
0
>
{}]);
if
constexpr
(
idx_1d
.
value
!=
num_access
es
-
1
)
if
constexpr
(
idx_1d
.
value
!=
num_access
-
1
)
{
{
constexpr
auto
forward_step
=
SpaceFillingCurve
::
GetForwardStep
(
idx_1d
);
constexpr
auto
forward_step
=
SpaceFillingCurve
::
GetForwardStep
(
idx_1d
);
...
@@ -180,15 +180,15 @@ struct ThreadwiseTensorSliceTransfer_v1r3
...
@@ -180,15 +180,15 @@ struct ThreadwiseTensorSliceTransfer_v1r3
DimAccessOrder
,
DimAccessOrder
,
remove_cv_t
<
decltype
(
dst_scalar_per_access
)
>>
;
remove_cv_t
<
decltype
(
dst_scalar_per_access
)
>>
;
constexpr
auto
num_access
es
=
SpaceFillingCurve
::
GetNumOfAccess
();
constexpr
auto
num_access
=
SpaceFillingCurve
::
GetNumOfAccess
();
if
constexpr
(
num_access
es
==
0
)
if
constexpr
(
num_access
==
0
)
{
{
return
typename
SpaceFillingCurve
::
Index
{};
return
typename
SpaceFillingCurve
::
Index
{};
}
}
else
else
{
{
constexpr
auto
reset_step
=
constexpr
auto
reset_step
=
SpaceFillingCurve
::
GetStepBetween
(
Number
<
num_access
es
-
1
>
{},
Number
<
0
>
{});
SpaceFillingCurve
::
GetStepBetween
(
Number
<
num_access
-
1
>
{},
Number
<
0
>
{});
return
reset_step
;
return
reset_step
;
}
}
...
@@ -291,9 +291,9 @@ struct ThreadwiseTensorSliceTransfer_v2
...
@@ -291,9 +291,9 @@ struct ThreadwiseTensorSliceTransfer_v2
remove_cv_t
<
decltype
(
src_scalar_per_access
)
>>
;
remove_cv_t
<
decltype
(
src_scalar_per_access
)
>>
;
// loop over tensor and copy
// loop over tensor and copy
constexpr
auto
num_access
es
=
SpaceFillingCurve
::
GetNumOfAccess
();
constexpr
auto
num_access
=
SpaceFillingCurve
::
GetNumOfAccess
();
static_for
<
0
,
num_access
es
,
1
>
{}([
&
](
auto
idx_1d
)
{
static_for
<
0
,
num_access
,
1
>
{}([
&
](
auto
idx_1d
)
{
typename
vector_type_maker
<
SrcData
,
SrcScalarPerVector
>::
type
src_vector
;
typename
vector_type_maker
<
SrcData
,
SrcScalarPerVector
>::
type
src_vector
;
using
src_vector_t
=
using
src_vector_t
=
...
@@ -316,7 +316,7 @@ struct ThreadwiseTensorSliceTransfer_v2
...
@@ -316,7 +316,7 @@ struct ThreadwiseTensorSliceTransfer_v2
dst_buf
(
Number
<
dst_offset
>
{})
=
src_vector
.
template
AsType
<
SrcData
>()[
i
];
dst_buf
(
Number
<
dst_offset
>
{})
=
src_vector
.
template
AsType
<
SrcData
>()[
i
];
});
});
if
constexpr
(
idx_1d
.
value
!=
num_access
es
-
1
)
if
constexpr
(
idx_1d
.
value
!=
num_access
-
1
)
{
{
constexpr
auto
forward_step
=
SpaceFillingCurve
::
GetForwardStep
(
idx_1d
);
constexpr
auto
forward_step
=
SpaceFillingCurve
::
GetForwardStep
(
idx_1d
);
...
@@ -344,15 +344,15 @@ struct ThreadwiseTensorSliceTransfer_v2
...
@@ -344,15 +344,15 @@ struct ThreadwiseTensorSliceTransfer_v2
DimAccessOrder
,
DimAccessOrder
,
remove_cv_t
<
decltype
(
src_scalar_per_access
)
>>
;
remove_cv_t
<
decltype
(
src_scalar_per_access
)
>>
;
constexpr
auto
num_access
es
=
SpaceFillingCurve
::
GetNumOfAccess
();
constexpr
auto
num_access
=
SpaceFillingCurve
::
GetNumOfAccess
();
if
constexpr
(
num_access
es
==
0
)
if
constexpr
(
num_access
==
0
)
{
{
return
typename
SpaceFillingCurve
::
Index
{};
return
typename
SpaceFillingCurve
::
Index
{};
}
}
else
else
{
{
constexpr
auto
reset_step
=
constexpr
auto
reset_step
=
SpaceFillingCurve
::
GetStepBetween
(
Number
<
num_access
es
-
1
>
{},
Number
<
0
>
{});
SpaceFillingCurve
::
GetStepBetween
(
Number
<
num_access
-
1
>
{},
Number
<
0
>
{});
return
reset_step
;
return
reset_step
;
}
}
...
...
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