Revert "[AMDGPU] Do not emit arch dependent macros with unspecified c…

…pu (#79660)" This reverts commit c9a6e99. 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.
llvm · Jan 29, 2024 · 72d4fc1 · 72d4fc1
1 parent f28430d
commit 72d4fc1
Show file tree

Hide file tree

Showing 3 changed files with 30 additions and 36 deletions.
diff --git 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
@@ -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
@@ -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 ----------------