[clang] 72d4fc1 - Revert "[AMDGPU] Do not emit arch dependent macros with unspecified cpu (#79660)"

[Joseph Huber via cfe-commits]

Author: Joseph Huber
Date: 2024-01-29T11:11:25-06:00
New Revision: 72d4fc1b4d5cfc4f7d50cc5cf1b315543c088f4d

URL: https://github.com/llvm/llvm-project/commit/72d4fc1b4d5cfc4f7d50cc5cf1b315543c088f4d
DIFF: https://github.com/llvm/llvm-project/commit/72d4fc1b4d5cfc4f7d50cc5cf1b315543c088f4d.diff

LOG: Revert "[AMDGPU] Do not emit arch dependent macros with unspecified cpu (#79660)"

This reverts commit c9a6e993f7b349405b6c8f9244cd9cf0f56a6a81.

This breaks HIP code that incorrectly depended on GPU-specific macros to
be set. The code is totally wrong as using `__WAVEFRTONSIZE__` on the
host is absolutely meaningless, but it seems this entire corner of the
toolchain is fundmentally broken. Reverting for now to avoid breakages.

Added: 
    

Modified: 
    clang/lib/Basic/Targets/AMDGPU.cpp
    clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl
    clang/test/Preprocessor/predefined-arch-macros.c

Removed: 
    


################################################################################
diff  --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp
index c0cd5dd13d7d53..6f3a4908623da7 100644
--- a/clang/lib/Basic/Targets/AMDGPU.cpp
+++ b/clang/lib/Basic/Targets/AMDGPU.cpp
@@ -274,29 +274,30 @@ void AMDGPUTargetInfo::getTargetDefines(const LangOptions &Opts,
   else
     Builder.defineMacro("__R600__");
 
-  if (GPUKind == llvm::AMDGPU::GK_NONE)
-    return;
-
-  StringRef CanonName = isAMDGCN(getTriple()) ? getArchNameAMDGCN(GPUKind)
-                                              : getArchNameR600(GPUKind);
-  Builder.defineMacro(Twine("__") + Twine(CanonName) + Twine("__"));
-  // Emit macros for gfx family e.g. gfx906 -> __GFX9__, gfx1030 -> __GFX10___
-  if (isAMDGCN(getTriple())) {
-    assert(CanonName.starts_with("gfx") && "Invalid amdgcn canonical name");
-    Builder.defineMacro(Twine("__") + Twine(CanonName.drop_back(2).upper()) +
-                        Twine("__"));
-    Builder.defineMacro("__amdgcn_processor__",
-                        Twine("\"") + Twine(CanonName) + Twine("\""));
-    Builder.defineMacro("__amdgcn_target_id__",
-                        Twine("\"") + Twine(*getTargetID()) + Twine("\""));
-    for (auto F : getAllPossibleTargetIDFeatures(getTriple(), CanonName)) {
-      auto Loc = OffloadArchFeatures.find(F);
-      if (Loc != OffloadArchFeatures.end()) {
-        std::string NewF = F.str();
-        std::replace(NewF.begin(), NewF.end(), '-', '_');
-        Builder.defineMacro(Twine("__amdgcn_feature_") + Twine(NewF) +
-                                Twine("__"),
-                            Loc->second ? "1" : "0");
+  if (GPUKind != llvm::AMDGPU::GK_NONE) {
+    StringRef CanonName = isAMDGCN(getTriple()) ?
+      getArchNameAMDGCN(GPUKind) : getArchNameR600(GPUKind);
+    Builder.defineMacro(Twine("__") + Twine(CanonName) + Twine("__"));
+    // Emit macros for gfx family e.g. gfx906 -> __GFX9__, gfx1030 -> __GFX10___
+    if (isAMDGCN(getTriple())) {
+      assert(CanonName.starts_with("gfx") && "Invalid amdgcn canonical name");
+      Builder.defineMacro(Twine("__") + Twine(CanonName.drop_back(2).upper()) +
+                          Twine("__"));
+    }
+    if (isAMDGCN(getTriple())) {
+      Builder.defineMacro("__amdgcn_processor__",
+                          Twine("\"") + Twine(CanonName) + Twine("\""));
+      Builder.defineMacro("__amdgcn_target_id__",
+                          Twine("\"") + Twine(*getTargetID()) + Twine("\""));
+      for (auto F : getAllPossibleTargetIDFeatures(getTriple(), CanonName)) {
+        auto Loc = OffloadArchFeatures.find(F);
+        if (Loc != OffloadArchFeatures.end()) {
+          std::string NewF = F.str();
+          std::replace(NewF.begin(), NewF.end(), '-', '_');
+          Builder.defineMacro(Twine("__amdgcn_feature_") + Twine(NewF) +
+                                  Twine("__"),
+                              Loc->second ? "1" : "0");
+        }
       }
     }
   }

diff  --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl
index 5875f6fef2f278..53f34c6a44ae7d 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl
@@ -44,6 +44,6 @@ void test_read_exec_hi(global ulong* out) {
   *out = __builtin_amdgcn_read_exec_hi();
 }
 
-#if defined(__AMDGCN_WAVEFRONT_SIZE__) && __AMDGCN_WAVEFRONT_SIZE__ != 64
+#if __AMDGCN_WAVEFRONT_SIZE != 64
 #error Wrong wavesize detected
 #endif

diff  --git a/clang/test/Preprocessor/predefined-arch-macros.c b/clang/test/Preprocessor/predefined-arch-macros.c
index 9879a61825fbbb..27c7b4a271fee8 100644
--- a/clang/test/Preprocessor/predefined-arch-macros.c
+++ b/clang/test/Preprocessor/predefined-arch-macros.c
@@ -4294,20 +4294,13 @@
 
 // Begin amdgcn tests ----------------
 
-// RUN: %clang -mcpu=gfx803 -E -dM %s -o - 2>&1 \
-// RUN:     -target amdgcn-unknown-unknown \
-// RUN:   | FileCheck -match-full-lines %s -check-prefixes=CHECK_AMDGCN,CHECK_AMDGCN_803
-// RUN: %clang -E -dM %s -o - 2>&1 \
+// RUN: %clang -march=amdgcn -E -dM %s -o - 2>&1 \
 // RUN:     -target amdgcn-unknown-unknown \
-// RUN:   | FileCheck -match-full-lines %s -check-prefixes=CHECK_AMDGCN,CHECK_AMDGCN_NONE
+// RUN:   | FileCheck -match-full-lines %s -check-prefix=CHECK_AMDGCN
 // CHECK_AMDGCN: #define __AMDGCN__ 1
-// CHECK_AMDGCN_803: #define __HAS_FMAF__ 1
-// CHECK_AMDGCN_803: #define __HAS_FP64__ 1
-// CHECK_AMDGCN_803: #define __HAS_LDEXPF__ 1
-// CHECK_AMDGCN_NONE-NOT: #define __HAS_FMAF__
-// CHECK_AMDGCN_NONE-NOT: #define __HAS_FP64__
-// CHECK_AMDGCN_NONE-NOT: #define __HAS_LDEXPF__
-// CHECK_AMDGCN_NONE-NOT: #define __AMDGCN_WAVEFRONT_SIZE__
+// CHECK_AMDGCN: #define __HAS_FMAF__ 1
+// CHECK_AMDGCN: #define __HAS_FP64__ 1
+// CHECK_AMDGCN: #define __HAS_LDEXPF__ 1
 
 // Begin r600 tests ----------------