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
7887c0ac
Commit
7887c0ac
authored
Mar 10, 2022
by
Jianfeng yan
Browse files
rename num_accesses to num_access
parent
f8301636
Changes
4
Show whitespace changes
Inline
Side-by-side
Showing
4 changed files
with
24 additions
and
24 deletions
+24
-24
include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v6r1.hpp
...tion/gpu/thread/threadwise_tensor_slice_transfer_v6r1.hpp
+6
-6
include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v6r2.hpp
...tion/gpu/thread/threadwise_tensor_slice_transfer_v6r2.hpp
+6
-6
include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v6r3.hpp
...tion/gpu/thread/threadwise_tensor_slice_transfer_v6r3.hpp
+6
-6
test/space_filling_curve/space_filling_curve.cpp
test/space_filling_curve/space_filling_curve.cpp
+6
-6
No files found.
include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v6r1.hpp
View file @
7887c0ac
...
@@ -82,9 +82,9 @@ struct ThreadwiseTensorSliceTransfer_v6r1
...
@@ -82,9 +82,9 @@ struct ThreadwiseTensorSliceTransfer_v6r1
remove_cv_t
<
decltype
(
scalar_per_access
)
>>
;
remove_cv_t
<
decltype
(
scalar_per_access
)
>>
;
// loop over space-filling curve
// loop over space-filling curve
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
)
{
using
src_vector_type
=
vector_type_maker_t
<
SrcData
,
ScalarPerVector
>
;
using
src_vector_type
=
vector_type_maker_t
<
SrcData
,
ScalarPerVector
>
;
using
src_vector_t
=
typename
src_vector_type
::
type
;
using
src_vector_t
=
typename
src_vector_type
::
type
;
...
@@ -116,7 +116,7 @@ struct ThreadwiseTensorSliceTransfer_v6r1
...
@@ -116,7 +116,7 @@ struct ThreadwiseTensorSliceTransfer_v6r1
dst_vector_container
.
template
AsType
<
dst_vector_t
>()[
I0
]);
dst_vector_container
.
template
AsType
<
dst_vector_t
>()[
I0
]);
// move coordinate
// move coordinate
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
);
move_tensor_coordinate
(
move_tensor_coordinate
(
...
@@ -153,15 +153,15 @@ struct ThreadwiseTensorSliceTransfer_v6r1
...
@@ -153,15 +153,15 @@ struct ThreadwiseTensorSliceTransfer_v6r1
DimAccessOrder
,
DimAccessOrder
,
remove_cv_t
<
decltype
(
scalar_per_access
)
>>
;
remove_cv_t
<
decltype
(
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
;
}
}
...
...
include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v6r2.hpp
View file @
7887c0ac
...
@@ -97,10 +97,10 @@ struct ThreadwiseTensorSliceTransfer_v6r2
...
@@ -97,10 +97,10 @@ struct ThreadwiseTensorSliceTransfer_v6r2
DimAccessOrder
,
DimAccessOrder
,
remove_cv_t
<
decltype
(
scalar_per_access
)
>>
;
remove_cv_t
<
decltype
(
scalar_per_access
)
>>
;
constexpr
auto
num_access
es
=
SpaceFillingCurve
::
GetNumOfAccess
();
constexpr
auto
num_access
=
SpaceFillingCurve
::
GetNumOfAccess
();
// loop over space-filling curve
// loop over space-filling curve
static_for
<
0
,
num_access
es
,
1
>
{}([
&
](
auto
idx_1d
)
{
static_for
<
0
,
num_access
,
1
>
{}([
&
](
auto
idx_1d
)
{
using
src0_vector_type
=
vector_type_maker_t
<
Src0Data
,
ScalarPerVector
>
;
using
src0_vector_type
=
vector_type_maker_t
<
Src0Data
,
ScalarPerVector
>
;
using
src0_vector_t
=
typename
src0_vector_type
::
type
;
using
src0_vector_t
=
typename
src0_vector_type
::
type
;
...
@@ -142,7 +142,7 @@ struct ThreadwiseTensorSliceTransfer_v6r2
...
@@ -142,7 +142,7 @@ struct ThreadwiseTensorSliceTransfer_v6r2
dst_vector_container
.
template
AsType
<
dst_vector_t
>()[
I0
]);
dst_vector_container
.
template
AsType
<
dst_vector_t
>()[
I0
]);
// move coordinate
// move coordinate
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
);
move_tensor_coordinate
(
move_tensor_coordinate
(
...
@@ -189,15 +189,15 @@ struct ThreadwiseTensorSliceTransfer_v6r2
...
@@ -189,15 +189,15 @@ struct ThreadwiseTensorSliceTransfer_v6r2
DimAccessOrder
,
DimAccessOrder
,
remove_cv_t
<
decltype
(
scalar_per_access
)
>>
;
remove_cv_t
<
decltype
(
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
;
}
}
...
...
include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v6r3.hpp
View file @
7887c0ac
...
@@ -112,10 +112,10 @@ struct ThreadwiseTensorSliceTransfer_v6r3
...
@@ -112,10 +112,10 @@ struct ThreadwiseTensorSliceTransfer_v6r3
DimAccessOrder
,
DimAccessOrder
,
remove_cv_t
<
decltype
(
scalar_per_access
)
>>
;
remove_cv_t
<
decltype
(
scalar_per_access
)
>>
;
constexpr
auto
num_access
es
=
SpaceFillingCurve
::
GetNumOfAccess
();
constexpr
auto
num_access
=
SpaceFillingCurve
::
GetNumOfAccess
();
// loop over space-filling curve
// loop over space-filling curve
static_for
<
0
,
num_access
es
,
1
>
{}([
&
](
auto
idx_1d
)
{
static_for
<
0
,
num_access
,
1
>
{}([
&
](
auto
idx_1d
)
{
using
src0_vector_type
=
vector_type_maker_t
<
Src0Data
,
ScalarPerVector
>
;
using
src0_vector_type
=
vector_type_maker_t
<
Src0Data
,
ScalarPerVector
>
;
using
src0_vector_t
=
typename
src0_vector_type
::
type
;
using
src0_vector_t
=
typename
src0_vector_type
::
type
;
...
@@ -166,7 +166,7 @@ struct ThreadwiseTensorSliceTransfer_v6r3
...
@@ -166,7 +166,7 @@ struct ThreadwiseTensorSliceTransfer_v6r3
dst_vector_container
.
template
AsType
<
dst_vector_t
>()[
I0
]);
dst_vector_container
.
template
AsType
<
dst_vector_t
>()[
I0
]);
// move coordinate
// move coordinate
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
);
move_tensor_coordinate
(
move_tensor_coordinate
(
...
@@ -223,15 +223,15 @@ struct ThreadwiseTensorSliceTransfer_v6r3
...
@@ -223,15 +223,15 @@ struct ThreadwiseTensorSliceTransfer_v6r3
DimAccessOrder
,
DimAccessOrder
,
remove_cv_t
<
decltype
(
scalar_per_access
)
>>
;
remove_cv_t
<
decltype
(
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
;
}
}
...
...
test/space_filling_curve/space_filling_curve.cpp
View file @
7887c0ac
...
@@ -95,13 +95,13 @@ void traverse_using_space_filling_curve()
...
@@ -95,13 +95,13 @@ void traverse_using_space_filling_curve()
make_tuple
(
12
,
2
,
6
),
make_tuple
(
12
,
2
,
6
),
make_tuple
(
12
,
0
,
6
));
make_tuple
(
12
,
0
,
6
));
constexpr
index_t
num_access
es
=
SpaceFillingCurve
::
GetNumOfAccess
();
constexpr
index_t
num_access
=
SpaceFillingCurve
::
GetNumOfAccess
();
static_assert
(
num_access
es
==
reduce_on_sequence
(
TensorLengths
{}
/
ScalarsPerAccess
{},
static_assert
(
num_access
==
reduce_on_sequence
(
TensorLengths
{}
/
ScalarsPerAccess
{},
math
::
multiplies
{},
math
::
multiplies
{},
Number
<
1
>
{}));
Number
<
1
>
{}));
static_for
<
1
,
num_access
es
,
1
>
{}([
&
](
auto
i
)
{
static_for
<
1
,
num_access
,
1
>
{}([
&
](
auto
i
)
{
constexpr
auto
idx_curr
=
SpaceFillingCurve
::
GetIndex
(
i
);
constexpr
auto
idx_curr
=
SpaceFillingCurve
::
GetIndex
(
i
);
static_assert
(
idx_curr
[
I0
]
==
expected
[
i
][
I0
]);
static_assert
(
idx_curr
[
I0
]
==
expected
[
i
][
I0
]);
...
@@ -115,7 +115,7 @@ void traverse_using_space_filling_curve()
...
@@ -115,7 +115,7 @@ void traverse_using_space_filling_curve()
static_assert
(
backward_step
[
I2
]
==
expected_step
[
I2
]);
static_assert
(
backward_step
[
I2
]
==
expected_step
[
I2
]);
});
});
static_for
<
0
,
num_access
es
-
1
,
1
>
{}([
&
](
auto
i
)
{
static_for
<
0
,
num_access
-
1
,
1
>
{}([
&
](
auto
i
)
{
constexpr
auto
idx_curr
=
SpaceFillingCurve
::
GetIndex
(
i
);
constexpr
auto
idx_curr
=
SpaceFillingCurve
::
GetIndex
(
i
);
static_assert
(
idx_curr
[
I0
]
==
expected
[
i
][
I0
]);
static_assert
(
idx_curr
[
I0
]
==
expected
[
i
][
I0
]);
...
...
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