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
97ec23bf
"...composable_kernel.git" did not exist on "b30edb4ce598901ce998b408338958764e6ab7a4"
Commit
97ec23bf
authored
Jun 13, 2022
by
Chao Liu
Browse files
refactor
parent
7fd5e9f5
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
32 additions
and
13 deletions
+32
-13
include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v7.hpp
...ration/gpu/thread/threadwise_tensor_slice_transfer_v7.hpp
+32
-13
No files found.
include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v7.hpp
View file @
97ec23bf
...
@@ -138,23 +138,42 @@ struct ThreadwiseTensorSliceTransfer_v7
...
@@ -138,23 +138,42 @@ struct ThreadwiseTensorSliceTransfer_v7
coordinate_has_valid_offset_assuming_visible_index_is_valid
(
src_descs
[
i
],
coordinate_has_valid_offset_assuming_visible_index_is_valid
(
src_descs
[
i
],
src_coords_
[
i
]);
src_coords_
[
i
]);
src_vectors
(
i
)
=
src_bufs
[
i
].
template
Get
<
src_vector_t
>(
src_coords_
[
i
].
GetOffset
(),
src_vectors
(
i
).
template
AsType
<
src_vector_t
>()(
I0
)
=
is_src_valid
);
src_bufs
[
i
].
template
Get
<
src_vector_t
>(
src_coords_
[
i
].
GetOffset
(),
is_src_valid
);
});
});
// apply pointwise function
// apply pointwise function
// FIXME: support tuple of arbitary size
static_for
<
0
,
ScalarPerVector
,
1
>
{}([
&
](
auto
i
)
{
static_for
<
0
,
ScalarPerVector
,
1
>
{}([
&
](
auto
i
)
{
using
SrcData0
=
remove_cvref_t
<
decltype
(
SrcDatas
{}[
I0
])
>
;
// get reference to src data
using
SrcData1
=
remove_cvref_t
<
decltype
(
SrcDatas
{}[
I1
])
>
;
const
auto
src_data_refs
=
generate_tie
(
using
SrcData2
=
remove_cvref_t
<
decltype
(
SrcDatas
{}[
I2
])
>
;
// return type should be lvalue
[
&
](
auto
iSrc
)
->
const
auto
&
{
using
DstData0
=
remove_cvref_t
<
decltype
(
DstDatas
{}[
I0
])
>
;
using
SrcData
=
remove_cvref_t
<
tuple_element_t
<
iSrc
.
value
,
SrcDatas
>>
;
element_op_
(
dst_vectors
(
I0
).
template
AsType
<
DstData0
>()(
i
),
return
src_vectors
[
iSrc
].
template
AsType
<
SrcData
>()[
i
];
src_vectors
[
I0
].
template
AsType
<
SrcData0
>()[
i
],
},
src_vectors
[
I1
].
template
AsType
<
SrcData1
>()[
i
],
Number
<
nSrc
>
{});
src_vectors
[
I2
].
template
AsType
<
SrcData2
>()[
i
]);
// get reference to dst data
auto
dst_data_refs
=
generate_tie
(
// return type should be lvalue
[
&
](
auto
iDst
)
->
auto
&
{
using
DstData
=
remove_cvref_t
<
tuple_element_t
<
iDst
.
value
,
DstDatas
>>
;
return
dst_vectors
(
iDst
).
template
AsType
<
DstData
>()(
i
);
},
Number
<
nDst
>
{});
// apply pointwise function
// pointwise function signature:
// element_op_(dst_data_refs[I0],
// dst_data_refs[I1],
// ...,
// src_data_refs[I0],
// src_data_refs[I1],
// ...)
unpack2
(
element_op_
,
dst_data_refs
,
src_data_refs
);
});
});
// copy data from buf_vectors into dst_bufs
// copy data from buf_vectors into dst_bufs
...
...
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