From 5482fa8102473decac780faef1f2972ef706aabf 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/clr commit: 13655df76e199e5d97ddacfe80bfe5ed4cb27d88] --- projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h index 9e7bb618fe..62be5794fb 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/projects/clr/hipamd/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 21d7bbab11c50222779c000891ef4f5a60235ea8 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/clr commit: ad11972f4721ac46ed83cf3ac30e10473e63a928] --- projects/clr/hipamd/src/hip_memory.cpp | 12 ++++++++---- 1 file changed, 8 insertions(+), 4 deletions(-) diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index d540ab782b..7cae2b2a10 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/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 0d47ae42039e10c56d453e213232d582bb1f996d 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/clr commit: 9b38380c03d4040712e48e24f1934e5252393383] --- projects/clr/hipamd/src/hip_memory.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 7cae2b2a10..e3823c504e 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/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 c5717a37d7adee2a7085abd63bdc32f7667b805e 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/clr commit: d355122bf975dc25b368f1a84189b50ed8923004] --- .../clr/hipamd/include/hip/hcc_detail/device_functions.h | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h b/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h index a24df22315..b88df2e168 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h +++ b/projects/clr/hipamd/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; }