[CRL] Add unified reduction kernel implementation and uniRunner for demo 2-rank AllReduce #315
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
|
|
yuejiewang
force-pushed
the
dev
branch
2 times, most recently
from
9fdadc7 to
027e885
Compare
MC952-arch
reviewed
Dec 8, 2025
|
|
||
| // (6) set completion flag | ||
| __syncthreads(); | ||
| FLAGCX_DEVICE_THREAD_FENCE(); |
Collaborator
There was a problem hiding this comment.
move FLAGCX_DEVICE_THREAD_FENCE() after t->setComplete()
yuejiewang
force-pushed
the
dev
branch
2 times, most recently
from
70003ed to
642accd
Compare
…hile loop in RedKernel
yuejiewang
force-pushed
the
dev
branch
2 times, most recently
from
b3d8a81 to
425bb29
Compare
MC952-arch
reviewed
Jan 4, 2026
flagcx/core/transport.cc
Outdated
| #define ENABLE_TIMER 0 | ||
| #include "timer.h" | ||
|
|
||
| FLAGCX_PARAM(RunUniRunnerAllReduce, "RUN_UNIRUNNER_ALLREDUCE", 0); |
MC952-arch
reviewed
Jan 4, 2026
| } | ||
|
|
||
| FLAGCX_HOST_DECORATOR uint64_t flagcxReduceTrigger::pollState() { | ||
| uint64_t curr_val = __atomic_load_n(&this->value[3], __ATOMIC_ACQUIRE); |
MC952-arch
reviewed
Jan 4, 2026
|
|
||
| FLAGCX_HOST_DECORATOR void flagcxReduceTrigger::setState(int state) { | ||
| uint64_t curr_val = __atomic_load_n(&this->value[3], __ATOMIC_ACQUIRE); | ||
| curr_val &= ~(flagcxTriggerMask(flagcxReduceTriggerBitsState) |
MC952-arch
reviewed
Jan 4, 2026
flagcx/core/transport.cc
Outdated
| FLAGCX_PARAM(P2pDisable, "P2P_DISABLE", 0); | ||
|
|
||
| static inline bool isSameNode(struct flagcxHeteroComm *comm, int peer) { | ||
| // force use network transport for unirunner allreduce |
Collaborator
There was a problem hiding this comment.
// force use net transport
MC952-arch
reviewed
Jan 4, 2026
| FLAGCX_DEVICE_INLINE_DECORATOR flagcxResult_t dequeue(volatile uint64_t *buffer, | ||
| int *idx) { | ||
| while (true) { | ||
| unsigned long long int old_c = *(buffer + 1); |
Collaborator
There was a problem hiding this comment.
old_c -> oldConsumed
cur_p -> curProduced
MC952-arch
reviewed
Jan 4, 2026
| uint64_t nthreads, uint64_t datatype, uint64_t redOp) { | ||
| // to be implemented by vendors | ||
| int tid = threadIdx.x; | ||
| float *fst_ptr = (float *)fst; |
MC952-arch
reviewed
Jan 4, 2026
| float *snd_ptr = (float *)snd; | ||
| float *out_ptr = (float *)out; | ||
| for (int i = tid; i < count; i += nthreads) { | ||
| out_ptr[i] = fst_ptr[i] + snd_ptr[i]; |
MC952-arch
reviewed
Jan 4, 2026
|
|
||
| FLAGCX_GLOBAL_DECORATOR void flagcxCollectiveKernel(void *fifoBuffer) { | ||
| volatile uint64_t *vBuf = (volatile uint64_t *)fifoBuffer; | ||
| int empty_iter = 0; // backoff counter |
Collaborator
There was a problem hiding this comment.
empty_iter -> emptyIter
MC952-arch
reviewed
Jan 4, 2026
| uint64_t redop; | ||
| int slot = myIdx & (*vBuf - 1); | ||
| if (tid == 0) { | ||
| // printf("block %d get work idx %d, slot %d\n", blockIdx.x, myIdx, slot); |
MC952-arch
reviewed
Jan 4, 2026
| t->setComplete(); | ||
| } | ||
| } | ||
| // FLAGCX_DEVICE_THREAD_FENCE(); |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
5 participants
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
PR Category
CRL
PR Types
New Features
PR Description