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

Use hipExtMallocWithFlags to allocate host memory on APU (#1149)

Browse Source 
Also use SM60 as CUDA compatibility level.

[ROCm/rccl commit: 220066197a]
This commit is contained in:
Wenkai Du
2024-04-17 16:56:38 -07:00
committed by GitHub
parent 34fb1007a7
commit 35f8d269f8
2 changed files with 15 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl/src/include/alloc.h
+10 -1
View File
@@ -27,7 +27,16 @@ ncclResult_t ncclCudaHostCallocDebug(T** ptr, size_t nelem, const char *filefunc
cudaStreamCaptureMode mode = cudaStreamCaptureModeRelaxed;
*ptr = nullptr;
CUDACHECK(cudaThreadExchangeStreamCaptureMode(&mode));
CUDACHECKGOTO(hipHostMalloc(ptr, nelem*sizeof(T), cudaHostAllocMapped), result, finish);
int managed = 0;
CUDACHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeDirectManagedMemAccessFromHost, 0));
if (managed) {
#if defined(HIP_UNCACHED_MEMORY)
CUDACHECKGOTO(hipExtMallocWithFlags((void**)ptr, nelem*sizeof(T), hipDeviceMallocUncached), result, finish);
#else
CUDACHECKGOTO(hipExtMallocWithFlags((void**)ptr, nelem*sizeof(T), hipDeviceMallocFinegrained), result, finish);
#endif
} else
CUDACHECKGOTO(hipHostMalloc(ptr, nelem*sizeof(T), cudaHostAllocMapped), result, finish);
memset(*ptr, 0, nelem*sizeof(T));
finish:
CUDACHECK(cudaThreadExchangeStreamCaptureMode(&mode));