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
d7206533
Commit
d7206533
authored
Jun 18, 2022
by
Chao Liu
Browse files
add comment
parent
6805df0e
Changes
2
Hide whitespace changes
Inline
Side-by-side
Showing
2 changed files
with
26 additions
and
14 deletions
+26
-14
include/ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_v7.hpp
...ation/gpu/block/thread_group_tensor_slice_transfer_v7.hpp
+12
-4
include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v7.hpp
...ration/gpu/thread/threadwise_tensor_slice_transfer_v7.hpp
+14
-10
No files found.
include/ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_v7.hpp
View file @
d7206533
...
@@ -8,10 +8,18 @@
...
@@ -8,10 +8,18 @@
namespace
ck
{
namespace
ck
{
// this version does following things to avoid scratch memory issue
// Thread-group level multi-source, multi-destination tensor slice data movement
// 1. Use StaticallyIndexedArray instead of C array for thread buffer
// Assume:
// 2. ThreadwiseTensorSliceTransfer_v3 does not keep reference to tensor descriptor
// 1. All sources and destinations are DynamicBuffer
// 3. ThreadwiseTensorSliceTransfer_v3::Run() does not construct new tensor coordinate
// 2. Same VectorDim and ScalerPerVector for all sources and destinations
// 3. DstInMemOps are per destination tensor
// 4. ThreadTransferSrcResetCoordinateAfterRunFlags are per source tensor
// 5. ThreadTransferDstResetCoordinateAfterRunFlags are per destination tensor
//
// Does following things to avoid scratch memory issue
// 1. Pass tensor descritpors by reference (or tuple of references)
// 2. Does not keep reference to tensor descriptor
// 3. Does not construct new tensor coordinate when call Run()
template
<
typename
ThreadGroup
,
template
<
typename
ThreadGroup
,
typename
SrcDatas
,
typename
SrcDatas
,
typename
DstDatas
,
typename
DstDatas
,
...
...
include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_v7.hpp
View file @
d7206533
...
@@ -7,17 +7,21 @@
...
@@ -7,17 +7,21 @@
namespace
ck
{
namespace
ck
{
// Thread-level multi-source, multi-destination tensor slice data movement
// Assume:
// Assume:
// 1. src_descs and dst_descs are not known at compile-time
// 1. All sources and destinations are DynamicBuffer
// 2. SrcBuffers and DstBuffers are DynamicBuffer
// 2. Same VectorDim and ScalerPerVector for all sources and destinations
// 3. src_slice_origins and dst_slice_origins are not known at compile-time,
// 3. DstInMemOps are per destination tensor
// Do following things to avoid "alloca" in LLVM-IR, which would cause scratch memory
// 4. ThreadTransferSrcResetCoordinateAfterRunFlags are per source tensor
// and sometimes useless instructions:
// 5. ThreadTransferDstResetCoordinateAfterRunFlags are per destination tensor
// 1. Don't save a reference to tensor descriptor in class, pass in tensor descriptor as argument
// 6. Does not need to know src_descs and dst_descs at compile-time
// instead
// 7. Does not need to know src_slice_origins and dst_slice_origins at compile-time,
// 2. Don't construct a new tensor coordinate everytime when using it, update and reuse the same
//
// tensor coordinate instead
// Does following things to avoid scratch memory issue
// 3. Don't use a pointer to VGPR buffer, use vector instead
// 1. Use StaticallyIndexedArray or vector_type instead of C array for thread buffer
// 2. Pass tensor descritpors by reference (or tuple of references)
// 3. Does not keep reference to tensor descriptor
// 4. Does not construct new tensor coordinate when call Run()
template
<
typename
SrcDatas
,
template
<
typename
SrcDatas
,
typename
DstDatas
,
typename
DstDatas
,
typename
SrcDescs
,
typename
SrcDescs
,
...
...
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