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

IFC mix build (#998)

Browse Source 
[ROCm/rccl commit: 7c0f49a878]
This commit is contained in:
Bertan Dogancay
2023-12-02 18:49:52 -07:00
committed by GitHub
parent b38b7fa3a2
commit ae0bdad45c
13 changed files with 21 additions and 19 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl/src/init.cc
+3 -1
View File
@@ -1688,6 +1688,7 @@ static ncclResult_t ncclCommInitRankFunc(struct ncclAsyncJob* job_) {
int* parentRanks = NULL;
int cudaArch;
int64_t stackSize = rcclParamStackSizeOverride() ? rcclParamStackSizeOverride() : maxLocalSizeBytes;
hipDeviceProp_t devProp;
CUDACHECKGOTO(cudaSetDevice(cudaDev), res, fail);
CUDACHECKGOTO(cudaDeviceGetAttribute(&archMajor, cudaDevAttrComputeCapabilityMajor, cudaDev), res, fail);
@@ -1698,7 +1699,8 @@ static ncclResult_t ncclCommInitRankFunc(struct ncclAsyncJob* job_) {
// Set the maximum kernel stack size of all kernels to avoid
// a CUDA memory reconfig on load (c.f. NVSHMEM issue)
#ifdef USE_INDIRECT_FUNCTION_CALL
if (stackSize > 0 && ncclParamSetStackSize() == 1) {
CUDACHECK(hipGetDeviceProperties(&devProp, 0));
if (stackSize > 0 && ncclParamSetStackSize() == 1 && devProp.gcnArch != 940 && devProp.gcnArch != 941 && devProp.gcnArch != 942) {
INFO(NCCL_INIT, "Setting cudaLimitStackSize to %zi maxLocalSizeBytes %zi", stackSize, maxLocalSizeBytes);
CUDACHECKIGNORE(cudaDeviceSetLimit(cudaLimitStackSize, stackSize));
}