Skip to content

Commit f2a78e6

Browse files
authored
[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.
1 parent 97d7283 commit f2a78e6

File tree

4 files changed

+50
-30
lines changed

4 files changed

+50
-30
lines changed

clang/docs/HIPSupport.rst

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -176,6 +176,10 @@ Predefined Macros
176176
* - ``HIP_API_PER_THREAD_DEFAULT_STREAM``
177177
- Alias to ``__HIP_API_PER_THREAD_DEFAULT_STREAM__``. Deprecated.
178178

179+
Note that some architecture specific AMDGPU macros will have default values when
180+
used from the HIP host compilation. Other :doc:`AMDGPU macros <AMDGPUSupport>`
181+
like ``__AMDGCN_WAVEFRONT_SIZE__`` will default to 64 for example.
182+
179183
Compilation Modes
180184
=================
181185

clang/lib/Basic/Targets/AMDGPU.cpp

Lines changed: 24 additions & 24 deletions
Original file line numberDiff line numberDiff line change
@@ -274,30 +274,30 @@ void AMDGPUTargetInfo::getTargetDefines(const LangOptions &Opts,
274274
else
275275
Builder.defineMacro("__R600__");
276276

277-
if (GPUKind != llvm::AMDGPU::GK_NONE) {
278-
StringRef CanonName = isAMDGCN(getTriple()) ?
279-
getArchNameAMDGCN(GPUKind) : getArchNameR600(GPUKind);
280-
Builder.defineMacro(Twine("__") + Twine(CanonName) + Twine("__"));
281-
// Emit macros for gfx family e.g. gfx906 -> __GFX9__, gfx1030 -> __GFX10___
282-
if (isAMDGCN(getTriple())) {
283-
assert(CanonName.starts_with("gfx") && "Invalid amdgcn canonical name");
284-
Builder.defineMacro(Twine("__") + Twine(CanonName.drop_back(2).upper()) +
285-
Twine("__"));
286-
}
287-
if (isAMDGCN(getTriple())) {
288-
Builder.defineMacro("__amdgcn_processor__",
289-
Twine("\"") + Twine(CanonName) + Twine("\""));
290-
Builder.defineMacro("__amdgcn_target_id__",
291-
Twine("\"") + Twine(*getTargetID()) + Twine("\""));
292-
for (auto F : getAllPossibleTargetIDFeatures(getTriple(), CanonName)) {
293-
auto Loc = OffloadArchFeatures.find(F);
294-
if (Loc != OffloadArchFeatures.end()) {
295-
std::string NewF = F.str();
296-
std::replace(NewF.begin(), NewF.end(), '-', '_');
297-
Builder.defineMacro(Twine("__amdgcn_feature_") + Twine(NewF) +
298-
Twine("__"),
299-
Loc->second ? "1" : "0");
300-
}
277+
// Legacy HIP host code relies on these default attributes to be defined.
278+
if (GPUKind == llvm::AMDGPU::GK_NONE && !(Opts.HIP && !Opts.CUDAIsDevice))
279+
return;
280+
281+
StringRef CanonName = isAMDGCN(getTriple()) ? getArchNameAMDGCN(GPUKind)
282+
: getArchNameR600(GPUKind);
283+
Builder.defineMacro(Twine("__") + Twine(CanonName) + Twine("__"));
284+
// Emit macros for gfx family e.g. gfx906 -> __GFX9__, gfx1030 -> __GFX10___
285+
if (isAMDGCN(getTriple()) && Opts.CUDAIsDevice) {
286+
assert(CanonName.starts_with("gfx") && "Invalid amdgcn canonical name");
287+
Builder.defineMacro(Twine("__") + Twine(CanonName.drop_back(2).upper()) +
288+
Twine("__"));
289+
Builder.defineMacro("__amdgcn_processor__",
290+
Twine("\"") + Twine(CanonName) + Twine("\""));
291+
Builder.defineMacro("__amdgcn_target_id__",
292+
Twine("\"") + Twine(*getTargetID()) + Twine("\""));
293+
for (auto F : getAllPossibleTargetIDFeatures(getTriple(), CanonName)) {
294+
auto Loc = OffloadArchFeatures.find(F);
295+
if (Loc != OffloadArchFeatures.end()) {
296+
std::string NewF = F.str();
297+
std::replace(NewF.begin(), NewF.end(), '-', '_');
298+
Builder.defineMacro(Twine("__amdgcn_feature_") + Twine(NewF) +
299+
Twine("__"),
300+
Loc->second ? "1" : "0");
301301
}
302302
}
303303
}

clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -44,6 +44,6 @@ void test_read_exec_hi(global ulong* out) {
4444
*out = __builtin_amdgcn_read_exec_hi();
4545
}
4646

47-
#if __AMDGCN_WAVEFRONT_SIZE != 64
47+
#if defined(__AMDGCN_WAVEFRONT_SIZE__) && __AMDGCN_WAVEFRONT_SIZE__ != 64
4848
#error Wrong wavesize detected
4949
#endif

clang/test/Preprocessor/predefined-arch-macros.c

Lines changed: 21 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -4294,13 +4294,20 @@
42944294

42954295
// Begin amdgcn tests ----------------
42964296

4297-
// RUN: %clang -march=amdgcn -E -dM %s -o - 2>&1 \
4297+
// RUN: %clang -mcpu=gfx803 -E -dM %s -o - 2>&1 \
4298+
// RUN: -target amdgcn-unknown-unknown \
4299+
// RUN: | FileCheck -match-full-lines %s -check-prefixes=CHECK_AMDGCN,CHECK_AMDGCN_803
4300+
// RUN: %clang -E -dM %s -o - 2>&1 \
42984301
// RUN: -target amdgcn-unknown-unknown \
4299-
// RUN: | FileCheck -match-full-lines %s -check-prefix=CHECK_AMDGCN
4302+
// RUN: | FileCheck -match-full-lines %s -check-prefixes=CHECK_AMDGCN,CHECK_AMDGCN_NONE
43004303
// CHECK_AMDGCN: #define __AMDGCN__ 1
4301-
// CHECK_AMDGCN: #define __HAS_FMAF__ 1
4302-
// CHECK_AMDGCN: #define __HAS_FP64__ 1
4303-
// CHECK_AMDGCN: #define __HAS_LDEXPF__ 1
4304+
// CHECK_AMDGCN_803: #define __HAS_FMAF__ 1
4305+
// CHECK_AMDGCN_803: #define __HAS_FP64__ 1
4306+
// CHECK_AMDGCN_803: #define __HAS_LDEXPF__ 1
4307+
// CHECK_AMDGCN_NONE-NOT: #define __HAS_FMAF__
4308+
// CHECK_AMDGCN_NONE-NOT: #define __HAS_FP64__
4309+
// CHECK_AMDGCN_NONE-NOT: #define __HAS_LDEXPF__
4310+
// CHECK_AMDGCN_NONE-NOT: #define __AMDGCN_WAVEFRONT_SIZE__
43044311

43054312
// Begin r600 tests ----------------
43064313

@@ -4316,6 +4323,15 @@
43164323
// CHECK_R600_FP64-DAG: #define __R600__ 1
43174324
// CHECK_R600_FP64-DAG: #define __HAS_FMAF__ 1
43184325

4326+
// Begin HIP host tests -----------
4327+
4328+
// RUN: %clang -x hip -E -dM %s -o - 2>&1 --offload-host-only \
4329+
// RUN: --offload-arch=gfx803 -target x86_64-unknown-linux \
4330+
// RUN: | FileCheck -match-full-lines %s -check-prefixes=CHECK_HIP_HOST
4331+
// CHECK_HIP_HOST: #define __AMDGCN_WAVEFRONT_SIZE__ 64
4332+
// CHECK_HIP_HOST: #define __AMDGPU__ 1
4333+
// CHECK_HIP_HOST: #define __AMD__ 1
4334+
43194335
// Begin avr tests ----------------
43204336

43214337
// RUN: %clang --target=avr -mmcu=atmega328 -E -dM %s -o - 2>&1 \

0 commit comments

Comments
 (0)