From 94bb45ff6dc30756a84b2bd6ed6ed6d108419be4 Mon Sep 17 00:00:00 2001
From: Satyanvesh Dittakavi <Satyanvesh.Dittakavi@amd.com>
Date: Wed, 14 Apr 2021 03:53:07 -0700
Subject: [PATCH] SWDEV-281504 - Rename structure names to use reserved names
 in HIP headers

HIP headers use few structure names as X, Y, and Z. This causes
compilation issues when the apps use similar names as macros.
Renamed the struct names to use reserved names such as
__X, __Y and __Z

Change-Id: I59416c3734f274e853c87d4856b7e616f6cee5f5


[ROCm/hip commit: b26ef8522ed6c8505cb0a3b2492a22a7da0d6844]
---
 .../hip/include/hip/amd_detail/hip_runtime.h  | 42 +++++++++----------
 1 file changed, 21 insertions(+), 21 deletions(-)

diff --git a/projects/hip/include/hip/amd_detail/hip_runtime.h b/projects/hip/include/hip/amd_detail/hip_runtime.h
index 35318cb262..ca88fda2f0 100644
--- a/projects/hip/include/hip/amd_detail/hip_runtime.h
+++ b/projects/hip/include/hip/amd_detail/hip_runtime.h
@@ -249,13 +249,13 @@ template <typename F>
 struct __HIP_Coordinates {
   using R = decltype(F{}(0));
 
-  struct X { __device__ operator R() const noexcept { return F{}(0); } };
-  struct Y { __device__ operator R() const noexcept { return F{}(1); } };
-  struct Z { __device__ operator R() const noexcept { return F{}(2); } };
+  struct __X { __device__ operator R() const noexcept { return F{}(0); } };
+  struct __Y { __device__ operator R() const noexcept { return F{}(1); } };
+  struct __Z { __device__ operator R() const noexcept { return F{}(2); } };
 
-  static constexpr X x{};
-  static constexpr Y y{};
-  static constexpr Z z{};
+  static constexpr __X x{};
+  static constexpr __Y y{};
+  static constexpr __Z z{};
 #ifdef __cplusplus
   __device__ operator dim3() const { return dim3(x, y, z); }
 #endif
@@ -265,53 +265,53 @@ template <typename F>
 #if !defined(_MSC_VER)
 __attribute__((weak))
 #endif
-constexpr typename __HIP_Coordinates<F>::X __HIP_Coordinates<F>::x;
+constexpr typename __HIP_Coordinates<F>::__X __HIP_Coordinates<F>::x;
 template <typename F>
 #if !defined(_MSC_VER)
 __attribute__((weak))
 #endif
-constexpr typename __HIP_Coordinates<F>::Y __HIP_Coordinates<F>::y;
+constexpr typename __HIP_Coordinates<F>::__Y __HIP_Coordinates<F>::y;
 template <typename F>
 #if !defined(_MSC_VER)
 __attribute__((weak))
 #endif
-constexpr typename __HIP_Coordinates<F>::Z __HIP_Coordinates<F>::z;
+constexpr typename __HIP_Coordinates<F>::__Z __HIP_Coordinates<F>::z;
 
 extern "C" __device__ __attribute__((const)) size_t __ockl_get_global_size(uint);
 inline
 __device__
-std::uint32_t operator*(__HIP_Coordinates<__HIP_GridDim>::X,
-                        __HIP_Coordinates<__HIP_BlockDim>::X) noexcept {
+std::uint32_t operator*(__HIP_Coordinates<__HIP_GridDim>::__X,
+                        __HIP_Coordinates<__HIP_BlockDim>::__X) noexcept {
   return __ockl_get_global_size(0);
 }
 inline
 __device__
-std::uint32_t operator*(__HIP_Coordinates<__HIP_BlockDim>::X,
-                        __HIP_Coordinates<__HIP_GridDim>::X) noexcept {
+std::uint32_t operator*(__HIP_Coordinates<__HIP_BlockDim>::__X,
+                        __HIP_Coordinates<__HIP_GridDim>::__X) noexcept {
   return __ockl_get_global_size(0);
 }
 inline
 __device__
-std::uint32_t operator*(__HIP_Coordinates<__HIP_GridDim>::Y,
-                        __HIP_Coordinates<__HIP_BlockDim>::Y) noexcept {
+std::uint32_t operator*(__HIP_Coordinates<__HIP_GridDim>::__Y,
+                        __HIP_Coordinates<__HIP_BlockDim>::__Y) noexcept {
   return __ockl_get_global_size(1);
 }
 inline
 __device__
-std::uint32_t operator*(__HIP_Coordinates<__HIP_BlockDim>::Y,
-                        __HIP_Coordinates<__HIP_GridDim>::Y) noexcept {
+std::uint32_t operator*(__HIP_Coordinates<__HIP_BlockDim>::__Y,
+                        __HIP_Coordinates<__HIP_GridDim>::__Y) noexcept {
   return __ockl_get_global_size(1);
 }
 inline
 __device__
-std::uint32_t operator*(__HIP_Coordinates<__HIP_GridDim>::Z,
-                        __HIP_Coordinates<__HIP_BlockDim>::Z) noexcept {
+std::uint32_t operator*(__HIP_Coordinates<__HIP_GridDim>::__Z,
+                        __HIP_Coordinates<__HIP_BlockDim>::__Z) noexcept {
   return __ockl_get_global_size(2);
 }
 inline
 __device__
-std::uint32_t operator*(__HIP_Coordinates<__HIP_BlockDim>::Z,
-                        __HIP_Coordinates<__HIP_GridDim>::Z) noexcept {
+std::uint32_t operator*(__HIP_Coordinates<__HIP_BlockDim>::__Z,
+                        __HIP_Coordinates<__HIP_GridDim>::__Z) noexcept {
   return __ockl_get_global_size(2);
 }