https://github.com/jhuber6 created https://github.com/llvm/llvm-project/pull/79660
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. >From 63d0328f9b17b833762ad599ee97381158958f3b Mon Sep 17 00:00:00 2001 From: Joseph Huber <hube...@outlook.com> Date: Fri, 26 Jan 2024 16:25:30 -0600 Subject: [PATCH] [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..1fe09c1ce67abc1 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 ---------------- _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits