AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Improve MSCCL algorithms (#1023)

Browse Source 
[ROCm/rccl commit: 0a53077c9c]
This commit is contained in:
Ziyue Yang
2024-01-04 06:51:34 +08:00
committed by GitHub
parent 0c1f773021
commit e3d45f9de4
11 changed files with 8601 additions and 6 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl/src/collectives/device/msccl_kernel_impl.h
+1 -1
View File
@@ -454,7 +454,7 @@ __device__ __forceinline__ void mscclRunInterpreter(
}
#endif
}
else if (fullOps && t->type == MSCCL_LOCAL_COPY)
else if (t->type == MSCCL_LOCAL_COPY)
prims.localCopy(srcPointer+srcOffset, dstPointer+dstOffset, thisNelem);
else
return;
projects/rccl/src/misc/msccl/msccl_setup.cc
+2 -3
View File
@@ -408,11 +408,10 @@ ncclResult_t mscclSetupKernel(const void* sendBuff, void* recvBuff, size_t count
uint8_t fullOpMask = (1<<MSCCL_RECV_COPY_SEND) |
(1<<MSCCL_RECV_REDUCE_SEND) |
(1<<MSCCL_RECV_REDUCE_COPY_SEND) |
(1<<MSCCL_RECV_REDUCE_COPY) |
(1<<MSCCL_LOCAL_COPY);
(1<<MSCCL_RECV_REDUCE_COPY);
//check if need full ops msccl kernel
if ((hostAlgo->typeMask & fullOpMask) || rcclParamMscclForceFullOps()) {
WARN("MSCCL: this version of MSCCL build doesn't support fill Ops");
WARN("MSCCL: this version of MSCCL build doesn't support full Ops");
return ncclInternalError;
}
projects/rccl/tools/msccl-algorithms/allgather-8n-1kb-4kb.xml
+1 -1
View File
@@ -1,6 +1,6 @@
<!-- Copyright (c) Microsoft Corporation. -->
<!-- Licensed under the MIT License. -->
<algo name="all_gather_llm" proto="LL" nchannels="2" nchunksperloop="16" ngpus="8" coll="allgather" inplace="1" outofplace="0" minBytes="0" maxBytes="4095">
<algo name="all_gather_llm" proto="LL" nchannels="2" nchunksperloop="16" ngpus="8" coll="allgather" inplace="1" outofplace="0" minBytes="1024" maxBytes="4095">
<gpu id="0" i_chunks="0" o_chunks="16" s_chunks="0">
<tb id="0" send="-1" recv="1" chan="0">
<step s="0" type="r" srcbuf="o" srcoff="2" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
projects/rccl/tools/msccl-algorithms/allgather-8n-1mb-16mb.xml → projects/rccl/tools/msccl-algorithms/allgather-8n-1mb-4mb.xml
+1 -1
View File
@@ -1,6 +1,6 @@
<!-- Copyright (c) Microsoft Corporation. -->
<!-- Licensed under the MIT License. -->
<algo name="all_gather_llm" proto="Simple" nchannels="32" nchunksperloop="32" ngpus="8" coll="allgather" inplace="1" outofplace="0" minBytes="1048576" maxBytes="16777216">
<algo name="all_gather_llm" proto="Simple" nchannels="32" nchunksperloop="32" ngpus="8" coll="allgather" inplace="1" outofplace="0" minBytes="1048576" maxBytes="4194303">
<gpu id="0" i_chunks="0" o_chunks="32" s_chunks="0">
<tb id="0" send="-1" recv="1" chan="8">
<step s="0" type="r" srcbuf="o" srcoff="4" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
projects/rccl/tools/msccl-algorithms/allgather-8n-4mb-16mb.xml
+1812
View File
File diff suppressed because it is too large Load Diff
projects/rccl/tools/msccl-algorithms/alltoall-8n-0-1kb.xml
+380
View File
@@ -0,0 +1,380 @@
<!-- Copyright (c) Microsoft Corporation. -->
<!-- Licensed under the MIT License. -->
<algo name="alltoall_llm" proto="LL" nchannels="1" nchunksperloop="8" ngpus="8" coll="alltoall" inplace="0" outofplace="1" minBytes="0" maxBytes="1023">
<gpu id="0" i_chunks="8" o_chunks="8" s_chunks="0">
<tb id="0" send="-1" recv="-1" chan="0">
<step s="0" type="cpy" srcbuf="i" srcoff="0" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="1" send="-1" recv="1" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="0" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="2" send="-1" recv="2" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="0" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="3" send="-1" recv="3" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="0" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="4" send="-1" recv="4" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="0" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="5" send="-1" recv="5" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="0" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="6" send="-1" recv="6" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="0" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="7" send="-1" recv="7" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="0" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="8" send="1" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="1" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="9" send="2" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="2" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="10" send="3" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="3" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="11" send="4" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="4" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="12" send="5" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="5" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="13" send="6" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="6" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="14" send="7" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="7" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
</gpu>
<gpu id="1" i_chunks="8" o_chunks="8" s_chunks="0">
<tb id="0" send="-1" recv="-1" chan="0">
<step s="0" type="cpy" srcbuf="i" srcoff="1" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="1" send="-1" recv="0" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="1" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="2" send="-1" recv="2" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="1" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="3" send="-1" recv="3" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="1" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="4" send="-1" recv="4" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="1" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="5" send="-1" recv="5" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="1" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="6" send="-1" recv="6" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="1" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="7" send="-1" recv="7" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="1" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="8" send="0" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="0" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="9" send="2" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="2" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="10" send="3" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="3" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="11" send="4" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="4" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="12" send="5" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="5" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="13" send="6" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="6" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="14" send="7" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="7" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
</gpu>
<gpu id="2" i_chunks="8" o_chunks="8" s_chunks="0">
<tb id="0" send="-1" recv="-1" chan="0">
<step s="0" type="cpy" srcbuf="i" srcoff="2" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="1" send="-1" recv="0" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="2" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="2" send="-1" recv="1" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="2" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="3" send="-1" recv="3" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="2" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="4" send="-1" recv="4" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="2" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="5" send="-1" recv="5" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="2" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="6" send="-1" recv="6" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="2" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="7" send="-1" recv="7" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="2" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="8" send="0" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="0" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="9" send="1" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="1" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="10" send="3" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="3" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="11" send="4" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="4" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="12" send="5" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="5" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="13" send="6" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="6" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="14" send="7" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="7" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
</gpu>
<gpu id="3" i_chunks="8" o_chunks="8" s_chunks="0">
<tb id="0" send="-1" recv="-1" chan="0">
<step s="0" type="cpy" srcbuf="i" srcoff="3" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="1" send="-1" recv="0" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="3" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="2" send="-1" recv="1" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="3" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="3" send="-1" recv="2" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="3" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="4" send="-1" recv="4" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="3" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="5" send="-1" recv="5" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="3" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="6" send="-1" recv="6" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="3" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="7" send="-1" recv="7" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="3" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="8" send="0" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="0" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="9" send="1" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="1" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="10" send="2" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="2" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="11" send="4" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="4" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="12" send="5" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="5" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="13" send="6" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="6" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="14" send="7" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="7" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
</gpu>
<gpu id="4" i_chunks="8" o_chunks="8" s_chunks="0">
<tb id="0" send="-1" recv="-1" chan="0">
<step s="0" type="cpy" srcbuf="i" srcoff="4" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="1" send="-1" recv="0" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="4" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="2" send="-1" recv="1" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="4" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="3" send="-1" recv="2" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="4" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="4" send="-1" recv="3" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="4" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="5" send="-1" recv="5" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="4" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="6" send="-1" recv="6" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="4" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="7" send="-1" recv="7" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="4" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="8" send="0" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="0" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="9" send="1" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="1" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="10" send="2" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="2" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="11" send="3" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="3" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="12" send="5" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="5" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="13" send="6" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="6" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="14" send="7" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="7" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
</gpu>
<gpu id="5" i_chunks="8" o_chunks="8" s_chunks="0">
<tb id="0" send="-1" recv="-1" chan="0">
<step s="0" type="cpy" srcbuf="i" srcoff="5" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="1" send="-1" recv="0" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="5" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="2" send="-1" recv="1" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="5" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="3" send="-1" recv="2" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="5" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="4" send="-1" recv="3" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="5" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="5" send="-1" recv="4" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="5" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="6" send="-1" recv="6" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="5" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="7" send="-1" recv="7" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="5" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="8" send="0" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="0" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="9" send="1" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="1" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="10" send="2" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="2" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="11" send="3" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="3" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="12" send="4" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="4" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="13" send="6" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="6" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="14" send="7" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="7" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
</gpu>
<gpu id="6" i_chunks="8" o_chunks="8" s_chunks="0">
<tb id="0" send="-1" recv="-1" chan="0">
<step s="0" type="cpy" srcbuf="i" srcoff="6" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="1" send="-1" recv="0" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="6" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="2" send="-1" recv="1" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="6" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="3" send="-1" recv="2" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="6" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="4" send="-1" recv="3" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="6" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="5" send="-1" recv="4" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="6" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="6" send="-1" recv="5" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="6" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="7" send="-1" recv="7" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="6" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="8" send="0" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="0" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="9" send="1" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="1" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="10" send="2" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="2" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="11" send="3" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="3" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="12" send="4" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="4" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="13" send="5" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="5" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="14" send="7" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="7" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
</gpu>
<gpu id="7" i_chunks="8" o_chunks="8" s_chunks="0">
<tb id="0" send="-1" recv="-1" chan="0">
<step s="0" type="cpy" srcbuf="i" srcoff="7" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="1" send="-1" recv="0" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="7" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="2" send="-1" recv="1" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="7" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="3" send="-1" recv="2" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="7" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="4" send="-1" recv="3" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="7" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="5" send="-1" recv="4" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="7" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="6" send="-1" recv="5" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="7" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="7" send="-1" recv="6" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="7" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="8" send="0" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="0" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="9" send="1" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="1" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="10" send="2" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="2" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="11" send="3" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="3" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="12" send="4" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="4" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="13" send="5" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="5" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="14" send="6" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="6" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
</gpu>
</algo>
projects/rccl/tools/msccl-algorithms/alltoall-8n-1kb-4kb.xml
+740
View File
@@ -0,0 +1,740 @@
<!-- Copyright (c) Microsoft Corporation. -->
<!-- Licensed under the MIT License. -->
<algo name="alltoall_llm" proto="LL" nchannels="2" nchunksperloop="16" ngpus="8" coll="alltoall" inplace="0" outofplace="1" minBytes="1024" maxBytes="4095">
<gpu id="0" i_chunks="16" o_chunks="16" s_chunks="0">
<tb id="0" send="-1" recv="-1" chan="0">
<step s="0" type="cpy" srcbuf="i" srcoff="0" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="1" send="-1" recv="-1" chan="1">
<step s="0" type="cpy" srcbuf="i" srcoff="1" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="2" send="-1" recv="1" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="0" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="3" send="-1" recv="1" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="1" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="4" send="-1" recv="2" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="0" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="5" send="-1" recv="2" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="1" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="6" send="-1" recv="3" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="0" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="7" send="-1" recv="3" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="1" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="8" send="-1" recv="4" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="0" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="9" send="-1" recv="4" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="1" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="10" send="-1" recv="5" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="0" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="11" send="-1" recv="5" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="1" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="12" send="-1" recv="6" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="0" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="13" send="-1" recv="6" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="1" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="14" send="-1" recv="7" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="0" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="15" send="-1" recv="7" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="1" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="16" send="1" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="2" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="17" send="1" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="3" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="18" send="2" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="4" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="19" send="2" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="5" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="20" send="3" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="6" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="21" send="3" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="7" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="22" send="4" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="8" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="23" send="4" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="9" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="24" send="5" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="10" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="25" send="5" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="11" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="26" send="6" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="12" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="27" send="6" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="13" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="28" send="7" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="14" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="29" send="7" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="15" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
</gpu>
<gpu id="1" i_chunks="16" o_chunks="16" s_chunks="0">
<tb id="0" send="-1" recv="-1" chan="0">
<step s="0" type="cpy" srcbuf="i" srcoff="2" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="1" send="-1" recv="-1" chan="1">
<step s="0" type="cpy" srcbuf="i" srcoff="3" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="2" send="-1" recv="0" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="2" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="3" send="-1" recv="0" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="3" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="4" send="-1" recv="2" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="2" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="5" send="-1" recv="2" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="3" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="6" send="-1" recv="3" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="2" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="7" send="-1" recv="3" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="3" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="8" send="-1" recv="4" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="2" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="9" send="-1" recv="4" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="3" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="10" send="-1" recv="5" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="2" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="11" send="-1" recv="5" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="3" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="12" send="-1" recv="6" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="2" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="13" send="-1" recv="6" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="3" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="14" send="-1" recv="7" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="2" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="15" send="-1" recv="7" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="3" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="16" send="0" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="0" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="17" send="0" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="1" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="18" send="2" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="4" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="19" send="2" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="5" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="20" send="3" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="6" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="21" send="3" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="7" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="22" send="4" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="8" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="23" send="4" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="9" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="24" send="5" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="10" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="25" send="5" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="11" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="26" send="6" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="12" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="27" send="6" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="13" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="28" send="7" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="14" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="29" send="7" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="15" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
</gpu>
<gpu id="2" i_chunks="16" o_chunks="16" s_chunks="0">
<tb id="0" send="-1" recv="-1" chan="0">
<step s="0" type="cpy" srcbuf="i" srcoff="4" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="1" send="-1" recv="-1" chan="1">
<step s="0" type="cpy" srcbuf="i" srcoff="5" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="2" send="-1" recv="0" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="4" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="3" send="-1" recv="0" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="5" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="4" send="-1" recv="1" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="4" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="5" send="-1" recv="1" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="5" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="6" send="-1" recv="3" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="4" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="7" send="-1" recv="3" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="5" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="8" send="-1" recv="4" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="4" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="9" send="-1" recv="4" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="5" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="10" send="-1" recv="5" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="4" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="11" send="-1" recv="5" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="5" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="12" send="-1" recv="6" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="4" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="13" send="-1" recv="6" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="5" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="14" send="-1" recv="7" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="4" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="15" send="-1" recv="7" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="5" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="16" send="0" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="0" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="17" send="0" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="1" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="18" send="1" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="2" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="19" send="1" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="3" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="20" send="3" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="6" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="21" send="3" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="7" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="22" send="4" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="8" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="23" send="4" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="9" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="24" send="5" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="10" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="25" send="5" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="11" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="26" send="6" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="12" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="27" send="6" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="13" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="28" send="7" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="14" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="29" send="7" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="15" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
</gpu>
<gpu id="3" i_chunks="16" o_chunks="16" s_chunks="0">
<tb id="0" send="-1" recv="-1" chan="0">
<step s="0" type="cpy" srcbuf="i" srcoff="6" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="1" send="-1" recv="-1" chan="1">
<step s="0" type="cpy" srcbuf="i" srcoff="7" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="2" send="-1" recv="0" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="6" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="3" send="-1" recv="0" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="7" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="4" send="-1" recv="1" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="6" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="5" send="-1" recv="1" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="7" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="6" send="-1" recv="2" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="6" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="7" send="-1" recv="2" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="7" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="8" send="-1" recv="4" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="6" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="9" send="-1" recv="4" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="7" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="10" send="-1" recv="5" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="6" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="11" send="-1" recv="5" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="7" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="12" send="-1" recv="6" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="6" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="13" send="-1" recv="6" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="7" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="14" send="-1" recv="7" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="6" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="15" send="-1" recv="7" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="7" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="16" send="0" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="0" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="17" send="0" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="1" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="18" send="1" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="2" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="19" send="1" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="3" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="20" send="2" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="4" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="21" send="2" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="5" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="22" send="4" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="8" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="23" send="4" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="9" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="24" send="5" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="10" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="25" send="5" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="11" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="26" send="6" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="12" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="27" send="6" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="13" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="28" send="7" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="14" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="29" send="7" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="15" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
</gpu>
<gpu id="4" i_chunks="16" o_chunks="16" s_chunks="0">
<tb id="0" send="-1" recv="-1" chan="0">
<step s="0" type="cpy" srcbuf="i" srcoff="8" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="1" send="-1" recv="-1" chan="1">
<step s="0" type="cpy" srcbuf="i" srcoff="9" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="2" send="-1" recv="0" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="8" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="3" send="-1" recv="0" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="9" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="4" send="-1" recv="1" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="8" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="5" send="-1" recv="1" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="9" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="6" send="-1" recv="2" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="8" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="7" send="-1" recv="2" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="9" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="8" send="-1" recv="3" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="8" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="9" send="-1" recv="3" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="9" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="10" send="-1" recv="5" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="8" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="11" send="-1" recv="5" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="9" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="12" send="-1" recv="6" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="8" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="13" send="-1" recv="6" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="9" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="14" send="-1" recv="7" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="8" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="15" send="-1" recv="7" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="9" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="16" send="0" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="0" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="17" send="0" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="1" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="18" send="1" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="2" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="19" send="1" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="3" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="20" send="2" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="4" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="21" send="2" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="5" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="22" send="3" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="6" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="23" send="3" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="7" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="24" send="5" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="10" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="25" send="5" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="11" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="26" send="6" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="12" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="27" send="6" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="13" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="28" send="7" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="14" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="29" send="7" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="15" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
</gpu>
<gpu id="5" i_chunks="16" o_chunks="16" s_chunks="0">
<tb id="0" send="-1" recv="-1" chan="0">
<step s="0" type="cpy" srcbuf="i" srcoff="10" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="1" send="-1" recv="-1" chan="1">
<step s="0" type="cpy" srcbuf="i" srcoff="11" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="2" send="-1" recv="0" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="10" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="3" send="-1" recv="0" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="11" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="4" send="-1" recv="1" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="10" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="5" send="-1" recv="1" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="11" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="6" send="-1" recv="2" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="10" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="7" send="-1" recv="2" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="11" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="8" send="-1" recv="3" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="10" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="9" send="-1" recv="3" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="11" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="10" send="-1" recv="4" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="10" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="11" send="-1" recv="4" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="11" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="12" send="-1" recv="6" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="10" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="13" send="-1" recv="6" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="11" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="14" send="-1" recv="7" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="10" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="15" send="-1" recv="7" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="11" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="16" send="0" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="0" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="17" send="0" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="1" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="18" send="1" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="2" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="19" send="1" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="3" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="20" send="2" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="4" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="21" send="2" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="5" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="22" send="3" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="6" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="23" send="3" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="7" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="24" send="4" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="8" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="25" send="4" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="9" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="26" send="6" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="12" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="27" send="6" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="13" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="28" send="7" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="14" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="29" send="7" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="15" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
</gpu>
<gpu id="6" i_chunks="16" o_chunks="16" s_chunks="0">
<tb id="0" send="-1" recv="-1" chan="0">
<step s="0" type="cpy" srcbuf="i" srcoff="12" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="1" send="-1" recv="-1" chan="1">
<step s="0" type="cpy" srcbuf="i" srcoff="13" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="2" send="-1" recv="0" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="12" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="3" send="-1" recv="0" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="13" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="4" send="-1" recv="1" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="12" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="5" send="-1" recv="1" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="13" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="6" send="-1" recv="2" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="12" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="7" send="-1" recv="2" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="13" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="8" send="-1" recv="3" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="12" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="9" send="-1" recv="3" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="13" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="10" send="-1" recv="4" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="12" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="11" send="-1" recv="4" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="13" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="12" send="-1" recv="5" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="12" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="13" send="-1" recv="5" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="13" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="14" send="-1" recv="7" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="12" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="15" send="-1" recv="7" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="13" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="16" send="0" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="0" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="17" send="0" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="1" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="18" send="1" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="2" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="19" send="1" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="3" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="20" send="2" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="4" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="21" send="2" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="5" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="22" send="3" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="6" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="23" send="3" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="7" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="24" send="4" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="8" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="25" send="4" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="9" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="26" send="5" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="10" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="27" send="5" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="11" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="28" send="7" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="14" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="29" send="7" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="15" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
</gpu>
<gpu id="7" i_chunks="16" o_chunks="16" s_chunks="0">
<tb id="0" send="-1" recv="-1" chan="0">
<step s="0" type="cpy" srcbuf="i" srcoff="14" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="1" send="-1" recv="-1" chan="1">
<step s="0" type="cpy" srcbuf="i" srcoff="15" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="2" send="-1" recv="0" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="14" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="3" send="-1" recv="0" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="15" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="4" send="-1" recv="1" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="14" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="5" send="-1" recv="1" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="15" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="6" send="-1" recv="2" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="14" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="7" send="-1" recv="2" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="15" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="8" send="-1" recv="3" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="14" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="9" send="-1" recv="3" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="15" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="10" send="-1" recv="4" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="14" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="11" send="-1" recv="4" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="15" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="12" send="-1" recv="5" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="14" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="13" send="-1" recv="5" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="15" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="14" send="-1" recv="6" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="14" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="15" send="-1" recv="6" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="15" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="16" send="0" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="0" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="17" send="0" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="1" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="18" send="1" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="2" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="19" send="1" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="3" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="20" send="2" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="4" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="21" send="2" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="5" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="22" send="3" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="6" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="23" send="3" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="7" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="24" send="4" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="8" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="25" send="4" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="9" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="26" send="5" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="10" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="27" send="5" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="11" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="28" send="6" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="12" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="29" send="6" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="13" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
</gpu>
</algo>
projects/rccl/tools/msccl-algorithms/alltoall-8n-1mb-4mb.xml
+1460
View File
File diff suppressed because it is too large Load Diff
projects/rccl/tools/msccl-algorithms/alltoall-8n-4kb-8kb.xml
+740
View File
@@ -0,0 +1,740 @@
<!-- Copyright (c) Microsoft Corporation. -->
<!-- Licensed under the MIT License. -->
<algo name="alltoall_llm" proto="LL" nchannels="16" nchunksperloop="16" ngpus="8" coll="alltoall" inplace="0" outofplace="1" minBytes="4096" maxBytes="8191">
<gpu id="0" i_chunks="16" o_chunks="16" s_chunks="0">
<tb id="0" send="-1" recv="-1" chan="0">
<step s="0" type="cpy" srcbuf="i" srcoff="0" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="1" send="-1" recv="-1" chan="8">
<step s="0" type="cpy" srcbuf="i" srcoff="1" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="2" send="-1" recv="1" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="0" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="3" send="-1" recv="1" chan="9">
<step s="0" type="r" srcbuf="i" srcoff="1" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="4" send="-1" recv="2" chan="2">
<step s="0" type="r" srcbuf="i" srcoff="0" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="5" send="-1" recv="2" chan="10">
<step s="0" type="r" srcbuf="i" srcoff="1" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="6" send="-1" recv="3" chan="3">
<step s="0" type="r" srcbuf="i" srcoff="0" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="7" send="-1" recv="3" chan="11">
<step s="0" type="r" srcbuf="i" srcoff="1" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="8" send="-1" recv="4" chan="4">
<step s="0" type="r" srcbuf="i" srcoff="0" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="9" send="-1" recv="4" chan="12">
<step s="0" type="r" srcbuf="i" srcoff="1" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="10" send="-1" recv="5" chan="5">
<step s="0" type="r" srcbuf="i" srcoff="0" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="11" send="-1" recv="5" chan="13">
<step s="0" type="r" srcbuf="i" srcoff="1" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="12" send="-1" recv="6" chan="6">
<step s="0" type="r" srcbuf="i" srcoff="0" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="13" send="-1" recv="6" chan="14">
<step s="0" type="r" srcbuf="i" srcoff="1" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="14" send="-1" recv="7" chan="7">
<step s="0" type="r" srcbuf="i" srcoff="0" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="15" send="-1" recv="7" chan="15">
<step s="0" type="r" srcbuf="i" srcoff="1" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="16" send="1" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="2" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="17" send="1" recv="-1" chan="8">
<step s="0" type="s" srcbuf="i" srcoff="3" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="18" send="2" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="4" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="19" send="2" recv="-1" chan="8">
<step s="0" type="s" srcbuf="i" srcoff="5" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="20" send="3" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="6" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="21" send="3" recv="-1" chan="8">
<step s="0" type="s" srcbuf="i" srcoff="7" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="22" send="4" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="8" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="23" send="4" recv="-1" chan="8">
<step s="0" type="s" srcbuf="i" srcoff="9" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="24" send="5" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="10" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="25" send="5" recv="-1" chan="8">
<step s="0" type="s" srcbuf="i" srcoff="11" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="26" send="6" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="12" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="27" send="6" recv="-1" chan="8">
<step s="0" type="s" srcbuf="i" srcoff="13" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="28" send="7" recv="-1" chan="0">
<step s="0" type="s" srcbuf="i" srcoff="14" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="29" send="7" recv="-1" chan="8">
<step s="0" type="s" srcbuf="i" srcoff="15" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
</gpu>
<gpu id="1" i_chunks="16" o_chunks="16" s_chunks="0">
<tb id="0" send="-1" recv="-1" chan="1">
<step s="0" type="cpy" srcbuf="i" srcoff="2" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="1" send="-1" recv="-1" chan="9">
<step s="0" type="cpy" srcbuf="i" srcoff="3" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="2" send="-1" recv="0" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="2" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="3" send="-1" recv="0" chan="8">
<step s="0" type="r" srcbuf="i" srcoff="3" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="4" send="-1" recv="2" chan="2">
<step s="0" type="r" srcbuf="i" srcoff="2" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="5" send="-1" recv="2" chan="10">
<step s="0" type="r" srcbuf="i" srcoff="3" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="6" send="-1" recv="3" chan="3">
<step s="0" type="r" srcbuf="i" srcoff="2" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="7" send="-1" recv="3" chan="11">
<step s="0" type="r" srcbuf="i" srcoff="3" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="8" send="-1" recv="4" chan="4">
<step s="0" type="r" srcbuf="i" srcoff="2" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="9" send="-1" recv="4" chan="12">
<step s="0" type="r" srcbuf="i" srcoff="3" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="10" send="-1" recv="5" chan="5">
<step s="0" type="r" srcbuf="i" srcoff="2" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="11" send="-1" recv="5" chan="13">
<step s="0" type="r" srcbuf="i" srcoff="3" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="12" send="-1" recv="6" chan="6">
<step s="0" type="r" srcbuf="i" srcoff="2" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="13" send="-1" recv="6" chan="14">
<step s="0" type="r" srcbuf="i" srcoff="3" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="14" send="-1" recv="7" chan="7">
<step s="0" type="r" srcbuf="i" srcoff="2" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="15" send="-1" recv="7" chan="15">
<step s="0" type="r" srcbuf="i" srcoff="3" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="16" send="0" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="0" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="17" send="0" recv="-1" chan="9">
<step s="0" type="s" srcbuf="i" srcoff="1" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="18" send="2" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="4" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="19" send="2" recv="-1" chan="9">
<step s="0" type="s" srcbuf="i" srcoff="5" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="20" send="3" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="6" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="21" send="3" recv="-1" chan="9">
<step s="0" type="s" srcbuf="i" srcoff="7" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="22" send="4" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="8" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="23" send="4" recv="-1" chan="9">
<step s="0" type="s" srcbuf="i" srcoff="9" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="24" send="5" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="10" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="25" send="5" recv="-1" chan="9">
<step s="0" type="s" srcbuf="i" srcoff="11" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="26" send="6" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="12" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="27" send="6" recv="-1" chan="9">
<step s="0" type="s" srcbuf="i" srcoff="13" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="28" send="7" recv="-1" chan="1">
<step s="0" type="s" srcbuf="i" srcoff="14" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="29" send="7" recv="-1" chan="9">
<step s="0" type="s" srcbuf="i" srcoff="15" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
</gpu>
<gpu id="2" i_chunks="16" o_chunks="16" s_chunks="0">
<tb id="0" send="-1" recv="-1" chan="2">
<step s="0" type="cpy" srcbuf="i" srcoff="4" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="1" send="-1" recv="-1" chan="10">
<step s="0" type="cpy" srcbuf="i" srcoff="5" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="2" send="-1" recv="0" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="4" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="3" send="-1" recv="0" chan="8">
<step s="0" type="r" srcbuf="i" srcoff="5" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="4" send="-1" recv="1" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="4" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="5" send="-1" recv="1" chan="9">
<step s="0" type="r" srcbuf="i" srcoff="5" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="6" send="-1" recv="3" chan="3">
<step s="0" type="r" srcbuf="i" srcoff="4" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="7" send="-1" recv="3" chan="11">
<step s="0" type="r" srcbuf="i" srcoff="5" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="8" send="-1" recv="4" chan="4">
<step s="0" type="r" srcbuf="i" srcoff="4" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="9" send="-1" recv="4" chan="12">
<step s="0" type="r" srcbuf="i" srcoff="5" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="10" send="-1" recv="5" chan="5">
<step s="0" type="r" srcbuf="i" srcoff="4" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="11" send="-1" recv="5" chan="13">
<step s="0" type="r" srcbuf="i" srcoff="5" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="12" send="-1" recv="6" chan="6">
<step s="0" type="r" srcbuf="i" srcoff="4" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="13" send="-1" recv="6" chan="14">
<step s="0" type="r" srcbuf="i" srcoff="5" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="14" send="-1" recv="7" chan="7">
<step s="0" type="r" srcbuf="i" srcoff="4" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="15" send="-1" recv="7" chan="15">
<step s="0" type="r" srcbuf="i" srcoff="5" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="16" send="0" recv="-1" chan="2">
<step s="0" type="s" srcbuf="i" srcoff="0" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="17" send="0" recv="-1" chan="10">
<step s="0" type="s" srcbuf="i" srcoff="1" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="18" send="1" recv="-1" chan="2">
<step s="0" type="s" srcbuf="i" srcoff="2" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="19" send="1" recv="-1" chan="10">
<step s="0" type="s" srcbuf="i" srcoff="3" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="20" send="3" recv="-1" chan="2">
<step s="0" type="s" srcbuf="i" srcoff="6" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="21" send="3" recv="-1" chan="10">
<step s="0" type="s" srcbuf="i" srcoff="7" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="22" send="4" recv="-1" chan="2">
<step s="0" type="s" srcbuf="i" srcoff="8" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="23" send="4" recv="-1" chan="10">
<step s="0" type="s" srcbuf="i" srcoff="9" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="24" send="5" recv="-1" chan="2">
<step s="0" type="s" srcbuf="i" srcoff="10" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="25" send="5" recv="-1" chan="10">
<step s="0" type="s" srcbuf="i" srcoff="11" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="26" send="6" recv="-1" chan="2">
<step s="0" type="s" srcbuf="i" srcoff="12" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="27" send="6" recv="-1" chan="10">
<step s="0" type="s" srcbuf="i" srcoff="13" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="28" send="7" recv="-1" chan="2">
<step s="0" type="s" srcbuf="i" srcoff="14" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="29" send="7" recv="-1" chan="10">
<step s="0" type="s" srcbuf="i" srcoff="15" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
</gpu>
<gpu id="3" i_chunks="16" o_chunks="16" s_chunks="0">
<tb id="0" send="-1" recv="-1" chan="3">
<step s="0" type="cpy" srcbuf="i" srcoff="6" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="1" send="-1" recv="-1" chan="11">
<step s="0" type="cpy" srcbuf="i" srcoff="7" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="2" send="-1" recv="0" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="6" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="3" send="-1" recv="0" chan="8">
<step s="0" type="r" srcbuf="i" srcoff="7" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="4" send="-1" recv="1" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="6" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="5" send="-1" recv="1" chan="9">
<step s="0" type="r" srcbuf="i" srcoff="7" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="6" send="-1" recv="2" chan="2">
<step s="0" type="r" srcbuf="i" srcoff="6" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="7" send="-1" recv="2" chan="10">
<step s="0" type="r" srcbuf="i" srcoff="7" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="8" send="-1" recv="4" chan="4">
<step s="0" type="r" srcbuf="i" srcoff="6" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="9" send="-1" recv="4" chan="12">
<step s="0" type="r" srcbuf="i" srcoff="7" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="10" send="-1" recv="5" chan="5">
<step s="0" type="r" srcbuf="i" srcoff="6" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="11" send="-1" recv="5" chan="13">
<step s="0" type="r" srcbuf="i" srcoff="7" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="12" send="-1" recv="6" chan="6">
<step s="0" type="r" srcbuf="i" srcoff="6" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="13" send="-1" recv="6" chan="14">
<step s="0" type="r" srcbuf="i" srcoff="7" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="14" send="-1" recv="7" chan="7">
<step s="0" type="r" srcbuf="i" srcoff="6" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="15" send="-1" recv="7" chan="15">
<step s="0" type="r" srcbuf="i" srcoff="7" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="16" send="0" recv="-1" chan="3">
<step s="0" type="s" srcbuf="i" srcoff="0" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="17" send="0" recv="-1" chan="11">
<step s="0" type="s" srcbuf="i" srcoff="1" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="18" send="1" recv="-1" chan="3">
<step s="0" type="s" srcbuf="i" srcoff="2" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="19" send="1" recv="-1" chan="11">
<step s="0" type="s" srcbuf="i" srcoff="3" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="20" send="2" recv="-1" chan="3">
<step s="0" type="s" srcbuf="i" srcoff="4" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="21" send="2" recv="-1" chan="11">
<step s="0" type="s" srcbuf="i" srcoff="5" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="22" send="4" recv="-1" chan="3">
<step s="0" type="s" srcbuf="i" srcoff="8" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="23" send="4" recv="-1" chan="11">
<step s="0" type="s" srcbuf="i" srcoff="9" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="24" send="5" recv="-1" chan="3">
<step s="0" type="s" srcbuf="i" srcoff="10" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="25" send="5" recv="-1" chan="11">
<step s="0" type="s" srcbuf="i" srcoff="11" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="26" send="6" recv="-1" chan="3">
<step s="0" type="s" srcbuf="i" srcoff="12" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="27" send="6" recv="-1" chan="11">
<step s="0" type="s" srcbuf="i" srcoff="13" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="28" send="7" recv="-1" chan="3">
<step s="0" type="s" srcbuf="i" srcoff="14" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="29" send="7" recv="-1" chan="11">
<step s="0" type="s" srcbuf="i" srcoff="15" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
</gpu>
<gpu id="4" i_chunks="16" o_chunks="16" s_chunks="0">
<tb id="0" send="-1" recv="-1" chan="4">
<step s="0" type="cpy" srcbuf="i" srcoff="8" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="1" send="-1" recv="-1" chan="12">
<step s="0" type="cpy" srcbuf="i" srcoff="9" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="2" send="-1" recv="0" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="8" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="3" send="-1" recv="0" chan="8">
<step s="0" type="r" srcbuf="i" srcoff="9" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="4" send="-1" recv="1" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="8" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="5" send="-1" recv="1" chan="9">
<step s="0" type="r" srcbuf="i" srcoff="9" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="6" send="-1" recv="2" chan="2">
<step s="0" type="r" srcbuf="i" srcoff="8" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="7" send="-1" recv="2" chan="10">
<step s="0" type="r" srcbuf="i" srcoff="9" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="8" send="-1" recv="3" chan="3">
<step s="0" type="r" srcbuf="i" srcoff="8" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="9" send="-1" recv="3" chan="11">
<step s="0" type="r" srcbuf="i" srcoff="9" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="10" send="-1" recv="5" chan="5">
<step s="0" type="r" srcbuf="i" srcoff="8" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="11" send="-1" recv="5" chan="13">
<step s="0" type="r" srcbuf="i" srcoff="9" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="12" send="-1" recv="6" chan="6">
<step s="0" type="r" srcbuf="i" srcoff="8" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="13" send="-1" recv="6" chan="14">
<step s="0" type="r" srcbuf="i" srcoff="9" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="14" send="-1" recv="7" chan="7">
<step s="0" type="r" srcbuf="i" srcoff="8" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="15" send="-1" recv="7" chan="15">
<step s="0" type="r" srcbuf="i" srcoff="9" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="16" send="0" recv="-1" chan="4">
<step s="0" type="s" srcbuf="i" srcoff="0" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="17" send="0" recv="-1" chan="12">
<step s="0" type="s" srcbuf="i" srcoff="1" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="18" send="1" recv="-1" chan="4">
<step s="0" type="s" srcbuf="i" srcoff="2" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="19" send="1" recv="-1" chan="12">
<step s="0" type="s" srcbuf="i" srcoff="3" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="20" send="2" recv="-1" chan="4">
<step s="0" type="s" srcbuf="i" srcoff="4" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="21" send="2" recv="-1" chan="12">
<step s="0" type="s" srcbuf="i" srcoff="5" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="22" send="3" recv="-1" chan="4">
<step s="0" type="s" srcbuf="i" srcoff="6" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="23" send="3" recv="-1" chan="12">
<step s="0" type="s" srcbuf="i" srcoff="7" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="24" send="5" recv="-1" chan="4">
<step s="0" type="s" srcbuf="i" srcoff="10" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="25" send="5" recv="-1" chan="12">
<step s="0" type="s" srcbuf="i" srcoff="11" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="26" send="6" recv="-1" chan="4">
<step s="0" type="s" srcbuf="i" srcoff="12" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="27" send="6" recv="-1" chan="12">
<step s="0" type="s" srcbuf="i" srcoff="13" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="28" send="7" recv="-1" chan="4">
<step s="0" type="s" srcbuf="i" srcoff="14" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="29" send="7" recv="-1" chan="12">
<step s="0" type="s" srcbuf="i" srcoff="15" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
</gpu>
<gpu id="5" i_chunks="16" o_chunks="16" s_chunks="0">
<tb id="0" send="-1" recv="-1" chan="5">
<step s="0" type="cpy" srcbuf="i" srcoff="10" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="1" send="-1" recv="-1" chan="13">
<step s="0" type="cpy" srcbuf="i" srcoff="11" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="2" send="-1" recv="0" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="10" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="3" send="-1" recv="0" chan="8">
<step s="0" type="r" srcbuf="i" srcoff="11" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="4" send="-1" recv="1" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="10" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="5" send="-1" recv="1" chan="9">
<step s="0" type="r" srcbuf="i" srcoff="11" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="6" send="-1" recv="2" chan="2">
<step s="0" type="r" srcbuf="i" srcoff="10" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="7" send="-1" recv="2" chan="10">
<step s="0" type="r" srcbuf="i" srcoff="11" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="8" send="-1" recv="3" chan="3">
<step s="0" type="r" srcbuf="i" srcoff="10" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="9" send="-1" recv="3" chan="11">
<step s="0" type="r" srcbuf="i" srcoff="11" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="10" send="-1" recv="4" chan="4">
<step s="0" type="r" srcbuf="i" srcoff="10" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="11" send="-1" recv="4" chan="12">
<step s="0" type="r" srcbuf="i" srcoff="11" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="12" send="-1" recv="6" chan="6">
<step s="0" type="r" srcbuf="i" srcoff="10" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="13" send="-1" recv="6" chan="14">
<step s="0" type="r" srcbuf="i" srcoff="11" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="14" send="-1" recv="7" chan="7">
<step s="0" type="r" srcbuf="i" srcoff="10" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="15" send="-1" recv="7" chan="15">
<step s="0" type="r" srcbuf="i" srcoff="11" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="16" send="0" recv="-1" chan="5">
<step s="0" type="s" srcbuf="i" srcoff="0" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="17" send="0" recv="-1" chan="13">
<step s="0" type="s" srcbuf="i" srcoff="1" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="18" send="1" recv="-1" chan="5">
<step s="0" type="s" srcbuf="i" srcoff="2" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="19" send="1" recv="-1" chan="13">
<step s="0" type="s" srcbuf="i" srcoff="3" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="20" send="2" recv="-1" chan="5">
<step s="0" type="s" srcbuf="i" srcoff="4" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="21" send="2" recv="-1" chan="13">
<step s="0" type="s" srcbuf="i" srcoff="5" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="22" send="3" recv="-1" chan="5">
<step s="0" type="s" srcbuf="i" srcoff="6" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="23" send="3" recv="-1" chan="13">
<step s="0" type="s" srcbuf="i" srcoff="7" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="24" send="4" recv="-1" chan="5">
<step s="0" type="s" srcbuf="i" srcoff="8" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="25" send="4" recv="-1" chan="13">
<step s="0" type="s" srcbuf="i" srcoff="9" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="26" send="6" recv="-1" chan="5">
<step s="0" type="s" srcbuf="i" srcoff="12" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="27" send="6" recv="-1" chan="13">
<step s="0" type="s" srcbuf="i" srcoff="13" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="28" send="7" recv="-1" chan="5">
<step s="0" type="s" srcbuf="i" srcoff="14" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="29" send="7" recv="-1" chan="13">
<step s="0" type="s" srcbuf="i" srcoff="15" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
</gpu>
<gpu id="6" i_chunks="16" o_chunks="16" s_chunks="0">
<tb id="0" send="-1" recv="-1" chan="6">
<step s="0" type="cpy" srcbuf="i" srcoff="12" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="1" send="-1" recv="-1" chan="14">
<step s="0" type="cpy" srcbuf="i" srcoff="13" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="2" send="-1" recv="0" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="12" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="3" send="-1" recv="0" chan="8">
<step s="0" type="r" srcbuf="i" srcoff="13" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="4" send="-1" recv="1" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="12" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="5" send="-1" recv="1" chan="9">
<step s="0" type="r" srcbuf="i" srcoff="13" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="6" send="-1" recv="2" chan="2">
<step s="0" type="r" srcbuf="i" srcoff="12" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="7" send="-1" recv="2" chan="10">
<step s="0" type="r" srcbuf="i" srcoff="13" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="8" send="-1" recv="3" chan="3">
<step s="0" type="r" srcbuf="i" srcoff="12" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="9" send="-1" recv="3" chan="11">
<step s="0" type="r" srcbuf="i" srcoff="13" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="10" send="-1" recv="4" chan="4">
<step s="0" type="r" srcbuf="i" srcoff="12" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="11" send="-1" recv="4" chan="12">
<step s="0" type="r" srcbuf="i" srcoff="13" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="12" send="-1" recv="5" chan="5">
<step s="0" type="r" srcbuf="i" srcoff="12" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="13" send="-1" recv="5" chan="13">
<step s="0" type="r" srcbuf="i" srcoff="13" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="14" send="-1" recv="7" chan="7">
<step s="0" type="r" srcbuf="i" srcoff="12" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="15" send="-1" recv="7" chan="15">
<step s="0" type="r" srcbuf="i" srcoff="13" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="16" send="0" recv="-1" chan="6">
<step s="0" type="s" srcbuf="i" srcoff="0" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="17" send="0" recv="-1" chan="14">
<step s="0" type="s" srcbuf="i" srcoff="1" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="18" send="1" recv="-1" chan="6">
<step s="0" type="s" srcbuf="i" srcoff="2" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="19" send="1" recv="-1" chan="14">
<step s="0" type="s" srcbuf="i" srcoff="3" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="20" send="2" recv="-1" chan="6">
<step s="0" type="s" srcbuf="i" srcoff="4" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="21" send="2" recv="-1" chan="14">
<step s="0" type="s" srcbuf="i" srcoff="5" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="22" send="3" recv="-1" chan="6">
<step s="0" type="s" srcbuf="i" srcoff="6" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="23" send="3" recv="-1" chan="14">
<step s="0" type="s" srcbuf="i" srcoff="7" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="24" send="4" recv="-1" chan="6">
<step s="0" type="s" srcbuf="i" srcoff="8" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="25" send="4" recv="-1" chan="14">
<step s="0" type="s" srcbuf="i" srcoff="9" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="26" send="5" recv="-1" chan="6">
<step s="0" type="s" srcbuf="i" srcoff="10" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="27" send="5" recv="-1" chan="14">
<step s="0" type="s" srcbuf="i" srcoff="11" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="28" send="7" recv="-1" chan="6">
<step s="0" type="s" srcbuf="i" srcoff="14" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="29" send="7" recv="-1" chan="14">
<step s="0" type="s" srcbuf="i" srcoff="15" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
</gpu>
<gpu id="7" i_chunks="16" o_chunks="16" s_chunks="0">
<tb id="0" send="-1" recv="-1" chan="7">
<step s="0" type="cpy" srcbuf="i" srcoff="14" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="1" send="-1" recv="-1" chan="15">
<step s="0" type="cpy" srcbuf="i" srcoff="15" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="2" send="-1" recv="0" chan="0">
<step s="0" type="r" srcbuf="i" srcoff="14" dstbuf="o" dstoff="0" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="3" send="-1" recv="0" chan="8">
<step s="0" type="r" srcbuf="i" srcoff="15" dstbuf="o" dstoff="1" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="4" send="-1" recv="1" chan="1">
<step s="0" type="r" srcbuf="i" srcoff="14" dstbuf="o" dstoff="2" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="5" send="-1" recv="1" chan="9">
<step s="0" type="r" srcbuf="i" srcoff="15" dstbuf="o" dstoff="3" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="6" send="-1" recv="2" chan="2">
<step s="0" type="r" srcbuf="i" srcoff="14" dstbuf="o" dstoff="4" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="7" send="-1" recv="2" chan="10">
<step s="0" type="r" srcbuf="i" srcoff="15" dstbuf="o" dstoff="5" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="8" send="-1" recv="3" chan="3">
<step s="0" type="r" srcbuf="i" srcoff="14" dstbuf="o" dstoff="6" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="9" send="-1" recv="3" chan="11">
<step s="0" type="r" srcbuf="i" srcoff="15" dstbuf="o" dstoff="7" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="10" send="-1" recv="4" chan="4">
<step s="0" type="r" srcbuf="i" srcoff="14" dstbuf="o" dstoff="8" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="11" send="-1" recv="4" chan="12">
<step s="0" type="r" srcbuf="i" srcoff="15" dstbuf="o" dstoff="9" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="12" send="-1" recv="5" chan="5">
<step s="0" type="r" srcbuf="i" srcoff="14" dstbuf="o" dstoff="10" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="13" send="-1" recv="5" chan="13">
<step s="0" type="r" srcbuf="i" srcoff="15" dstbuf="o" dstoff="11" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="14" send="-1" recv="6" chan="6">
<step s="0" type="r" srcbuf="i" srcoff="14" dstbuf="o" dstoff="12" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="15" send="-1" recv="6" chan="14">
<step s="0" type="r" srcbuf="i" srcoff="15" dstbuf="o" dstoff="13" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="16" send="0" recv="-1" chan="7">
<step s="0" type="s" srcbuf="i" srcoff="0" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="17" send="0" recv="-1" chan="15">
<step s="0" type="s" srcbuf="i" srcoff="1" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="18" send="1" recv="-1" chan="7">
<step s="0" type="s" srcbuf="i" srcoff="2" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="19" send="1" recv="-1" chan="15">
<step s="0" type="s" srcbuf="i" srcoff="3" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="20" send="2" recv="-1" chan="7">
<step s="0" type="s" srcbuf="i" srcoff="4" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="21" send="2" recv="-1" chan="15">
<step s="0" type="s" srcbuf="i" srcoff="5" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="22" send="3" recv="-1" chan="7">
<step s="0" type="s" srcbuf="i" srcoff="6" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="23" send="3" recv="-1" chan="15">
<step s="0" type="s" srcbuf="i" srcoff="7" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="24" send="4" recv="-1" chan="7">
<step s="0" type="s" srcbuf="i" srcoff="8" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="25" send="4" recv="-1" chan="15">
<step s="0" type="s" srcbuf="i" srcoff="9" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="26" send="5" recv="-1" chan="7">
<step s="0" type="s" srcbuf="i" srcoff="10" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="27" send="5" recv="-1" chan="15">
<step s="0" type="s" srcbuf="i" srcoff="11" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="28" send="6" recv="-1" chan="7">
<step s="0" type="s" srcbuf="i" srcoff="12" dstbuf="o" dstoff="14" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
<tb id="29" send="6" recv="-1" chan="15">
<step s="0" type="s" srcbuf="i" srcoff="13" dstbuf="o" dstoff="15" cnt="1" depid="-1" deps="-1" hasdep="0"/>
</tb>
</gpu>
</algo>
projects/rccl/tools/msccl-algorithms/alltoall-8n-4mb-16mb.xml
+2004
View File
File diff suppressed because it is too large Load Diff
projects/rccl/tools/msccl-algorithms/alltoall-8n-8kb-1mb.xml
+1460
View File
File diff suppressed because it is too large Load Diff