[clang] f2a78e6 - [AMDGPU] Do not emit arch dependent macros with unspecified cpu (#80035)
via cfe-commits
cfe-commits at lists.llvm.org
Tue Jan 30 11:05:33 PST 2024
Author: Joseph Huber
Date: 2024-01-30T13:05:29-06:00
New Revision: f2a78e68eee53646327f71c475c7f18a28b7f576
URL: https://github.com/llvm/llvm-project/commit/f2a78e68eee53646327f71c475c7f18a28b7f576
DIFF: https://github.com/llvm/llvm-project/commit/f2a78e68eee53646327f71c475c7f18a28b7f576.diff
LOG: [AMDGPU] Do not emit arch dependent macros with unspecified cpu (#80035)
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.
Added:
Modified:
clang/docs/HIPSupport.rst
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/docs/HIPSupport.rst b/clang/docs/HIPSupport.rst
index 84cee45e83ba3..0c93d722ff14e 100644
--- a/clang/docs/HIPSupport.rst
+++ b/clang/docs/HIPSupport.rst
@@ -176,6 +176,10 @@ Predefined Macros
* - ``HIP_API_PER_THREAD_DEFAULT_STREAM``
- Alias to ``__HIP_API_PER_THREAD_DEFAULT_STREAM__``. Deprecated.
+Note that some architecture specific AMDGPU macros will have default values when
+used from the HIP host compilation. Other :doc:`AMDGPU macros <AMDGPUSupport>`
+like ``__AMDGCN_WAVEFRONT_SIZE__`` will default to 64 for example.
+
Compilation Modes
=================
diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp
index 6f3a4908623da..d2ef3dafc6b3e 100644
--- a/clang/lib/Basic/Targets/AMDGPU.cpp
+++ b/clang/lib/Basic/Targets/AMDGPU.cpp
@@ -274,30 +274,30 @@ 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");
- }
+ // Legacy HIP host code relies on these default attributes to be defined.
+ if (GPUKind == llvm::AMDGPU::GK_NONE && !(Opts.HIP && !Opts.CUDAIsDevice))
+ 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()) && Opts.CUDAIsDevice) {
+ 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");
}
}
}
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl
index 53f34c6a44ae7..5875f6fef2f27 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 27c7b4a271fee..1b5c2eaf8e855 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 ----------------
@@ -4316,6 +4323,15 @@
// CHECK_R600_FP64-DAG: #define __R600__ 1
// CHECK_R600_FP64-DAG: #define __HAS_FMAF__ 1
+// Begin HIP host tests -----------
+
+// RUN: %clang -x hip -E -dM %s -o - 2>&1 --offload-host-only \
+// RUN: --offload-arch=gfx803 -target x86_64-unknown-linux \
+// RUN: | FileCheck -match-full-lines %s -check-prefixes=CHECK_HIP_HOST
+// CHECK_HIP_HOST: #define __AMDGCN_WAVEFRONT_SIZE__ 64
+// CHECK_HIP_HOST: #define __AMDGPU__ 1
+// CHECK_HIP_HOST: #define __AMD__ 1
+
// Begin avr tests ----------------
// RUN: %clang --target=avr -mmcu=atmega328 -E -dM %s -o - 2>&1 \
More information about the cfe-commits
mailing list