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
8720ef1a
Commit
8720ef1a
authored
Sep 14, 2023
by
Jing Zhang
Browse files
merge element_wise op with data read
parent
65f984f0
Changes
1
Show whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
15 additions
and
3 deletions
+15
-3
include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v7r2.hpp
...tion/gpu/thread/threadwise_tensor_slice_transfer_v7r2.hpp
+15
-3
No files found.
include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v7r2.hpp
View file @
8720ef1a
...
@@ -143,7 +143,7 @@ struct ThreadwiseTensorSliceTransfer_v7r2
...
@@ -143,7 +143,7 @@ struct ThreadwiseTensorSliceTransfer_v7r2
auto
src_vectors
=
generate_vectors
<
SrcDatas
,
SrcScalarPerVector
>
();
auto
src_vectors
=
generate_vectors
<
SrcDatas
,
SrcScalarPerVector
>
();
auto
dst_vectors
=
generate_vectors
<
DstDatas
,
DstScalarPerVector
>
();
auto
dst_vectors
=
generate_vectors
<
DstDatas
,
DstScalarPerVector
>
();
// copy data from src_bufs into src_vectors
#if 0
static_for<0, nSrc, 1>{}([&](auto i) {
static_for<0, nSrc, 1>{}([&](auto i) {
using src_vector_t = typename remove_cvref_t<decltype(src_vectors[i])>::type;
using src_vector_t = typename remove_cvref_t<decltype(src_vectors[i])>::type;
...
@@ -155,13 +155,25 @@ struct ThreadwiseTensorSliceTransfer_v7r2
...
@@ -155,13 +155,25 @@ struct ThreadwiseTensorSliceTransfer_v7r2
src_bufs[i].template Get<src_vector_t>(src_coords_[i].GetOffset(),
src_bufs[i].template Get<src_vector_t>(src_coords_[i].GetOffset(),
is_src_valid);
is_src_valid);
});
});
#endif
// apply pointwise function
static_for
<
0
,
SrcScalarPerVector
,
1
>
{}([
&
](
auto
i
)
{
static_for
<
0
,
SrcScalarPerVector
,
1
>
{}([
&
](
auto
i
)
{
// get reference to src data
const
auto
src_data_refs
=
generate_tie
(
const
auto
src_data_refs
=
generate_tie
(
// return type should be lvalue
// return type should be lvalue
[
&
](
auto
iSrc
)
->
const
auto
&
{
[
&
](
auto
iSrc
)
->
const
auto
&
{
// copy data from src_bufs into src_vectors
using
src_vector_t
=
typename
remove_cvref_t
<
decltype
(
src_vectors
[
iSrc
])
>::
type
;
const
bool
is_src_valid
=
coordinate_has_valid_offset_assuming_visible_index_is_valid
(
src_descs
[
iSrc
],
src_coords_
[
iSrc
]);
src_vectors
(
iSrc
).
template
AsType
<
src_vector_t
>()(
I0
)
=
src_bufs
[
iSrc
].
template
Get
<
src_vector_t
>(
src_coords_
[
iSrc
].
GetOffset
(),
is_src_valid
);
// get reference to src data
using
SrcData
=
remove_cvref_t
<
tuple_element_t
<
iSrc
.
value
,
SrcDatas
>>
;
using
SrcData
=
remove_cvref_t
<
tuple_element_t
<
iSrc
.
value
,
SrcDatas
>>
;
return
src_vectors
[
iSrc
].
template
AsType
<
SrcData
>()[
i
];
return
src_vectors
[
iSrc
].
template
AsType
<
SrcData
>()[
i
];
...
...
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