Promote till commit 'f1c05e902656c15d7a6c229e0b400a1656dbcb5f'

Change-Id: I1cf4c3b8c671c69dcca9b5f3a453415b48e511ad

Author: satyanveshd

CHANGELOG.md

@@ -7,10 +7,6 @@ Full documentation for HIP is available at [rocm.docs.amd.com](https://rocm.docs
 ### Changed
 * Added new environment variable
     - `DEBUG_HIP_7_PREVIEW` This is used for enabling the backward incompatible changes before the next major ROCm release 7.0. By default this is set to 0. Users can set this variable to 0x1, to match the behavior of hipGetLastError with its corresponding CUDA API.
-* New HIP APIs
-  - The `_sync()` version of crosslane builtins such as `shfl_sync()`,
-    `__all_sync()` and `__any_sync()`, are enabled by default. These can be
-    disabled by setting the preprocessor macro `HIP_DISABLE_WARP_SYNC_BUILTINS`.
 
 ## HIP 6.3 for ROCm 6.3
 

hipamd/hip-config-amd.cmake.in

@@ -66,9 +66,16 @@ if(NOT WIN32)
   find_dependency(AMDDeviceLibs HINTS ${ROCM_PATH} PATHS "/opt/rocm")
 endif()
 
-# If AMDGPU_TARGETS is not defined by the app, amdgpu-arch is run to find the gpu archs
+if(DEFINED AMDGPU_TARGETS AND NOT DEFINED GPU_TARGETS)
+  message(AUTHOR_WARNING "AMDGPU_TARGETS is deprecated. Please use GPU_TARGETS instead.")
+
+  # Set GPU_TARGETS to the value of AMDGPU_TARGETS
+  set(GPU_TARGETS "${AMDGPU_TARGETS}")
+endif()
+
+# If GPU_TARGETS is not defined by the app, amdgpu-arch is run to find the gpu archs
 # of all the devices present in the machine
-if(NOT AMDGPU_TARGETS)
+if(NOT GPU_TARGETS)
   if(@BUILD_SHARED_LIBS@)
     if (WIN32)
       set(AMDGPU_ARCH "${HIP_CLANG_ROOT}/bin/amdgpu-arch.exe")
@@ -89,7 +96,7 @@ if(NOT AMDGPU_TARGETS)
 
   if(AMDGPU_ARCH_ERROR)
     message(AUTHOR_WARNING
-        " AMDGPU_TARGETS was not set, and system GPU detection was unsuccsesful.\n \n"
+        " GPU_TARGETS was not set, and system GPU detection was unsuccsesful.\n \n"
         " The amdgpu-arch tool failed:\n"
         " Error: '${AMDGPU_ARCH_ERROR}'\n"
         " Output: '${AMDGPU_ARCH_OUTPUT}'\n \n"
@@ -102,17 +109,17 @@ if(NOT AMDGPU_TARGETS)
     string(REPLACE "gfx000\n" "" AMDGPU_ARCH_OUTPUT "${AMDGPU_ARCH_OUTPUT}")
     if (NOT AMDGPU_ARCH_OUTPUT STREQUAL "")
       string(REPLACE "\n" ";" AMDGPU_ARCH_OUTPUT ${AMDGPU_ARCH_OUTPUT})
-      set(AMDGPU_TARGETS ${AMDGPU_ARCH_OUTPUT} CACHE STRING "AMD GPU targets to compile for")
+      set(GPU_TARGETS ${AMDGPU_ARCH_OUTPUT} CACHE STRING "AMD GPU targets to compile for")
     endif()
   endif()
 endif()
 
-if (NOT AMDGPU_TARGETS AND NOT @BUILD_SHARED_LIBS@)
+if (NOT GPU_TARGETS AND NOT @BUILD_SHARED_LIBS@)
   # The default architecture is gfx942 for static build
-  set(AMDGPU_TARGETS "gfx942" CACHE STRING "AMD GPU targets to compile for")
+  set(GPU_TARGETS "gfx942" CACHE STRING "AMD GPU targets to compile for")
 endif()
 
-set(GPU_TARGETS "${AMDGPU_TARGETS}" CACHE STRING "GPU targets to compile for")
+set(GPU_BUILD_TARGETS "${GPU_TARGETS}" CACHE STRING "GPU targets to compile for")
 if(NOT WIN32)
   find_dependency(amd_comgr HINTS ${ROCM_PATH} PATHS "/opt/rocm")
 endif()
@@ -153,7 +160,7 @@ endif()
 
 hip_add_interface_link_flags(hip::device --hip-link)
 
-foreach(GPU_TARGET ${GPU_TARGETS})
+foreach(GPU_TARGET ${GPU_BUILD_TARGETS})
     if (NOT compilePropIsSet)
       hip_add_interface_compile_flags(hip::device --offload-arch=${GPU_TARGET})
     endif()

hipamd/include/hip/amd_detail/amd_hip_bf16.h

@@ -681,7 +681,7 @@ __BF16_HOST_DEVICE_STATIC__ __hip_bfloat16 __ushort_as_bfloat16(const unsigned s
   return u.bf16;
 }
 
-#if !defined(HIP_DISABLE_WARP_SYNC_BUILTINS)
+#ifdef HIP_ENABLE_WARP_SYNC_BUILTINS
 /**
  * \ingroup HIP_INTRINSIC_BFLOAT16_MOVE
  * \brief shfl down warp intrinsic for bfloat16
@@ -789,7 +789,7 @@ __BF16_DEVICE_STATIC__ __hip_bfloat162 __shfl_xor_sync(const unsigned long long
   u.ui = __shfl_xor_sync(mask, u.ui, delta, width);
   return u.bf162;
 }
-#endif // HIP_DISABLE_WARP_SYNC_BUILTINS
+#endif
 
 /**
  * \ingroup HIP_INTRINSIC_BFLOAT16_ARITH

hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h

@@ -473,7 +473,7 @@ class coalesced_group : public thread_group {
 
     return __shfl(var, lane, __AMDGCN_WAVEFRONT_SIZE);
   }
-#if !defined(HIP_DISABLE_WARP_SYNC_BUILTINS)
+#ifdef HIP_ENABLE_WARP_SYNC_BUILTINS
    __CG_QUALIFIER__ unsigned long long ballot(int pred) const {
      return internal::helper::adjust_mask(
          coalesced_info.member_mask,
@@ -500,7 +500,7 @@ class coalesced_group : public thread_group {
          __match_all_sync(static_cast<unsigned long long>(coalesced_info.member_mask), value,
                           &pred));
    }
-#endif // HIP_DISABLE_WARP_SYNC_BUILTINS
+#endif
 };
 
 /** \brief   User exposed API to create coalesced groups.
@@ -665,12 +665,12 @@ template <unsigned int size> class thread_block_tile_base : public tile_base<siz
   friend __CG_QUALIFIER__ coalesced_group
   binary_partition(const thread_block_tile<fsize, fparent>& tgrp, bool pred);
 
-#if !defined(HIP_DISABLE_WARP_SYNC_BUILTINS)
+#ifdef HIP_ENABLE_WARP_SYNC_BUILTINS
   __CG_QUALIFIER__ unsigned long long build_mask() const {
     unsigned long long mask = ~0ull >> (64 - numThreads);
     return mask << ((internal::workgroup::thread_rank() / numThreads) * numThreads);
   }
-#endif // HIP_DISABLE_WARP_SYNC_BUILTINS
+#endif
 
  public:
   __CG_STATIC_QUALIFIER__ void sync() {
@@ -697,7 +697,7 @@ template <unsigned int size> class thread_block_tile_base : public tile_base<siz
     return (__shfl_xor(var, laneMask, numThreads));
   }
 
-#if !defined(HIP_DISABLE_WARP_SYNC_BUILTINS)
+#ifdef HIP_ENABLE_WARP_SYNC_BUILTINS
   __CG_QUALIFIER__ unsigned long long ballot(int pred) const {
     const auto mask = build_mask();
     return internal::helper::adjust_mask(mask, __ballot_sync(mask, pred));
@@ -716,7 +716,7 @@ template <unsigned int size> class thread_block_tile_base : public tile_base<siz
     const auto mask = build_mask();
     return internal::helper::adjust_mask(mask, __match_all_sync(mask, value, &pred));
   }
-#endif // HIP_DISABLE_WARP_SYNC_BUILTINS
+#endif
 };
 
 /** \brief   User exposed API that captures the state of the parent group pre-partition
@@ -901,7 +901,7 @@ __CG_QUALIFIER__ thread_block_tile<size, ParentCGTy> tiled_partition(const Paren
   return impl::tiled_partition_internal<size, ParentCGTy>(g);
 }
 
-#if !defined(HIP_DISABLE_WARP_SYNC_BUILTINS)
+#ifdef HIP_ENABLE_WARP_SYNC_BUILTINS
 /** \brief Binary partition
  *
  *  \details This splits the input thread group into two partitions determined by predicate
@@ -927,7 +927,7 @@ __CG_QUALIFIER__ coalesced_group binary_partition(const thread_block_tile<size,
     return coalesced_group(tgrp.build_mask() ^ mask);
   }
 }
-#endif // HIP_DISABLE_WARP_SYNC_BUILTINS
+#endif
 }  // namespace cooperative_groups
 
 #endif  // __cplusplus