[clang] [clang][AMDGPU] Don't define feature macros on host code (PR #83558)
Pierre van Houtryve via cfe-commits
cfe-commits at lists.llvm.org
Fri Mar 1 04:56:57 PST 2024
https://github.com/Pierre-vh updated https://github.com/llvm/llvm-project/pull/83558
>From 3730631ac58425f559f4bc3cfe3da89e6367c1c5 Mon Sep 17 00:00:00 2001
From: pvanhout <pierre.vanhoutryve at amd.com>
Date: Fri, 1 Mar 2024 12:43:55 +0100
Subject: [PATCH 1/2] [clang][AMDGPU] Don't define feature macros on host code
Those macros are unreliable because our features are mostly uninitialized at that stage, so any macro we define is unreliable.
Fixes SWDEV-447308
---
clang/lib/Basic/Targets/AMDGPU.cpp | 8 +++++++-
clang/test/Preprocessor/predefined-arch-macros.c | 2 +-
2 files changed, 8 insertions(+), 2 deletions(-)
diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp
index 5742885df0461b..df9a5855068ed3 100644
--- a/clang/lib/Basic/Targets/AMDGPU.cpp
+++ b/clang/lib/Basic/Targets/AMDGPU.cpp
@@ -292,8 +292,14 @@ void AMDGPUTargetInfo::getTargetDefines(const LangOptions &Opts,
}
Builder.defineMacro(Twine("__") + Twine(CanonName) + Twine("__"));
+
+ // Don't emit feature macros in host code because in such cases the
+ // feature list is not accurate.
+ if (IsHIPHost)
+ return;
+
// Emit macros for gfx family e.g. gfx906 -> __GFX9__, gfx1030 -> __GFX10___
- if (isAMDGCN(getTriple()) && !IsHIPHost) {
+ if (isAMDGCN(getTriple())) {
assert(StringRef(CanonName).starts_with("gfx") &&
"Invalid amdgcn canonical name");
StringRef CanonFamilyName = getArchFamilyNameAMDGCN(GPUKind);
diff --git a/clang/test/Preprocessor/predefined-arch-macros.c b/clang/test/Preprocessor/predefined-arch-macros.c
index ca51f2fc22c517..8904bcea1a1f68 100644
--- a/clang/test/Preprocessor/predefined-arch-macros.c
+++ b/clang/test/Preprocessor/predefined-arch-macros.c
@@ -4340,7 +4340,7 @@
// RUN: %clang -x hip -E -dM %s -o - 2>&1 --offload-host-only -nogpulib \
// RUN: -nogpuinc --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-NOT: #define __AMDGCN_WAVEFRONT_SIZE__ 64
// CHECK_HIP_HOST: #define __AMDGPU__ 1
// CHECK_HIP_HOST: #define __AMD__ 1
>From a60d9fa16876b90a69b60de429261a3d10404f7a Mon Sep 17 00:00:00 2001
From: pvanhout <pierre.vanhoutryve at amd.com>
Date: Fri, 1 Mar 2024 13:56:46 +0100
Subject: [PATCH 2/2] use CudaIssDevice
---
clang/lib/Basic/Targets/AMDGPU.cpp | 2 +-
.../CodeGenOpenCL/builtins-amdgcn-wave32.cl | 2 +-
clang/test/Driver/amdgpu-macros.cl | 212 +++++++++---------
clang/test/Driver/target-id-macros.cl | 10 +-
.../Preprocessor/predefined-arch-macros.c | 6 +-
5 files changed, 116 insertions(+), 116 deletions(-)
diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp
index df9a5855068ed3..0c5f6bb13ec2eb 100644
--- a/clang/lib/Basic/Targets/AMDGPU.cpp
+++ b/clang/lib/Basic/Targets/AMDGPU.cpp
@@ -295,7 +295,7 @@ void AMDGPUTargetInfo::getTargetDefines(const LangOptions &Opts,
// Don't emit feature macros in host code because in such cases the
// feature list is not accurate.
- if (IsHIPHost)
+ if (!Opts.CUDAIsDevice)
return;
// Emit macros for gfx family e.g. gfx906 -> __GFX9__, gfx1030 -> __GFX10___
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-wave32.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-wave32.cl
index da1ae244431556..de3020fdb6f98f 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-wave32.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-wave32.cl
@@ -42,6 +42,6 @@ void test_read_exec_hi(global uint* out) {
*out = __builtin_amdgcn_read_exec_hi();
}
-#if __AMDGCN_WAVEFRONT_SIZE != 32
+#if defined(__AMDGCN_WAVEFRONT_SIZE) && __AMDGCN_WAVEFRONT_SIZE != 32
#error Wrong wavesize detected
#endif
diff --git a/clang/test/Driver/amdgpu-macros.cl b/clang/test/Driver/amdgpu-macros.cl
index 004619321b271f..1f03ccc6ab9223 100644
--- a/clang/test/Driver/amdgpu-macros.cl
+++ b/clang/test/Driver/amdgpu-macros.cl
@@ -6,32 +6,32 @@
// R600-based processors.
//
-// RUN: %clang -E -dM -target r600 -mcpu=r600 %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,R600 %s -DCPU=r600
-// RUN: %clang -E -dM -target r600 -mcpu=rv630 %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,R600 %s -DCPU=r600
-// RUN: %clang -E -dM -target r600 -mcpu=rv635 %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,R600 %s -DCPU=r600
-// RUN: %clang -E -dM -target r600 -mcpu=r630 %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,R630 %s -DCPU=r630
-// RUN: %clang -E -dM -target r600 -mcpu=rs780 %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,RS880 %s -DCPU=rs880
-// RUN: %clang -E -dM -target r600 -mcpu=rs880 %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,RS880 %s -DCPU=rs880
-// RUN: %clang -E -dM -target r600 -mcpu=rv610 %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,RS880 %s -DCPU=rs880
-// RUN: %clang -E -dM -target r600 -mcpu=rv620 %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,RS880 %s -DCPU=rs880
-// RUN: %clang -E -dM -target r600 -mcpu=rv670 %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,RV670 %s -DCPU=rv670
-// RUN: %clang -E -dM -target r600 -mcpu=rv710 %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,RV710 %s -DCPU=rv710
-// RUN: %clang -E -dM -target r600 -mcpu=rv730 %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,RV730 %s -DCPU=rv730
-// RUN: %clang -E -dM -target r600 -mcpu=rv740 %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,RV770 %s -DCPU=rv770
-// RUN: %clang -E -dM -target r600 -mcpu=rv770 %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,RV770 %s -DCPU=rv770
-// RUN: %clang -E -dM -target r600 -mcpu=cedar %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,CEDAR %s -DCPU=cedar
-// RUN: %clang -E -dM -target r600 -mcpu=palm %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,CEDAR %s -DCPU=cedar
-// RUN: %clang -E -dM -target r600 -mcpu=cypress %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,CYPRESS %s -DCPU=cypress
-// RUN: %clang -E -dM -target r600 -mcpu=hemlock %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,CYPRESS %s -DCPU=cypress
-// RUN: %clang -E -dM -target r600 -mcpu=juniper %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,JUNIPER %s -DCPU=juniper
-// RUN: %clang -E -dM -target r600 -mcpu=redwood %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,REDWOOD %s -DCPU=redwood
-// RUN: %clang -E -dM -target r600 -mcpu=sumo %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,SUMO %s -DCPU=sumo
-// RUN: %clang -E -dM -target r600 -mcpu=sumo2 %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,SUMO %s -DCPU=sumo
-// RUN: %clang -E -dM -target r600 -mcpu=barts %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,BARTS %s -DCPU=barts
-// RUN: %clang -E -dM -target r600 -mcpu=caicos %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,CAICOS %s -DCPU=caicos
-// RUN: %clang -E -dM -target r600 -mcpu=aruba %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,CAYMAN %s -DCPU=cayman
-// RUN: %clang -E -dM -target r600 -mcpu=cayman %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,CAYMAN %s -DCPU=cayman
-// RUN: %clang -E -dM -target r600 -mcpu=turks %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,TURKS %s -DCPU=turks
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target r600 -mcpu=r600 %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,R600 %s -DCPU=r600
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target r600 -mcpu=rv630 %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,R600 %s -DCPU=r600
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target r600 -mcpu=rv635 %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,R600 %s -DCPU=r600
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target r600 -mcpu=r630 %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,R630 %s -DCPU=r630
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target r600 -mcpu=rs780 %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,RS880 %s -DCPU=rs880
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target r600 -mcpu=rs880 %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,RS880 %s -DCPU=rs880
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target r600 -mcpu=rv610 %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,RS880 %s -DCPU=rs880
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target r600 -mcpu=rv620 %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,RS880 %s -DCPU=rs880
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target r600 -mcpu=rv670 %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,RV670 %s -DCPU=rv670
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target r600 -mcpu=rv710 %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,RV710 %s -DCPU=rv710
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target r600 -mcpu=rv730 %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,RV730 %s -DCPU=rv730
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target r600 -mcpu=rv740 %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,RV770 %s -DCPU=rv770
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target r600 -mcpu=rv770 %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,RV770 %s -DCPU=rv770
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target r600 -mcpu=cedar %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,CEDAR %s -DCPU=cedar
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target r600 -mcpu=palm %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,CEDAR %s -DCPU=cedar
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target r600 -mcpu=cypress %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,CYPRESS %s -DCPU=cypress
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target r600 -mcpu=hemlock %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,CYPRESS %s -DCPU=cypress
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target r600 -mcpu=juniper %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,JUNIPER %s -DCPU=juniper
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target r600 -mcpu=redwood %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,REDWOOD %s -DCPU=redwood
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target r600 -mcpu=sumo %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,SUMO %s -DCPU=sumo
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target r600 -mcpu=sumo2 %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,SUMO %s -DCPU=sumo
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target r600 -mcpu=barts %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,BARTS %s -DCPU=barts
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target r600 -mcpu=caicos %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,CAICOS %s -DCPU=caicos
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target r600 -mcpu=aruba %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,CAYMAN %s -DCPU=cayman
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target r600 -mcpu=cayman %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,CAYMAN %s -DCPU=cayman
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target r600 -mcpu=turks %s 2>&1 | FileCheck --check-prefixes=ARCH-R600,TURKS %s -DCPU=turks
// ARCH-R600-NOT: #define FP_FAST_FMA 1
// ARCH-R600-NOT: #define FP_FAST_FMAF 1
@@ -67,74 +67,74 @@
// AMDGCN-based processors.
//
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx600 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx600 -DFAMILY=GFX6
-// RUN: %clang -E -dM -target amdgcn -mcpu=tahiti %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx600 -DFAMILY=GFX6
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx601 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,SLOW_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx601 -DFAMILY=GFX6
-// RUN: %clang -E -dM -target amdgcn -mcpu=pitcairn %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx601 -DFAMILY=GFX6
-// RUN: %clang -E -dM -target amdgcn -mcpu=verde %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx601 -DFAMILY=GFX6
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx602 %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx602 -DFAMILY=GFX6
-// RUN: %clang -E -dM -target amdgcn -mcpu=hainan %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx602 -DFAMILY=GFX6
-// RUN: %clang -E -dM -target amdgcn -mcpu=oland %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx602 -DFAMILY=GFX6
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx700 %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx700 -DFAMILY=GFX7
-// RUN: %clang -E -dM -target amdgcn -mcpu=kaveri %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx700 -DFAMILY=GFX7
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx701 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx701 -DFAMILY=GFX7
-// RUN: %clang -E -dM -target amdgcn -mcpu=hawaii %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx701 -DFAMILY=GFX7
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx702 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx702 -DFAMILY=GFX7
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx703 %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx703 -DFAMILY=GFX7
-// RUN: %clang -E -dM -target amdgcn -mcpu=kabini %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx703 -DFAMILY=GFX7
-// RUN: %clang -E -dM -target amdgcn -mcpu=mullins %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx703 -DFAMILY=GFX7
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx704 %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx704 -DFAMILY=GFX7
-// RUN: %clang -E -dM -target amdgcn -mcpu=bonaire %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx704 -DFAMILY=GFX7
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx705 %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx705 -DFAMILY=GFX7
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx801 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx801 -DFAMILY=GFX8
-// RUN: %clang -E -dM -target amdgcn -mcpu=carrizo %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx801 -DFAMILY=GFX8
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx802 %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx802 -DFAMILY=GFX8
-// RUN: %clang -E -dM -target amdgcn -mcpu=iceland %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx802 -DFAMILY=GFX8
-// RUN: %clang -E -dM -target amdgcn -mcpu=tonga %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx802 -DFAMILY=GFX8
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx803 %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx803 -DFAMILY=GFX8
-// RUN: %clang -E -dM -target amdgcn -mcpu=fiji %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx803 -DFAMILY=GFX8
-// RUN: %clang -E -dM -target amdgcn -mcpu=polaris10 %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx803 -DFAMILY=GFX8
-// RUN: %clang -E -dM -target amdgcn -mcpu=polaris11 %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx803 -DFAMILY=GFX8
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx805 %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx805 -DFAMILY=GFX8
-// RUN: %clang -E -dM -target amdgcn -mcpu=tongapro %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx805 -DFAMILY=GFX8
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx810 %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx810 -DFAMILY=GFX8
-// RUN: %clang -E -dM -target amdgcn -mcpu=stoney %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx810 -DFAMILY=GFX8
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx900 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx900 -DFAMILY=GFX9
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx902 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx902 -DFAMILY=GFX9
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx904 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx904 -DFAMILY=GFX9
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx906 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx906 -DFAMILY=GFX9
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx908 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx908 -DFAMILY=GFX9
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx908 -munsafe-fp-atomics %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,UNSAFEFPATOMIC %s -DWAVEFRONT_SIZE=64 -DCPU=gfx908 -DFAMILY=GFX9
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx909 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx909 -DFAMILY=GFX9
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx90a %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx90a -DFAMILY=GFX9
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx90c %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx90c -DFAMILY=GFX9
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx940 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx940 -DFAMILY=GFX9
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx941 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx941 -DFAMILY=GFX9
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx942 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx942 -DFAMILY=GFX9
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1010 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1010 -DFAMILY=GFX10
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1011 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1011 -DFAMILY=GFX10
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1012 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1012 -DFAMILY=GFX10
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1013 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1013 -DFAMILY=GFX10
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1030 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1030 -DFAMILY=GFX10
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1031 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1031 -DFAMILY=GFX10
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1032 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1032 -DFAMILY=GFX10
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1033 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1033 -DFAMILY=GFX10
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1034 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1034 -DFAMILY=GFX10
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1035 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1035 -DFAMILY=GFX10
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1036 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1036 -DFAMILY=GFX10
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1100 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1100 -DFAMILY=GFX11
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1101 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1101 -DFAMILY=GFX11
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1102 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1102 -DFAMILY=GFX11
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1103 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1103 -DFAMILY=GFX11
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1150 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1150 -DFAMILY=GFX11
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1151 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1151 -DFAMILY=GFX11
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1200 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1200 -DFAMILY=GFX12
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1201 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1201 -DFAMILY=GFX12
-
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx9-generic %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx9_generic -DFAMILY=GFX9
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx10-1-generic %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx10_1_generic -DFAMILY=GFX10
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx10-3-generic %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx10_3_generic -DFAMILY=GFX10
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx11-generic %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx11_generic -DFAMILY=GFX11
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx600 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx600 -DFAMILY=GFX6
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=tahiti %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx600 -DFAMILY=GFX6
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx601 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,SLOW_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx601 -DFAMILY=GFX6
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=pitcairn %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx601 -DFAMILY=GFX6
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=verde %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx601 -DFAMILY=GFX6
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx602 %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx602 -DFAMILY=GFX6
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=hainan %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx602 -DFAMILY=GFX6
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=oland %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx602 -DFAMILY=GFX6
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx700 %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx700 -DFAMILY=GFX7
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=kaveri %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx700 -DFAMILY=GFX7
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx701 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx701 -DFAMILY=GFX7
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=hawaii %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx701 -DFAMILY=GFX7
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx702 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx702 -DFAMILY=GFX7
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx703 %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx703 -DFAMILY=GFX7
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=kabini %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx703 -DFAMILY=GFX7
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=mullins %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx703 -DFAMILY=GFX7
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx704 %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx704 -DFAMILY=GFX7
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=bonaire %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx704 -DFAMILY=GFX7
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx705 %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx705 -DFAMILY=GFX7
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx801 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx801 -DFAMILY=GFX8
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=carrizo %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx801 -DFAMILY=GFX8
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx802 %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx802 -DFAMILY=GFX8
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=iceland %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx802 -DFAMILY=GFX8
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=tonga %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx802 -DFAMILY=GFX8
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx803 %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx803 -DFAMILY=GFX8
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=fiji %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx803 -DFAMILY=GFX8
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=polaris10 %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx803 -DFAMILY=GFX8
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=polaris11 %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx803 -DFAMILY=GFX8
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx805 %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx805 -DFAMILY=GFX8
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=tongapro %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx805 -DFAMILY=GFX8
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx810 %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx810 -DFAMILY=GFX8
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=stoney %s 2>&1 | FileCheck --check-prefix=ARCH-GCN %s -DWAVEFRONT_SIZE=64 -DCPU=gfx810 -DFAMILY=GFX8
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx900 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx900 -DFAMILY=GFX9
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx902 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx902 -DFAMILY=GFX9
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx904 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx904 -DFAMILY=GFX9
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx906 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx906 -DFAMILY=GFX9
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx908 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx908 -DFAMILY=GFX9
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx908 -munsafe-fp-atomics %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,UNSAFEFPATOMIC %s -DWAVEFRONT_SIZE=64 -DCPU=gfx908 -DFAMILY=GFX9
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx909 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx909 -DFAMILY=GFX9
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx90a %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx90a -DFAMILY=GFX9
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx90c %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx90c -DFAMILY=GFX9
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx940 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx940 -DFAMILY=GFX9
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx941 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx941 -DFAMILY=GFX9
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx942 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx942 -DFAMILY=GFX9
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx1010 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1010 -DFAMILY=GFX10
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx1011 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1011 -DFAMILY=GFX10
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx1012 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1012 -DFAMILY=GFX10
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx1013 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1013 -DFAMILY=GFX10
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx1030 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1030 -DFAMILY=GFX10
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx1031 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1031 -DFAMILY=GFX10
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx1032 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1032 -DFAMILY=GFX10
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx1033 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1033 -DFAMILY=GFX10
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx1034 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1034 -DFAMILY=GFX10
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx1035 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1035 -DFAMILY=GFX10
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx1036 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1036 -DFAMILY=GFX10
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx1100 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1100 -DFAMILY=GFX11
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx1101 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1101 -DFAMILY=GFX11
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx1102 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1102 -DFAMILY=GFX11
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx1103 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1103 -DFAMILY=GFX11
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx1150 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1150 -DFAMILY=GFX11
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx1151 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1151 -DFAMILY=GFX11
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx1200 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1200 -DFAMILY=GFX12
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx1201 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1201 -DFAMILY=GFX12
+
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx9-generic %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx9_generic -DFAMILY=GFX9
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx10-1-generic %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx10_1_generic -DFAMILY=GFX10
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx10-3-generic %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx10_3_generic -DFAMILY=GFX10
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx11-generic %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx11_generic -DFAMILY=GFX11
// ARCH-GCN-DAG: #define FP_FAST_FMA 1
@@ -153,32 +153,32 @@
// ARCH-GCN-DAG: #define __AMDGCN_WAVEFRONT_SIZE [[WAVEFRONT_SIZE]]
// UNSAFEFPATOMIC-DAG: #define __AMDGCN_UNSAFE_FP_ATOMICS__ 1
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx906 -mwavefrontsize64 \
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx906 -mwavefrontsize64 \
// RUN: %s 2>&1 | FileCheck --check-prefix=WAVE64 %s
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1010 -mwavefrontsize64 \
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx1010 -mwavefrontsize64 \
// RUN: %s 2>&1 | FileCheck --check-prefix=WAVE64 %s
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx906 -mwavefrontsize64 \
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx906 -mwavefrontsize64 \
// RUN: -mno-wavefrontsize64 %s 2>&1 | FileCheck --check-prefix=WAVE64 %s
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1010 -mwavefrontsize64 \
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx1010 -mwavefrontsize64 \
// RUN: -mno-wavefrontsize64 %s 2>&1 | FileCheck --check-prefix=WAVE32 %s
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx906 -mno-wavefrontsize64 \
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx906 -mno-wavefrontsize64 \
// RUN: -mwavefrontsize64 %s 2>&1 | FileCheck --check-prefix=WAVE64 %s
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1010 -mno-wavefrontsize64 \
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx1010 -mno-wavefrontsize64 \
// RUN: -mwavefrontsize64 %s 2>&1 | FileCheck --check-prefix=WAVE64 %s
// WAVE64-DAG: #define __AMDGCN_WAVEFRONT_SIZE 64
// WAVE32-DAG: #define __AMDGCN_WAVEFRONT_SIZE 32
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx906 \
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx906 \
// RUN: %s 2>&1 | FileCheck --check-prefix=CUMODE-ON %s
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx906 -mcumode \
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx906 -mcumode \
// RUN: %s 2>&1 | FileCheck --check-prefix=CUMODE-ON %s
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx906 -mno-cumode \
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx906 -mno-cumode \
// RUN: %s 2>&1 | FileCheck --check-prefixes=CUMODE-ON,WARN-CUMODE %s
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1030 \
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx1030 \
// RUN: %s 2>&1 | FileCheck --check-prefix=CUMODE-OFF %s
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1030 -mcumode \
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx1030 -mcumode \
// RUN: %s 2>&1 | FileCheck --check-prefix=CUMODE-ON %s
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1030 -mno-cumode \
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn -mcpu=gfx1030 -mno-cumode \
// RUN: %s 2>&1 | FileCheck --check-prefix=CUMODE-OFF %s
// WARN-CUMODE-DAG: warning: ignoring '-mno-cumode' option as it is not currently supported for processor 'gfx906' [-Woption-ignored]
// CUMODE-ON-DAG: #define __AMDGCN_CUMODE__ 1
diff --git a/clang/test/Driver/target-id-macros.cl b/clang/test/Driver/target-id-macros.cl
index 71c33a5ebff1c3..0a4281719bdc65 100644
--- a/clang/test/Driver/target-id-macros.cl
+++ b/clang/test/Driver/target-id-macros.cl
@@ -1,23 +1,23 @@
// REQUIRES: x86-registered-target
// REQUIRES: amdgpu-registered-target
-// RUN: %clang -E -dM -target amdgcn-amd-amdhsa \
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn-amd-amdhsa \
// RUN: -mcpu=gfx908:xnack+:sramecc- -nogpulib -o - %s 2>&1 \
// RUN: | FileCheck -check-prefixes=PROC,ID1 %s
-// RUN: %clang -E -dM -target amdgcn-amd-amdpal \
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn-amd-amdpal \
// RUN: -mcpu=gfx908:xnack+:sramecc- -nogpulib -o - %s 2>&1 \
// RUN: | FileCheck -check-prefixes=PROC,ID1 %s
-// RUN: %clang -E -dM -target amdgcn--mesa3d \
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn--mesa3d \
// RUN: -mcpu=gfx908:xnack+:sramecc- -nogpulib -o - %s 2>&1 \
// RUN: | FileCheck -check-prefixes=PROC,ID1 %s
-// RUN: %clang -E -dM -target amdgcn-amd-amdhsa \
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn-amd-amdhsa \
// RUN: -mcpu=gfx908 -nogpulib -o - %s 2>&1 \
// RUN: | FileCheck -check-prefixes=PROC,ID2 %s
-// RUN: %clang -E -dM -target amdgcn-amd-amdhsa \
+// RUN: %clang -E -dM -Xclang -fcuda-is-device -target amdgcn-amd-amdhsa \
// RUN: -nogpulib -o - %s 2>&1 \
// RUN: | FileCheck -check-prefixes=NONE %s
diff --git a/clang/test/Preprocessor/predefined-arch-macros.c b/clang/test/Preprocessor/predefined-arch-macros.c
index 8904bcea1a1f68..f6fdf95ea9cb4e 100644
--- a/clang/test/Preprocessor/predefined-arch-macros.c
+++ b/clang/test/Preprocessor/predefined-arch-macros.c
@@ -4306,10 +4306,10 @@
// Begin amdgcn tests ----------------
-// RUN: %clang -mcpu=gfx803 -E -dM %s -o - 2>&1 \
+// RUN: %clang -mcpu=gfx803 -E -dM -Xclang -fcuda-is-device %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: %clang -E -dM -Xclang -fcuda-is-device %s -o - 2>&1 \
// RUN: -target amdgcn-unknown-unknown \
// RUN: | FileCheck -match-full-lines %s -check-prefixes=CHECK_AMDGCN,CHECK_AMDGCN_NONE
// CHECK_AMDGCN: #define __AMDGCN__ 1
@@ -4329,7 +4329,7 @@
// CHECK_R600: #define __R600__ 1
// CHECK_R600-NOT: #define __HAS_FMAF__ 1
-// RUN: %clang -march=amdgcn -mcpu=cypress -E -dM %s -o - 2>&1 \
+// RUN: %clang -Xclang -fcuda-is-device -march=amdgcn -mcpu=cypress -E -dM %s -o - 2>&1 \
// RUN: -target r600-unknown-unknown \
// RUN: | FileCheck -match-full-lines %s -check-prefix=CHECK_R600_FP64
// CHECK_R600_FP64-DAG: #define __R600__ 1
More information about the cfe-commits
mailing list