msccl: allocate scratch as ext-scope fine-grained (#968)
[ROCm/rccl commit: 569d3f7d59]
이 커밋은 다음에 포함됨:
@@ -315,11 +315,7 @@ __device__ __forceinline__ void mscclRunInterpreter(
|
||||
int16_t dependentStep = mscclShmem.mscclTB.dependentStep[dependentPointer+tid];
|
||||
uint64_t goalFlag = COMPUTE_FLAG(workIndex, iter, dependentStep);
|
||||
while (true){
|
||||
#if defined(__gfx942__)
|
||||
uint64_t curFlag = __atomic_load_n(&(mscclFlags + dependentBid)->flag, __ATOMIC_ACQUIRE);
|
||||
#else
|
||||
uint64_t curFlag = __atomic_load_n(&(mscclFlags + dependentBid)->flag, __ATOMIC_RELAXED);
|
||||
#endif
|
||||
if (curFlag >= goalFlag && GET_WORKINDEX_FROM_FLAG(curFlag) == workIndex) break;
|
||||
}
|
||||
}
|
||||
@@ -463,13 +459,8 @@ __device__ __forceinline__ void mscclRunInterpreter(
|
||||
else
|
||||
return;
|
||||
}
|
||||
if (t->hasDependence && tid == nthreads-1){
|
||||
#if defined(__gfx942__)
|
||||
__atomic_store_n(&mscclFlags[bid].flag, (uint64_t) COMPUTE_FLAG(workIndex, iter, step), __ATOMIC_RELEASE);
|
||||
#else
|
||||
if (t->hasDependence && tid == nthreads-1)
|
||||
__atomic_store_n(&mscclFlags[bid].flag, (uint64_t) COMPUTE_FLAG(workIndex, iter, step), __ATOMIC_RELAXED);
|
||||
#endif
|
||||
}
|
||||
step++;
|
||||
}
|
||||
}
|
||||
|
||||
@@ -73,7 +73,7 @@ ncclResult_t mscclSetupScratch(struct mscclAlgo* hostAlgo, hipStream_t stream) {
|
||||
size_t sizeNeeded = (status.nBytes * (size_t)(hostAlgo->nScratchChunks)) / (size_t)(hostAlgo->nChunksPerLoop);
|
||||
if (sizeNeeded > status.scratchBufferSize){
|
||||
NCCLCHECK(ncclCudaFree(status.scratchBuffer));
|
||||
NCCLCHECK(ncclCudaMalloc((char**)&status.scratchBuffer, sizeNeeded));
|
||||
NCCLCHECK(ncclCudaMalloc((char**)&status.scratchBuffer, sizeNeeded, true));
|
||||
status.scratchBufferSize = sizeNeeded;
|
||||
}
|
||||
return ncclSuccess;
|
||||
|
||||
새 이슈에서 참조
사용자 차단