Skip to content

Commit 421990c

Browse files
MC952-archMC952-arch
committed
Fix critical issues
1 parent 355b104 commit 421990c

2 files changed

Lines changed: 63 additions & 69 deletions

File tree

‎flagcx/adaptor/include/device_api/default_comm_traits.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -819,7 +819,7 @@ struct Barrier<DefaultBackend<P>, flagcxTeamTagIntra, Coop> {
819819
_myRank(team.rank), _nBarriers(dc.nBarriers), _ctaIndex(index),
820820
_epochBuffer(dc.epochBuffer),
821821
_epoch(Atomic::load(&dc.epochBuffer[index],
822-
flagcxDeviceMemoryOrderRelaxed)) {}
822+
flagcxDeviceMemoryOrderAcquire)) {}
823823

824824
// arrive: thread-striped store epoch+1 to each peer's inbox slot for me
825825
FLAGCX_DEVICE_INLINE_DECORATOR void

‎test/unittest/device_api/test_device_ir_intra.cpp‎

Lines changed: 62 additions & 68 deletions
Original file line numberDiff line numberDiff line change
@@ -49,6 +49,13 @@ int main(int argc, char *argv[]) {
4949

5050
int worldSize = 1, worldRank = 0;
5151
int totalProcs = 1, proc = 0;
52+
// Print with rank prefix from all ranks, flushing immediately
53+
#define RPRINTF(...) \
54+
do { \
55+
printf("[R%02d] ", proc); \
56+
printf(__VA_ARGS__); \
57+
fflush(stdout); \
58+
} while (0)
5259
MPI_Comm splitComm;
5360
uint64_t splitMask = 0;
5461
int color = 0;
@@ -122,10 +129,10 @@ int main(int argc, char *argv[]) {
122129
(hostResults[2] == proc) && // single-node: intraRank == rank
123130
(hostResults[3] == totalProcs);
124131

125-
if (proc == 0) {
126-
printf("K1 CommQueries: %s (rank=%d size=%d intraRank=%d intraSize=%d)\n",
127-
k1Pass ? "PASS" : "FAIL", hostResults[0], hostResults[1],
128-
hostResults[2], hostResults[3]);
132+
{
133+
RPRINTF("K1 CommQueries: %s (rank=%d size=%d intraRank=%d intraSize=%d)\n",
134+
k1Pass ? "PASS" : "FAIL", hostResults[0], hostResults[1],
135+
hostResults[2], hostResults[3]);
129136
}
130137

131138
// -------------------------------------------------------------------------
@@ -151,9 +158,7 @@ int main(int argc, char *argv[]) {
151158
}
152159
}
153160

154-
if (proc == 0) {
155-
printf("K2 CoopGroup: %s\n", k2Pass ? "PASS" : "FAIL");
156-
}
161+
{ RPRINTF("K2 CoopGroup: %s\n", k2Pass ? "PASS" : "FAIL"); }
157162

158163
// -------------------------------------------------------------------------
159164
// Test K3: Team Queries
@@ -170,9 +175,7 @@ int main(int argc, char *argv[]) {
170175

171176
bool k3Pass = (hostResults[1] == proc); // worldRank should match proc
172177

173-
if (proc == 0) {
174-
printf("K3 TeamQueries: %s\n", k3Pass ? "PASS" : "FAIL");
175-
}
178+
{ RPRINTF("K3 TeamQueries: %s\n", k3Pass ? "PASS" : "FAIL"); }
176179

177180
// -------------------------------------------------------------------------
178181
// Test K4: Local Pointer
@@ -189,9 +192,7 @@ int main(int argc, char *argv[]) {
189192

190193
bool k4Pass = (hostResults[0] == 1); // Should match raw buffer
191194

192-
if (proc == 0) {
193-
printf("K4 LocalPointer: %s\n", k4Pass ? "PASS" : "FAIL");
194-
}
195+
{ RPRINTF("K4 LocalPointer: %s\n", k4Pass ? "PASS" : "FAIL"); }
195196

196197
// -------------------------------------------------------------------------
197198
// Test K5: Intra Pointer (LSA read)
@@ -235,9 +236,7 @@ int main(int argc, char *argv[]) {
235236
}
236237
}
237238

238-
if (proc == 0) {
239-
printf("K5 IntraPointer: %s\n", k5Pass ? "PASS" : "FAIL");
240-
}
239+
{ RPRINTF("K5 IntraPointer: %s\n", k5Pass ? "PASS" : "FAIL"); }
241240

242241
delete[] hostOutput;
243242
FLAGCXCHECK(devHandle->deviceFree(devOutput, flagcxMemDevice, NULL));
@@ -261,9 +260,7 @@ int main(int argc, char *argv[]) {
261260
(hostResults[3] == 4) && // int32
262261
(hostResults[4] == 8); // uint64
263262

264-
if (proc == 0) {
265-
printf("K6 DataTypeSize: %s\n", k6Pass ? "PASS" : "FAIL");
266-
}
263+
{ RPRINTF("K6 DataTypeSize: %s\n", k6Pass ? "PASS" : "FAIL"); }
267264

268265
// -------------------------------------------------------------------------
269266
// Test K7: Intra Barrier Sync
@@ -298,9 +295,7 @@ int main(int argc, char *argv[]) {
298295
}
299296
}
300297

301-
if (proc == 0) {
302-
printf("K7 IntraBarrierSync: %s\n", k7Pass ? "PASS" : "FAIL");
303-
}
298+
{ RPRINTF("K7 IntraBarrierSync: %s\n", k7Pass ? "PASS" : "FAIL"); }
304299

305300
delete[] hostBarrierResult;
306301
FLAGCXCHECK(devHandle->deviceFree(k7Output, flagcxMemDevice, NULL));
@@ -338,9 +333,7 @@ int main(int argc, char *argv[]) {
338333
}
339334
}
340335

341-
if (proc == 0) {
342-
printf("K8 IntraBarrierArriveWait: %s\n", k8Pass ? "PASS" : "FAIL");
343-
}
336+
{ RPRINTF("K8 IntraBarrierArriveWait: %s\n", k8Pass ? "PASS" : "FAIL"); }
344337

345338
delete[] hostArriveWaitResult;
346339
FLAGCXCHECK(devHandle->deviceFree(k8Output, flagcxMemDevice, NULL));
@@ -350,9 +343,7 @@ int main(int argc, char *argv[]) {
350343
// Scalar IR Tests (S1 - S6)
351344
// =========================================================================
352345

353-
if (proc == 0) {
354-
printf("\n--- Scalar IR Tests ---\n");
355-
}
346+
{ RPRINTF("\n--- Scalar IR Tests ---\n"); }
356347

357348
// -------------------------------------------------------------------------
358349
// S1: Cooperative Group (Scalar)
@@ -382,9 +373,7 @@ int main(int argc, char *argv[]) {
382373
break;
383374
}
384375
}
385-
if (proc == 0) {
386-
printf("S1 CoopGroup(Scalar): %s\n", s1Pass ? "PASS" : "FAIL");
387-
}
376+
{ RPRINTF("S1 CoopGroup(Scalar): %s\n", s1Pass ? "PASS" : "FAIL"); }
388377
delete[] hostS1;
389378
FLAGCXCHECK(devHandle->deviceFree(s1Results, flagcxMemDevice, NULL));
390379

@@ -404,9 +393,7 @@ int main(int argc, char *argv[]) {
404393

405394
// intraRank -> worldRank via INTRA team should give back our proc rank
406395
bool s2Pass = (hostS2[1] == proc);
407-
if (proc == 0) {
408-
printf("S2 TeamQueries(Scalar): %s\n", s2Pass ? "PASS" : "FAIL");
409-
}
396+
{ RPRINTF("S2 TeamQueries(Scalar): %s\n", s2Pass ? "PASS" : "FAIL"); }
410397
FLAGCXCHECK(devHandle->deviceFree(s2Results, flagcxMemDevice, NULL));
411398

412399
// -------------------------------------------------------------------------
@@ -424,9 +411,7 @@ int main(int argc, char *argv[]) {
424411
flagcxMemcpyDeviceToHost, NULL));
425412

426413
bool s3Pass = (hostS3 == 1);
427-
if (proc == 0) {
428-
printf("S3 LocalPointer(Scalar): %s\n", s3Pass ? "PASS" : "FAIL");
429-
}
414+
{ RPRINTF("S3 LocalPointer(Scalar): %s\n", s3Pass ? "PASS" : "FAIL"); }
430415
FLAGCXCHECK(devHandle->deviceFree(s3Results, flagcxMemDevice, NULL));
431416

432417
// -------------------------------------------------------------------------
@@ -463,13 +448,14 @@ int main(int argc, char *argv[]) {
463448
break;
464449
}
465450
}
466-
if (proc == 0) {
467-
printf("S4 IntraPointer(Scalar): %s\n", s4Pass ? "PASS" : "FAIL");
468-
}
451+
{ RPRINTF("S4 IntraPointer(Scalar): %s\n", s4Pass ? "PASS" : "FAIL"); }
469452
delete[] hostInitS4;
470453
delete[] hostS4;
471454
FLAGCXCHECK(devHandle->deviceFree(s4Output, flagcxMemDevice, NULL));
472455

456+
// Ensure all ranks have finished S4 reads before any rank modifies regBuff
457+
MPI_Barrier(MPI_COMM_WORLD);
458+
473459
// -------------------------------------------------------------------------
474460
// S5: Intra Barrier Sync (Scalar)
475461
// -------------------------------------------------------------------------
@@ -499,12 +485,19 @@ int main(int argc, char *argv[]) {
499485
break;
500486
}
501487
}
502-
if (proc == 0) {
503-
printf("S5 IntraBarrierSync(Scalar): %s\n", s5Pass ? "PASS" : "FAIL");
488+
{ RPRINTF("S5 IntraBarrierSync(Scalar): %s\n", s5Pass ? "PASS" : "FAIL"); }
489+
if (!s5Pass) {
490+
RPRINTF(" S5 debug: expected=%f peer=%d intraRank=%d intraSize=%d\n",
491+
expectedS5, peer, comm->localRank, comm->localRanks);
492+
RPRINTF(" S5 values: [%f %f %f %f %f %f %f %f]\n", hostS5[0], hostS5[1],
493+
hostS5[2], hostS5[3], hostS5[4], hostS5[5], hostS5[6], hostS5[7]);
504494
}
505495
delete[] hostS5;
506496
FLAGCXCHECK(devHandle->deviceFree(s5Output, flagcxMemDevice, NULL));
507497

498+
// Ensure all ranks have finished S5 reads before any rank modifies regBuff
499+
MPI_Barrier(MPI_COMM_WORLD);
500+
508501
// -------------------------------------------------------------------------
509502
// S6: Intra Barrier Arrive/Wait (Scalar)
510503
// -------------------------------------------------------------------------
@@ -534,12 +527,16 @@ int main(int argc, char *argv[]) {
534527
break;
535528
}
536529
}
537-
if (proc == 0) {
538-
printf("S6 IntraBarrierArriveWait(Scalar): %s\n", s6Pass ? "PASS" : "FAIL");
530+
{
531+
RPRINTF("S6 IntraBarrierArriveWait(Scalar): %s\n",
532+
s6Pass ? "PASS" : "FAIL");
539533
}
540534
delete[] hostS6;
541535
FLAGCXCHECK(devHandle->deviceFree(s6Output, flagcxMemDevice, NULL));
542536

537+
// Ensure all ranks have finished S6 reads before any rank modifies regBuff
538+
MPI_Barrier(MPI_COMM_WORLD);
539+
543540
// -------------------------------------------------------------------------
544541
// K7b: Intra Barrier Sync(AcqRel)
545542
// -------------------------------------------------------------------------
@@ -569,12 +566,13 @@ int main(int argc, char *argv[]) {
569566
break;
570567
}
571568
}
572-
if (proc == 0) {
573-
printf("K7b IntraBarrierSync(AcqRel): %s\n", k7bPass ? "PASS" : "FAIL");
574-
}
569+
{ RPRINTF("K7b IntraBarrierSync(AcqRel): %s\n", k7bPass ? "PASS" : "FAIL"); }
575570
delete[] hostK7b;
576571
FLAGCXCHECK(devHandle->deviceFree(k7bOutput, flagcxMemDevice, NULL));
577572

573+
// Ensure all ranks have finished K7b reads before any rank modifies regBuff
574+
MPI_Barrier(MPI_COMM_WORLD);
575+
578576
// -------------------------------------------------------------------------
579577
// K8b: Arrive(Release) + Wait(AcqRel)
580578
// -------------------------------------------------------------------------
@@ -604,13 +602,16 @@ int main(int argc, char *argv[]) {
604602
break;
605603
}
606604
}
607-
if (proc == 0) {
608-
printf("K8b IntraBarrierArriveWait(AcqRel): %s\n",
609-
k8bPass ? "PASS" : "FAIL");
605+
{
606+
RPRINTF("K8b IntraBarrierArriveWait(AcqRel): %s\n",
607+
k8bPass ? "PASS" : "FAIL");
610608
}
611609
delete[] hostK8b;
612610
FLAGCXCHECK(devHandle->deviceFree(k8bOutput, flagcxMemDevice, NULL));
613611

612+
// Ensure all ranks have finished K8b reads before any rank modifies regBuff
613+
MPI_Barrier(MPI_COMM_WORLD);
614+
614615
// -------------------------------------------------------------------------
615616
// S5b: ArriveS(Release) + WaitS(Acquire)
616617
// -------------------------------------------------------------------------
@@ -640,13 +641,16 @@ int main(int argc, char *argv[]) {
640641
break;
641642
}
642643
}
643-
if (proc == 0) {
644-
printf("S5b IntraBarrierArriveWait(Split): %s\n",
645-
s5bPass ? "PASS" : "FAIL");
644+
{
645+
RPRINTF("S5b IntraBarrierArriveWait(Split): %s\n",
646+
s5bPass ? "PASS" : "FAIL");
646647
}
647648
delete[] hostS5b;
648649
FLAGCXCHECK(devHandle->deviceFree(s5bOutput, flagcxMemDevice, NULL));
649650

651+
// Ensure all ranks have finished S5b reads before any rank modifies regBuff
652+
MPI_Barrier(MPI_COMM_WORLD);
653+
650654
// -------------------------------------------------------------------------
651655
// S5c: SyncS(Release) + read + SyncS(Acquire)
652656
// -------------------------------------------------------------------------
@@ -676,18 +680,14 @@ int main(int argc, char *argv[]) {
676680
break;
677681
}
678682
}
679-
if (proc == 0) {
680-
printf("S5c IntraBarrierSync(Split): %s\n", s5cPass ? "PASS" : "FAIL");
681-
}
683+
{ RPRINTF("S5c IntraBarrierSync(Split): %s\n", s5cPass ? "PASS" : "FAIL"); }
682684
delete[] hostS5c;
683685
FLAGCXCHECK(devHandle->deviceFree(s5cOutput, flagcxMemDevice, NULL));
684686

685687
// -------------------------------------------------------------------------
686688
// S7: TILE_SPAN Cooperative Group
687689
// -------------------------------------------------------------------------
688-
if (proc == 0) {
689-
printf("\n--- Extended Coop Tests ---\n");
690-
}
690+
{ RPRINTF("\n--- Extended Coop Tests ---\n"); }
691691

692692
int nBlocksS7 = 4, nThreadsS7 = 128;
693693
int totalThreadsS7 = nBlocksS7 * nThreadsS7;
@@ -715,9 +715,7 @@ int main(int argc, char *argv[]) {
715715
break;
716716
}
717717
}
718-
if (proc == 0) {
719-
printf("S7 CoopTileSpan: %s\n", s7Pass ? "PASS" : "FAIL");
720-
}
718+
{ RPRINTF("S7 CoopTileSpan: %s\n", s7Pass ? "PASS" : "FAIL"); }
721719
delete[] hostS7;
722720
FLAGCXCHECK(devHandle->deviceFree(s7Results, flagcxMemDevice, NULL));
723721

@@ -743,9 +741,7 @@ int main(int argc, char *argv[]) {
743741
break;
744742
}
745743
}
746-
if (proc == 0) {
747-
printf("S8 CoopLanes: %s\n", s8Pass ? "PASS" : "FAIL");
748-
}
744+
{ RPRINTF("S8 CoopLanes: %s\n", s8Pass ? "PASS" : "FAIL"); }
749745
delete[] hostS8;
750746
FLAGCXCHECK(devHandle->deviceFree(s8Results, flagcxMemDevice, NULL));
751747

@@ -761,9 +757,7 @@ int main(int argc, char *argv[]) {
761757
int globalPass = 0;
762758
MPI_Allreduce(&allPass, &globalPass, 1, MPI_INT, MPI_MIN, MPI_COMM_WORLD);
763759

764-
if (proc == 0) {
765-
printf("\n=== Overall: %s ===\n", globalPass ? "PASS" : "FAIL");
766-
}
760+
{ RPRINTF("\n=== Overall: %s ===\n", globalPass ? "PASS" : "FAIL"); }
767761

768762
// Cleanup
769763
FLAGCXCHECK(devHandle->deviceFree(devResults, flagcxMemDevice, NULL));

0 commit comments

Comments
 (0)