[clang] [AMDGPU] Do not emit arch dependent macros with unspecified cpu (PR #79660)
Joseph Huber via cfe-commits
cfe-commits at lists.llvm.org
Fri Jan 26 15:42:45 PST 2024
https://github.com/jhuber6 updated https://github.com/llvm/llvm-project/pull/79660
>From ba04b20709cbf76ef6f1490081aecc125bdafec7 Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Fri, 26 Jan 2024 16:25:30 -0600
Subject: [PATCH 1/2] [AMDGPU] Do not emit arch dependent macros with
unspecified cpu
Summary:
Currently, the AMDGPU toolchain accepts not passing `-mcpu` as a means
to create a sort of "generic" IR. The resulting IR will not contain any
target dependent attributes and can then be inserted into another
program via `-mlink-builtin-bitcode` to inherit its attributes.
However, there are a handful of macros that can leak incorrect
information when compiling for an unspecified architecture. Currently,
things like the wavefront size will default to 64, which is actually
variable. We should not expose these macros unless it is known.
---
clang/lib/Basic/Targets/AMDGPU.cpp | 49 ++++++++++---------
.../CodeGenOpenCL/builtins-amdgcn-wave64.cl | 2 +-
.../Preprocessor/predefined-arch-macros.c | 17 +++++--
3 files changed, 38 insertions(+), 30 deletions(-)
diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp
index 6f3a4908623da77..dcdbc1ebe2f570b 100644
--- a/clang/lib/Basic/Targets/AMDGPU.cpp
+++ b/clang/lib/Basic/Targets/AMDGPU.cpp
@@ -274,30 +274,31 @@ void AMDGPUTargetInfo::getTargetDefines(const LangOptions &Opts,
else
Builder.defineMacro("__R600__");
- 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");
- }
+ 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("__"));
+ }
+ 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 53f34c6a44ae7dc..5875f6fef2f2787 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 __AMDGCN_WAVEFRONT_SIZE != 64
+#if defined(__AMDGCN_WAVEFRONT_SIZE__) && __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 27c7b4a271fee80..9879a61825fbbb1 100644
--- a/clang/test/Preprocessor/predefined-arch-macros.c
+++ b/clang/test/Preprocessor/predefined-arch-macros.c
@@ -4294,13 +4294,20 @@
// Begin amdgcn tests ----------------
-// RUN: %clang -march=amdgcn -E -dM %s -o - 2>&1 \
+// 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: -target amdgcn-unknown-unknown \
-// RUN: | FileCheck -match-full-lines %s -check-prefix=CHECK_AMDGCN
+// RUN: | FileCheck -match-full-lines %s -check-prefixes=CHECK_AMDGCN,CHECK_AMDGCN_NONE
// CHECK_AMDGCN: #define __AMDGCN__ 1
-// CHECK_AMDGCN: #define __HAS_FMAF__ 1
-// CHECK_AMDGCN: #define __HAS_FP64__ 1
-// CHECK_AMDGCN: #define __HAS_LDEXPF__ 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__
// Begin r600 tests ----------------
>From 0ee2a139db0f88981abdc20986e949ef90c67c76 Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Fri, 26 Jan 2024 17:42:35 -0600
Subject: [PATCH 2/2] Address comments
---
clang/lib/Basic/Targets/AMDGPU.cpp | 2 --
1 file changed, 2 deletions(-)
diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp
index dcdbc1ebe2f570b..c0cd5dd13d7d533 100644
--- a/clang/lib/Basic/Targets/AMDGPU.cpp
+++ b/clang/lib/Basic/Targets/AMDGPU.cpp
@@ -285,8 +285,6 @@ void AMDGPUTargetInfo::getTargetDefines(const LangOptions &Opts,
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__",
More information about the cfe-commits
mailing list