From 94b7856a1ad9e95f22c9f96f5927c59f6df32cd4 Mon Sep 17 00:00:00 2001 From: Michael LIAO Date: Mon, 25 Mar 2019 12:06:46 -0400 Subject: [PATCH 1/4] [hip] Fix typo in macro `hipLaunchKernel` [ROCm/hip commit: b3497824ec29deb8975faf837662b05fbdbe0bec] --- projects/hip/include/hip/hcc_detail/hip_runtime.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime.h b/projects/hip/include/hip/hcc_detail/hip_runtime.h index 9e7bb618fe..62be5794fb 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime.h @@ -336,7 +336,7 @@ typedef int hipLaunchParm; #define hipLaunchKernel(kernelName, numblocks, numthreads, memperblock, streamId, ...) \ do { \ - kernelName<<<(numblocks), (numthreads), (memperblock), (streamId)>>>(hipLaunchParam{}, ##__VA_ARGS__); \ + kernelName<<<(numblocks), (numthreads), (memperblock), (streamId)>>>(hipLaunchParm{}, ##__VA_ARGS__); \ } while (0) #define hipLaunchKernelGGL(kernelName, numblocks, numthreads, memperblock, streamId, ...) \ From c6ef78546477bb40fa7277983dd816cb84492cb4 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Mon, 25 Mar 2019 23:07:05 +0530 Subject: [PATCH 2/4] Avoid double mapping of devices to hostMalloc buffer [ROCm/hip commit: 5e917d70f38b76b2a6f81784b762328b2ed30be6] --- projects/hip/src/hip_memory.cpp | 12 ++++++++---- 1 file changed, 8 insertions(+), 4 deletions(-) diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index d540ab782b..7cae2b2a10 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -339,15 +339,19 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { hip_status = hipErrorInvalidValue; } else { auto device = ctx->getWriteableDevice(); - +#if (__hcc_workweek__ >= 19115) + //Avoid mapping host pinned memory to all devices by HCC + unsigned amFlags = amHostUnmapped; +#else unsigned amFlags = 0; +#endif if (flags & hipHostMallocCoherent) { - amFlags = amHostCoherent; + amFlags |= amHostCoherent; } else if (flags & hipHostMallocNonCoherent) { - amFlags = amHostNonCoherent; + amFlags |= amHostNonCoherent; } else { // depends on env variables: - amFlags = HIP_HOST_COHERENT ? amHostCoherent : amHostNonCoherent; + amFlags |= HIP_HOST_COHERENT ? amHostCoherent : amHostNonCoherent; } From d98d5ca12a07a581f829720d75de7ae0b877d9e1 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Tue, 26 Mar 2019 10:19:13 +0530 Subject: [PATCH 3/4] Let hipHostMalloc always share/map pinned host ptr [ROCm/hip commit: f0af07379360a8c79a3c7fe70ab5a533750484a1] --- projects/hip/src/hip_memory.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index 7cae2b2a10..e3823c504e 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -357,7 +357,7 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { *ptr = hip_internal::allocAndSharePtr( (amFlags & amHostCoherent) ? "finegrained_host" : "pinned_host", sizeBytes, ctx, - (trueFlags & hipHostMallocPortable) /*shareWithAll*/, amFlags, flags, 0); + true /*shareWithAll*/, amFlags, flags, 0); if (sizeBytes && (*ptr == NULL)) { hip_status = hipErrorMemoryAllocation; From a9a397bc4b1ef90b628c00269b5f4248a842d534 Mon Sep 17 00:00:00 2001 From: Michael LIAO Date: Tue, 26 Mar 2019 15:11:42 -0400 Subject: [PATCH 4/4] SWDEV-184380 Fix hcc compilation - `hcc` has no builtin. Need to invoke LLVM intrinsic directly. [ROCm/hip commit: 939d153c1a175da38e944000fbff4f3ca65f30a5] --- projects/hip/include/hip/hcc_detail/device_functions.h | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/projects/hip/include/hip/hcc_detail/device_functions.h b/projects/hip/include/hip/hcc_detail/device_functions.h index a24df22315..b88df2e168 100644 --- a/projects/hip/include/hip/hcc_detail/device_functions.h +++ b/projects/hip/include/hip/hcc_detail/device_functions.h @@ -254,14 +254,22 @@ __device__ static inline float __hip_ds_permutef(int index, float src) { template __device__ static inline unsigned __hip_ds_swizzle_N(unsigned int src) { union { int i; unsigned u; float f; } tmp; tmp.u = src; +#if defined(__HCC__) + tmp.i = __llvm_amdgcn_ds_swizzle(tmp.i, pattern); +#else tmp.i = __builtin_amdgcn_ds_swizzle(tmp.i, pattern); +#endif return tmp.u; } template __device__ static inline float __hip_ds_swizzlef_N(float src) { union { int i; unsigned u; float f; } tmp; tmp.f = src; +#if defined(__HCC__) + tmp.i = __llvm_amdgcn_ds_swizzle(tmp.i, pattern); +#else tmp.i = __builtin_amdgcn_ds_swizzle(tmp.i, pattern); +#endif return tmp.f; }