Merge 'master' into 'amd-master'

Change-Id: Id7c377bd22abccbdffd9d1f8bd31a482102813ce


[ROCm/hip commit: e21df30587]

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 <int pattern>
 __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 <int pattern>
 __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;
 }
 

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, ...)          \

projects/hip/src/hip_memory.cpp

@@ -339,21 +339,25 @@ 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;
             }
 
 
             *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;