From 5e20bb6370c51c7b524aab9795b39adebfd83708 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Mon, 24 Oct 2016 17:39:43 -0500 Subject: [PATCH] Don't call allow-access if allocating device's only peer is self. Change-Id: Iac58e6c3e460675833f10b1e8b2e393de223654d [ROCm/clr commit: f8b356dce8e5bac16c50c1bfcdb22493ab938c12] --- projects/clr/hipamd/src/hip_memory.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 01f940a408..4f45370c5a 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -125,7 +125,8 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) hc::am_memtracker_update(*ptr, device->_deviceId, 0); { LockedAccessor_CtxCrit_t crit(ctx->criticalData()); - if (crit->peerCnt()) { + // the peerCnt always stores self so make sure the trace actually + if (crit->peerCnt() > 1) { hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr); } }