Skip to content

Commit

Permalink
[AMDGPU] Do not emit arch dependent macros with unspecified cpu (#80035)
Browse files Browse the repository at this point in the history
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.
  • Loading branch information
jhuber6 committed Jan 30, 2024
1 parent 97d7283 commit f2a78e6
Show file tree
Hide file tree
Showing 4 changed files with 50 additions and 30 deletions.
4 changes: 4 additions & 0 deletions clang/docs/HIPSupport.rst
Original file line number Diff line number Diff line change
Expand Up @@ -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
=================

Expand Down
48 changes: 24 additions & 24 deletions clang/lib/Basic/Targets/AMDGPU.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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");
}
}
}
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl
Original file line number Diff line number Diff line change
Expand Up @@ -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
26 changes: 21 additions & 5 deletions clang/test/Preprocessor/predefined-arch-macros.c
Original file line number Diff line number Diff line change
Expand Up @@ -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 ----------------

Expand All @@ -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 \
Expand Down

0 comments on commit f2a78e6

Please sign in to comment.