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

gfx11: don't use LL for sendrecv (#853)

Browse Source 
* gfx11: don't use LL for sendrecv

* Use builtin instead of inline asm

[ROCm/rccl commit: f70e3e569b]
This commit is contained in:
Wenkai Du authored and GitHub committed 2023-08-17 08:50:51 -07:00
1 parent 84c77f7652
commit 6fdb4103b7
3 files changed
+10 -7
Download Patch File Download Diff File Expand all files Collapse all files

No files matched your search

projects/rccl/src/enqueue.cc
+2 -1
View File
@@ -318,7 +318,8 @@ static ncclResult_t addP2pToPlan(
// 1 is connIndex
struct ncclConnInfo* conn = isSendNotRecv ?
&comm->channels[channelId].peers[peer]->send[1].conn : &comm->channels[channelId].peers[peer]->recv[1].conn;
info.protocol = ((conn->buffs[NCCL_PROTO_LL] != nullptr) && bytes <= ncclParamP2pLLThreshold()) ? NCCL_PROTO_LL : NCCL_PROTO_SIMPLE;
// do not use LL on gfx11
info.protocol = ((conn->buffs[NCCL_PROTO_LL] != nullptr) && bytes <= ncclParamP2pLLThreshold() && comm->topo->nodes[GPU].nodes[0].gpu.gcn/100 != 11) ? NCCL_PROTO_LL : NCCL_PROTO_SIMPLE;
struct ncclProxyOp proxyOp = {};
NCCLCHECK(ncclProxyComputeP2p(&info, &proxyOp));