Skip to content

Commit e74699e

Browse files
MC952-archMC952-arch
committed
Fix S21 Intra issues
1 parent 916d33e commit e74699e

3 files changed

Lines changed: 72 additions & 98 deletions

File tree

‎test/kernel/nvidia/device_ir.cu‎

Lines changed: 62 additions & 88 deletions
Original file line numberDiff line numberDiff line change
@@ -1776,141 +1776,74 @@ __global__ void kernelDevSignalStandaloneIntraWorldS(const void *devCommPtr,
17761776

17771777
#define S20_INTRA_COMBO(comboIdx, slot, teamKind, peer) \
17781778
do { \
1779-
if (FLAGCX_THREAD_IDX_X == 0) { \
1780-
printf("[rank %d block %d ctx %d] S20 combo=%d slot=%d peer=%d: enter\n",\
1781-
worldRank, myBlockIdx, (int)contextId, (comboIdx), (int)(slot), \
1782-
(int)(peer)); \
1783-
} \
1784-
/* Reset and verify zero */ \
1779+
/* Reset local signal (local op), then verify zero, then barrier */ \
17851780
if (FLAGCX_THREAD_IDX_X == 0) { \
17861781
flagcxDevResetSignal(devCommPtr, contextId, (flagcxDevSignal_t)(slot)); \
1787-
uint64_t v0 = flagcxDevReadSignal(devCommPtr, (flagcxDevSignal_t)(slot), \
1788-
64, contextId, \
1789-
flagcxDeviceMemoryOrderAcquire); \
1790-
printf("[rank %d block %d ctx %d] S20 combo=%d slot=%d: RESET-A " \
1791-
"post=%llu\n", \
1792-
worldRank, myBlockIdx, (int)contextId, (comboIdx), (int)(slot), \
1793-
(unsigned long long)v0); \
17941782
} \
17951783
flagcxCoopSyncS(FLAGCX_COOP_BLOCK); \
1796-
flagcxDevBarrierSync(devCommPtr, FLAGCX_TEAM_INTRA, myBlockIdx, \
1797-
contextId, FLAGCX_COOP_BLOCK, \
1798-
flagcxDeviceMemoryOrderAcqRel, \
1799-
flagcxDeviceScopeSystem); \
18001784
if (FLAGCX_THREAD_IDX_X == 0) { \
18011785
uint64_t v = flagcxDevReadSignal(devCommPtr, (flagcxDevSignal_t)(slot), \
18021786
64, contextId, \
18031787
flagcxDeviceMemoryOrderAcquire); \
1804-
printf("[rank %d block %d ctx %d] S20 combo=%d slot=%d: READ " \
1805-
"post-barrier-A value=%llu (expect 0)\n", \
1806-
worldRank, myBlockIdx, (int)contextId, (comboIdx), (int)(slot), \
1807-
(unsigned long long)v); \
18081788
if (v != 0) ok = false; \
18091789
} \
18101790
flagcxCoopSyncS(FLAGCX_COOP_BLOCK); \
1811-
if (FLAGCX_THREAD_IDX_X == 0) { \
1812-
printf("[rank %d block %d ctx %d] S20 combo=%d slot=%d: leg A " \
1813-
"pre-wait\n", \
1814-
worldRank, myBlockIdx, (int)contextId, (comboIdx), (int)(slot)); \
1815-
} \
1791+
/* Barrier: all ranks have reset before any rank sends signal */ \
1792+
flagcxDevBarrierSync(devCommPtr, FLAGCX_TEAM_INTRA, myBlockIdx, \
1793+
contextId, FLAGCX_COOP_BLOCK, \
1794+
flagcxDeviceMemoryOrderAcqRel, \
1795+
flagcxDeviceScopeSystem); \
18161796
/* Leg A: SignalInc */ \
18171797
if (FLAGCX_THREAD_IDX_X == 0) { \
18181798
flagcxDevSignalInc(devCommPtr, teamKind, peer, \
18191799
(flagcxDevSignal_t)(slot), contextId, \
18201800
FLAGCX_COOP_THREAD, flagcxDeviceScopeSystem); \
1821-
printf("[rank %d block %d ctx %d] S20 combo=%d slot=%d: SIGNAL-INC " \
1822-
"peer=%d\n", \
1823-
worldRank, myBlockIdx, (int)contextId, (comboIdx), (int)(slot), \
1824-
(int)(peer)); \
18251801
} \
18261802
flagcxCoopSyncS(FLAGCX_COOP_BLOCK); \
1827-
if (FLAGCX_THREAD_IDX_X == 0) { \
1828-
uint64_t v = flagcxDevReadSignal(devCommPtr, (flagcxDevSignal_t)(slot), \
1829-
64, contextId, \
1830-
flagcxDeviceMemoryOrderAcquire); \
1831-
printf("[rank %d block %d ctx %d] S20 combo=%d slot=%d: WAIT-A pre " \
1832-
"least=%llu cur=%llu\n", \
1833-
worldRank, myBlockIdx, (int)contextId, (comboIdx), (int)(slot), \
1834-
(unsigned long long)expectedInc, (unsigned long long)v); \
1835-
} \
18361803
flagcxDevWaitSignal(devCommPtr, (flagcxDevSignal_t)(slot), expectedInc, 64,\
18371804
contextId, FLAGCX_COOP_BLOCK, \
18381805
flagcxDeviceMemoryOrderAcquire); \
18391806
if (FLAGCX_THREAD_IDX_X == 0) { \
1840-
printf("[rank %d block %d ctx %d] S20 combo=%d slot=%d: leg A " \
1841-
"post-wait\n", \
1842-
worldRank, myBlockIdx, (int)contextId, (comboIdx), (int)(slot)); \
1807+
uint64_t v = flagcxDevReadSignal(devCommPtr, (flagcxDevSignal_t)(slot), \
1808+
64, contextId, \
1809+
flagcxDeviceMemoryOrderAcquire); \
1810+
if (v != expectedInc) ok = false; \
18431811
} \
1844-
/* Reset before Leg B */ \
1812+
flagcxCoopSyncS(FLAGCX_COOP_BLOCK); \
1813+
/* Reset before Leg B (local), then verify, then barrier */ \
18451814
if (FLAGCX_THREAD_IDX_X == 0) { \
18461815
flagcxDevResetSignal(devCommPtr, contextId, (flagcxDevSignal_t)(slot)); \
1847-
uint64_t v0 = flagcxDevReadSignal(devCommPtr, (flagcxDevSignal_t)(slot), \
1848-
64, contextId, \
1849-
flagcxDeviceMemoryOrderAcquire); \
1850-
printf("[rank %d block %d ctx %d] S20 combo=%d slot=%d: RESET-B " \
1851-
"post=%llu\n", \
1852-
worldRank, myBlockIdx, (int)contextId, (comboIdx), (int)(slot), \
1853-
(unsigned long long)v0); \
18541816
} \
1855-
flagcxCoopSyncS(FLAGCX_COOP_BLOCK); \
1817+
flagcxCoopSyncS(FLAGCX_COOP_BLOCK); \
18561818
if (FLAGCX_THREAD_IDX_X == 0) { \
18571819
uint64_t v = flagcxDevReadSignal(devCommPtr, (flagcxDevSignal_t)(slot), \
18581820
64, contextId, \
18591821
flagcxDeviceMemoryOrderAcquire); \
1860-
printf("[rank %d block %d ctx %d] S20 combo=%d slot=%d: READ " \
1861-
"post-barrier-B value=%llu (expect 0)\n", \
1862-
worldRank, myBlockIdx, (int)contextId, (comboIdx), (int)(slot), \
1863-
(unsigned long long)v); \
1822+
if (v != 0) ok = false; \
18641823
} \
18651824
flagcxCoopSyncS(FLAGCX_COOP_BLOCK); \
1866-
if (FLAGCX_THREAD_IDX_X == 0) { \
1867-
printf("[rank %d block %d ctx %d] S20 combo=%d slot=%d: leg B " \
1868-
"pre-wait\n", \
1869-
worldRank, myBlockIdx, (int)contextId, (comboIdx), (int)(slot)); \
1870-
} \
1825+
/* Barrier: all ranks have reset before any rank sends Leg B signal */ \
1826+
flagcxDevBarrierSync(devCommPtr, FLAGCX_TEAM_INTRA, myBlockIdx, \
1827+
contextId, FLAGCX_COOP_BLOCK, \
1828+
flagcxDeviceMemoryOrderAcqRel, \
1829+
flagcxDeviceScopeSystem); \
18711830
/* Leg B: SignalAdd(value=5) */ \
18721831
if (FLAGCX_THREAD_IDX_X == 0) { \
18731832
flagcxDevSignalAdd(devCommPtr, teamKind, peer, \
18741833
(flagcxDevSignal_t)(slot), (uint64_t)5, contextId, \
18751834
FLAGCX_COOP_THREAD, flagcxDeviceScopeSystem); \
1876-
printf("[rank %d block %d ctx %d] S20 combo=%d slot=%d: SIGNAL-ADD " \
1877-
"peer=%d value=5\n", \
1878-
worldRank, myBlockIdx, (int)contextId, (comboIdx), (int)(slot), \
1879-
(int)(peer)); \
18801835
} \
18811836
flagcxCoopSyncS(FLAGCX_COOP_BLOCK); \
1882-
if (FLAGCX_THREAD_IDX_X == 0) { \
1883-
uint64_t v = flagcxDevReadSignal(devCommPtr, (flagcxDevSignal_t)(slot), \
1884-
64, contextId, \
1885-
flagcxDeviceMemoryOrderAcquire); \
1886-
printf("[rank %d block %d ctx %d] S20 combo=%d slot=%d: WAIT-B pre " \
1887-
"least=%llu cur=%llu\n", \
1888-
worldRank, myBlockIdx, (int)contextId, (comboIdx), (int)(slot), \
1889-
(unsigned long long)expectedAdd, (unsigned long long)v); \
1890-
} \
18911837
flagcxDevWaitSignal(devCommPtr, (flagcxDevSignal_t)(slot), expectedAdd, 64,\
18921838
contextId, FLAGCX_COOP_BLOCK, \
18931839
flagcxDeviceMemoryOrderAcquire); \
1894-
if (FLAGCX_THREAD_IDX_X == 0) { \
1895-
printf("[rank %d block %d ctx %d] S20 combo=%d slot=%d: leg B " \
1896-
"post-wait\n", \
1897-
worldRank, myBlockIdx, (int)contextId, (comboIdx), (int)(slot)); \
1898-
} \
18991840
if (FLAGCX_THREAD_IDX_X == 0) { \
19001841
uint64_t v = flagcxDevReadSignal(devCommPtr, (flagcxDevSignal_t)(slot), \
19011842
64, contextId, \
19021843
flagcxDeviceMemoryOrderAcquire); \
1903-
printf("[rank %d block %d ctx %d] S20 combo=%d slot=%d: WAIT-B post " \
1904-
"value=%llu (expect %llu)\n", \
1905-
worldRank, myBlockIdx, (int)contextId, (comboIdx), (int)(slot), \
1906-
(unsigned long long)v, (unsigned long long)expectedAdd); \
19071844
if (v != expectedAdd) ok = false; \
19081845
} \
19091846
flagcxCoopSyncS(FLAGCX_COOP_BLOCK); \
1910-
if (FLAGCX_THREAD_IDX_X == 0) { \
1911-
printf("[rank %d block %d ctx %d] S20 combo=%d slot=%d: exit\n", \
1912-
worldRank, myBlockIdx, (int)contextId, (comboIdx), (int)(slot)); \
1913-
} \
19141847
} while (0)
19151848

19161849
// combo 0: THREAD + INTRA (slot 0)
@@ -2092,34 +2025,75 @@ __global__ void kernelDevPutSignalWaitIntraWorldS(const void *devCommPtr,
20922025

20932026
#define S21_INTRA_COMBO(slot, teamKind, peer, expected) \
20942027
do { \
2028+
if (myBlockIdx == 0 && FLAGCX_THREAD_IDX_X == 0) { \
2029+
printf("[rank %d blk %d ctx %d] S21 slot=%d team=%d peer=%d: enter\n", \
2030+
worldRank, myBlockIdx, (int)contextId, (int)(slot), \
2031+
(int)(teamKind), (int)(peer)); \
2032+
} \
20952033
if (myBlockIdx == 0 && FLAGCX_THREAD_IDX_X == 0) \
20962034
flagcxDevResetSignal(devCommPtr, contextId, (flagcxDevSignal_t)(slot)); \
20972035
flagcxCoopSyncS(FLAGCX_COOP_BLOCK); \
2036+
if (myBlockIdx == 0 && FLAGCX_THREAD_IDX_X == 0) { \
2037+
uint64_t v0 = flagcxDevReadSignal(devCommPtr, (flagcxDevSignal_t)(slot), \
2038+
64, contextId, \
2039+
flagcxDeviceMemoryOrderAcquire); \
2040+
printf("[rank %d blk %d ctx %d] S21 slot=%d: RESET sig=%llu\n", \
2041+
worldRank, myBlockIdx, (int)contextId, (int)(slot), \
2042+
(unsigned long long)v0); \
2043+
} \
2044+
flagcxCoopSyncS(FLAGCX_COOP_BLOCK); \
20982045
if (myBlockIdx == 0 && FLAGCX_THREAD_IDX_X == 0) { \
20992046
size_t off = (slot)*bytes; \
2047+
printf("[rank %d blk %d ctx %d] S21 slot=%d: PUT off=%zu bytes=%zu\n", \
2048+
worldRank, myBlockIdx, (int)contextId, (int)(slot), \
2049+
off, bytes); \
21002050
flagcxDevPut(devCommPtr, dstMemPtr, off, srcMemPtr, off, bytes, \
21012051
teamKind, peer, contextId, FLAGCX_COOP_THREAD, \
21022052
flagcxDeviceScopeSystem, flagcxDeviceMemoryOrderRelease); \
2103-
if ((slot) % 2 == 0) \
2053+
printf("[rank %d blk %d ctx %d] S21 slot=%d: PUT done, now signal\n", \
2054+
worldRank, myBlockIdx, (int)contextId, (int)(slot)); \
2055+
if ((slot) % 2 == 0) { \
21042056
flagcxDevSignalInc(devCommPtr, teamKind, peer, \
21052057
(flagcxDevSignal_t)(slot), contextId, \
21062058
FLAGCX_COOP_THREAD, flagcxDeviceScopeSystem); \
2107-
else \
2059+
printf("[rank %d blk %d ctx %d] S21 slot=%d: SIGNAL-INC done\n", \
2060+
worldRank, myBlockIdx, (int)contextId, (int)(slot)); \
2061+
} else { \
21082062
flagcxDevSignalAdd(devCommPtr, teamKind, peer, \
21092063
(flagcxDevSignal_t)(slot), (uint64_t)3, contextId, \
21102064
FLAGCX_COOP_THREAD, flagcxDeviceScopeSystem); \
2065+
printf("[rank %d blk %d ctx %d] S21 slot=%d: SIGNAL-ADD(3) done\n", \
2066+
worldRank, myBlockIdx, (int)contextId, (int)(slot)); \
2067+
} \
21112068
} \
21122069
flagcxCoopSyncS(FLAGCX_COOP_BLOCK); \
2070+
if (myBlockIdx == 0 && FLAGCX_THREAD_IDX_X == 0) { \
2071+
uint64_t v_pre = flagcxDevReadSignal(devCommPtr, (flagcxDevSignal_t)(slot),\
2072+
64, contextId, \
2073+
flagcxDeviceMemoryOrderAcquire); \
2074+
printf("[rank %d blk %d ctx %d] S21 slot=%d: pre-WAIT sig=%llu expect=%llu\n",\
2075+
worldRank, myBlockIdx, (int)contextId, (int)(slot), \
2076+
(unsigned long long)v_pre, (unsigned long long)expected); \
2077+
} \
21132078
flagcxDevWaitSignal(devCommPtr, (flagcxDevSignal_t)(slot), expected, 64, \
21142079
contextId, FLAGCX_COOP_BLOCK, \
21152080
flagcxDeviceMemoryOrderAcquire); \
21162081
if (myBlockIdx == 0 && FLAGCX_THREAD_IDX_X == 0) { \
2082+
printf("[rank %d blk %d ctx %d] S21 slot=%d: post-WAIT\n", \
2083+
worldRank, myBlockIdx, (int)contextId, (int)(slot)); \
21172084
uint64_t v = flagcxDevReadSignal(devCommPtr, (flagcxDevSignal_t)(slot), \
21182085
64, contextId, \
21192086
flagcxDeviceMemoryOrderAcquire); \
2087+
printf("[rank %d blk %d ctx %d] S21 slot=%d: READ sig=%llu expect=%llu\n",\
2088+
worldRank, myBlockIdx, (int)contextId, (int)(slot), \
2089+
(unsigned long long)v, (unsigned long long)expected); \
21202090
if (v != expected) ok = false; \
21212091
} \
21222092
flagcxCoopSyncS(FLAGCX_COOP_BLOCK); \
2093+
if (myBlockIdx == 0 && FLAGCX_THREAD_IDX_X == 0) { \
2094+
printf("[rank %d blk %d ctx %d] S21 slot=%d: exit\n", \
2095+
worldRank, myBlockIdx, (int)contextId, (int)(slot)); \
2096+
} \
21232097
} while (0)
21242098

21252099
// combo 0: THREAD + INTRA (even → Inc, expected=1)

0 commit comments

Comments
 (0)