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

Fix ROCTX call for MSCCL (#1502)

Browse Source 
[ROCm/rccl commit: dd185f26d2]
This commit is contained in:
Bertan Dogancay
2025-01-23 18:00:07 -05:00
committed by GitHub
parent 309de1f068
commit e261b0ef93
1 changed files with 6 additions and 5 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl/src/msccl.cc
+6 -5
View File
@@ -46,15 +46,16 @@ ncclResult_t mscclRunAlgo_impl(
size_t count, ncclDataType_t dataType, int root, int peer, ncclRedOp_t op, size_t count, ncclDataType_t dataType, int root, int peer, ncclRedOp_t op,
mscclAlgoHandle_t mscclAlgoHandle, ncclComm_t comm, hipStream_t stream) { mscclAlgoHandle_t mscclAlgoHandle, ncclComm_t comm, hipStream_t stream) {
struct NvtxParamsMsccl { struct NvtxParamsMsccl {
size_t sendbytes; size_t bytes;
size_t recvbytes; ncclRedOp_t op;
}; };
// Just pass the size of one send/recv messages and not the total bytes sent/received. // Just pass the size of one send/recv messages and not the total bytes sent/received.
constexpr nvtxPayloadSchemaEntry_t MscclSchema[] = { constexpr nvtxPayloadSchemaEntry_t MscclSchema[] = {
{0, NVTX_PAYLOAD_ENTRY_TYPE_SIZE, "Message size [bytes] (Send)"}, {0, NVTX_PAYLOAD_ENTRY_TYPE_SIZE, "Message size [bytes]"},
{0, NVTX_PAYLOAD_ENTRY_TYPE_SIZE, "Message size [bytes] (Recv)"} {0, NVTX_PAYLOAD_ENTRY_NCCL_REDOP, "Reduction operation", nullptr, 0,
offsetof(NvtxParamsMsccl, op)}
}; };
NvtxParamsMsccl payload{sendCounts[comm->rank] * ncclTypeSize(dataType), recvCounts[comm->rank] * ncclTypeSize(dataType)}; NvtxParamsMsccl payload{count * ncclTypeSize(dataType), op};
NVTX3_FUNC_WITH_PARAMS(MSCCL, MscclSchema, payload) NVTX3_FUNC_WITH_PARAMS(MSCCL, MscclSchema, payload)
mscclStatus& status = mscclGetStatus(comm->rank); mscclStatus& status = mscclGetStatus(comm->rank);