[HIP] Diagnose compiling kernel without offload arch

[yaxunl]

Summary

AMDGPU does not have a common processor (GPU arch). A HIP kernel
must be compiled with a specified processor to be able to be launched
on that processor.

However we cannot simply diagnose missing --offload-arch in clang
driver since a valid HIP program can contain no kernel, which can be
compiled without specifying offload arch and executed on machines
without AMDGPU.

Therefore only HIP programs containing kernels should be diagnosed
when compiled without offload arch.

This patch changes clang driver so that when offload arch is not specified
for HIP, no target CPU is specified for clang -cc1. If HIP program contains
kernel, FE will diagnose it as a fatal error so that the diagnostics will be
emitted only once. This way, we allow HIP programs without kernels to
be compiled without offload arch whereas forbid HIP programs with
kernels to be compiled without offload arch.

[aaron.ballman]

Drive-by comment on the diagnostic wording.

clang/include/clang/Basic/DiagnosticSemaKinds.td

8260

[yaxunl]

revised error msg by Aaron's comments

[tra]

Enforcing explicit GPU target makes sense.

However, I think that singling out a __global__ as the trigger is not sufficient for the intended purpose.

If we can't generate a usable GPU-side binary, then we should produce an error if we need to generate *anything* during GPU-side compilation.
Using __global__ as a proxy would not catch some use cases and, possibly, will produce false positives in others.

E.g. what if I have a TU which only has a __device__ int var = 42; combined with a host-side code to memcpy to/from it? It would still be a valid, if not very useful code, but it would still suffer from runtime being unable to load it on a GPU unless that variable is in a GPU binary compiled with a valid target.

__device__ functions in TUs compiled with -fgpu-rdc would have a similar problem. They would eventually be linked into a GPU binary which will be useless if it's not compiled for correct GPU. Granted, __device__ functions will eventually need to be called from a kernel, so we will error out on a __global__ *somewhere*, but it will miss the problem when such TU does not get to the linking stage (e.g. maybe the user wants to link them at runtime).

clang/include/clang/Basic/DiagnosticSemaKinds.td

8260

How about compiling a file with __device__ functions with -fgpu-rdc? If a kernel with no-arch is an error, then this should be an error, too.

clang/lib/Sema/SemaDeclAttr.cpp

4431

Will this fire if we have an uninstantiated kernel template?

clang/test/SemaCUDA/kernel-no-gpu.cu

7

We'll need few more test cases.

E.g. these should be fine.

template <typename T> __global__ void kernel(T arg ) {};
__global__ void kernel(T arg );

clang/include/clang/Basic/DiagnosticSemaKinds.td

 def err_capture_bad_target_this_ptr : Error<
   "capture host side class data member by this pointer in device or host device lambda function">;
 def warn_kern_is_method : Extension<
   "kernel function %0 is a member function; this may not be accepted by nvcc">,
   InGroup<CudaCompat>;
 def warn_kern_is_inline : Warning<
   "ignored 'inline' attribute on kernel function %0">,
   InGroup<CudaCompat>;
+def err_hip_kern_without_gpu : Error<
+  "compiling a HIP kernel without specifying an offload arch is not allowed">,
+  DefaultFatal;
 def err_variadic_device_fn : Error<
   "CUDA device code does not support variadic functions">;
 def err_va_arg_in_device : Error<
   "CUDA device code does not support va_arg">;
 def err_alias_not_supported_on_nvptx : Error<"CUDA does not support aliases">;
 def err_cuda_unattributed_constexpr_cannot_overload_device : Error<
   "constexpr function %0 without __host__ or __device__ attributes cannot "
   "overload __device__ function with same signature.  Add a __host__ "

clang/lib/Driver/ToolChains/HIP.cpp

HIPToolChain::TranslateArgs(const llvm::opt::DerivedArgList &Args,
 
   for (Arg *A : Args) {
     if (!shouldSkipArgument(A))
       DAL->append(A);
   }
 
   if (!BoundArch.empty()) {
     DAL->eraseArg(options::OPT_mcpu_EQ);
-    DAL->AddJoinedArg(nullptr, Opts.getOption(options::OPT_mcpu_EQ), BoundArch);
+    if (BoundArch != "unknown") {
+      DAL->AddJoinedArg(nullptr, Opts.getOption(options::OPT_mcpu_EQ),
+                        BoundArch);
       checkTargetID(*DAL);
     }
+  }
 
   return DAL;
 }
 
 Tool *HIPToolChain::buildLinker() const {
   assert(getTriple().getArch() == llvm::Triple::amdgcn);
   return new tools::AMDGCN::Linker(*this);
 }

 VersionTuple HIPToolChain::computeMSVCVersion(const Driver *D,
                                                const ArgList &Args) const {
   return HostTC.computeMSVCVersion(D, Args);
 }
 
 llvm::SmallVector<std::string, 12>
 HIPToolChain::getHIPDeviceLibs(const llvm::opt::ArgList &DriverArgs) const {
   llvm::SmallVector<std::string, 12> BCLibs;
-  if (DriverArgs.hasArg(options::OPT_nogpulib))
+  StringRef GpuArch = getGPUArch(DriverArgs);
+  if (DriverArgs.hasArg(options::OPT_nogpulib) || GpuArch.empty())
     return {};
   ArgStringList LibraryPaths;
 
   // Find in --hip-device-lib-path and HIP_LIBRARY_PATH.
   for (auto Path : RocmInstallation.getRocmDeviceLibPathArg())
     LibraryPaths.push_back(DriverArgs.MakeArgString(Path));
 
   addDirectoryList(DriverArgs, LibraryPaths, "", "HIP_DEVICE_LIB_PATH");
llvm::for_each(BCLibArgs, [&](StringRef BCName) {
       }
       getDriver().Diag(diag::err_drv_no_such_file) << BCName;
     });
   } else {
     if (!RocmInstallation.hasDeviceLibrary()) {
       getDriver().Diag(diag::err_drv_no_rocm_device_lib) << 0;
       return {};
     }
-    StringRef GpuArch = getGPUArch(DriverArgs);
     assert(!GpuArch.empty() && "Must have an explicit GPU arch.");
     (void)GpuArch;
     auto Kind = llvm::AMDGPU::parseArchAMDGCN(GpuArch);
     const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind);
 
     std::string LibDeviceFile = RocmInstallation.getLibDeviceFile(CanonArch);
     if (LibDeviceFile.empty()) {
       getDriver().Diag(diag::err_drv_no_rocm_device_lib) << 1 << GpuArch;

clang/lib/Sema/SemaDeclAttr.cpp

static void handleGlobalAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
   if (const auto *Method = dyn_cast<CXXMethodDecl>(FD)) {
     if (Method->isInstance()) {
       S.Diag(Method->getBeginLoc(), diag::err_kern_is_nonstatic_method)
           << Method;
       return;
     }
     S.Diag(Method->getBeginLoc(), diag::warn_kern_is_method) << Method;
   }
+  if (S.getASTContext().getTargetInfo().getTargetOpts().CPU.empty() &&
+      S.getLangOpts().HIP && S.getLangOpts().CUDAIsDevice) {
+    S.Diag(FD->getBeginLoc(), diag::err_hip_kern_without_gpu);
+    return;
+  }
   // Only warn for "inline" when compiling for host, to cut down on noise.
   if (FD->isInlineSpecified() && !S.getLangOpts().CUDAIsDevice)
     S.Diag(FD->getBeginLoc(), diag::warn_kern_is_inline) << FD;
 
   D->addAttr(::new (S.Context) CUDAGlobalAttr(S.Context, AL));
   // In host compilation the kernel is emitted as a stub function, which is
   // a helper function for launching the kernel. The instructions in the helper
   // function has nothing to do with the source code of the kernel. Do not emit

clang/test/CodeGenCUDA/amdgpu-hip-implicit-kernarg.cu

 // REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm -x hip -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx906 -fcuda-is-device -emit-llvm -x hip -o - %s | FileCheck %s
 #include "Inputs/cuda.h"
 
 __global__ void hip_kernel_temp() {
 }
 
 // CHECK: attributes {{.*}} = {{.*}} "amdgpu-implicitarg-num-bytes"="56"

clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu

 // REQUIRES: x86-registered-target
 // REQUIRES: amdgpu-registered-target
 
-// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm -x hip %s -o - | FileCheck --check-prefixes=COMMON,CHECK %s
-// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm -x hip %s -disable-O0-optnone -o - | opt -S -O2 | FileCheck %s --check-prefixes=COMMON,OPT
-// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -x hip %s -o - | FileCheck -check-prefix=HOST %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx906 \
+// RUN:   -fcuda-is-device -emit-llvm -x hip %s -o - \
+// RUN:   | FileCheck --check-prefixes=COMMON,CHECK %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx906 \
+// RUN:   -fcuda-is-device -emit-llvm -x hip %s \
+// RUN:   -disable-O0-optnone -o - | opt -S -O2 \
+// RUN:   | FileCheck %s --check-prefixes=COMMON,OPT
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm \
+// RUN:   -x hip %s -o - | FileCheck -check-prefix=HOST %s
 
 #include "Inputs/cuda.h"
 
 // Coerced struct from `struct S` without all generic pointers lowered into
 // global ones.
 
 // On the host-side compilation, generic pointer won't be coerced.
 // HOST-NOT: %struct.S.coerce

clang/test/CodeGenCUDA/amdgpu-kernel-attrs.cu

-// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx906 \
 // RUN:     -fcuda-is-device -emit-llvm -o - -x hip %s \
 // RUN:     | FileCheck -check-prefixes=CHECK,DEFAULT %s
 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa --gpu-max-threads-per-block=1024 \
-// RUN:     -fcuda-is-device -emit-llvm -o - -x hip %s \
+// RUN:     -target-cpu gfx906 -fcuda-is-device -emit-llvm -o - -x hip %s \
 // RUN:     | FileCheck -check-prefixes=CHECK,MAX1024 %s
 // RUN: %clang_cc1 -triple nvptx \
 // RUN:     -fcuda-is-device -emit-llvm -o - %s | FileCheck %s \
 // RUN:     -check-prefix=NAMD
 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm \
 // RUN:     -verify -o - -x hip %s | FileCheck -check-prefix=NAMD %s
 
 #include "Inputs/cuda.h"

clang/test/CodeGenCUDA/kernel-amdgcn.cu

-// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -emit-llvm -x hip %s -o - | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn -target-cpu gfx906 -fcuda-is-device \
+// RUN:    -emit-llvm -x hip %s -o - | FileCheck %s
 #include "Inputs/cuda.h"
 
 // CHECK: define{{.*}} amdgpu_kernel void @_ZN1A6kernelEv
 class A {
 public:
   static __global__ void kernel(){}
 };
 

clang/test/CodeGenCUDA/kernel-args.cu

 // RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -fcuda-is-device \
-// RUN:     -emit-llvm %s -o - | FileCheck -check-prefix=AMDGCN %s
+// RUN:     -target-cpu gfx906 -emit-llvm %s -o - \
+// RUN:     | FileCheck -check-prefix=AMDGCN %s
 // RUN: %clang_cc1 -x cuda -triple nvptx64-nvidia-cuda- -fcuda-is-device \
 // RUN:     -emit-llvm %s -o - | FileCheck -check-prefix=NVPTX %s
 #include "Inputs/cuda.h"
 
 struct A {
   int a[32];
   float *p;
 };

clang/test/CodeGenCUDA/kernel-dbg-info.cu

 // RUN: echo "GPU binary would be here" > %t
 
 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -O0 \
 // RUN:   -fcuda-include-gpubinary %t -debug-info-kind=limited \
 // RUN:   -o - -x hip | FileCheck -check-prefixes=CHECK,O0 %s
 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm %s -O0 \
 // RUN:   -fcuda-include-gpubinary %t -debug-info-kind=limited \
-// RUN:   -o - -x hip -fcuda-is-device | FileCheck -check-prefix=DEV %s
+// RUN:   -target-cpu gfx906 -o - -x hip -fcuda-is-device \
+// RUN:   | FileCheck -check-prefix=DEV %s
 
 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -O0 \
 // RUN:   -fcuda-include-gpubinary %t -debug-info-kind=limited \
 // RUN:   -o - -x hip -debugger-tuning=gdb -dwarf-version=4 \
 // RUN:   | FileCheck -check-prefixes=CHECK,O0 %s
 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm %s -O0 \
 // RUN:   -fcuda-include-gpubinary %t -debug-info-kind=limited \
 // RUN:   -o - -x hip -debugger-tuning=gdb -dwarf-version=4 \
-// RUN:   -fcuda-is-device | FileCheck -check-prefix=DEV %s
+// RUN:   -target-cpu gfx906 -fcuda-is-device \
+// RUN:   | FileCheck -check-prefix=DEV %s
 
 // RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -O3 \
 // RUN:   -fcuda-include-gpubinary %t -debug-info-kind=limited \
 // RUN:   -o - -x hip -debugger-tuning=gdb -dwarf-version=4 | FileCheck %s
 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm %s -O3 \
 // RUN:   -fcuda-include-gpubinary %t -debug-info-kind=limited \
 // RUN:   -o - -x hip -debugger-tuning=gdb -dwarf-version=4 \
-// RUN:   -fcuda-is-device | FileCheck -check-prefix=DEV %s
+// RUN:   -target-cpu gfx906 -fcuda-is-device \
+// RUN:   | FileCheck -check-prefix=DEV %s
 
 #include "Inputs/cuda.h"
 
 extern "C" __global__ void ckernel(int *a) {
   *a = 1;
 }
 
 // Kernel symbol for launching kernel.

clang/test/CodeGenCUDA/lambda-reference-var.cu

 // RUN: %clang_cc1 -x hip -emit-llvm -std=c++11 %s -o - \
 // RUN:   -triple x86_64-linux-gnu \
 // RUN:   | FileCheck -check-prefix=HOST %s
 // RUN: %clang_cc1 -x hip -emit-llvm -std=c++11 %s -o - \
 // RUN:   -triple amdgcn-amd-amdhsa -fcuda-is-device \
-// RUN:   | FileCheck -check-prefix=DEV %s
+// RUN:   -target-cpu gfx906 | FileCheck -check-prefix=DEV %s
 
 #include "Inputs/cuda.h"
 
 // HOST: %[[T1:.*]] = type <{ i32*, i32, [4 x i8] }>
 // HOST: %[[T2:.*]] = type { i32*, i32** }
 // HOST: %[[T3:.*]] = type <{ i32*, i32, [4 x i8] }>
 // DEV: %[[T1:.*]] = type { i32* }
 // DEV: %[[T2:.*]] = type { i32** }

clang/test/CodeGenCUDA/lambda.cu

 // RUN: %clang_cc1 -x hip -emit-llvm -std=c++11 %s -o - \
 // RUN:   -triple x86_64-linux-gnu \
 // RUN:   | FileCheck -check-prefix=HOST %s
 // RUN: %clang_cc1 -x hip -emit-llvm -std=c++11 %s -o - \
 // RUN:   -triple amdgcn-amd-amdhsa -fcuda-is-device \
-// RUN:   | FileCheck -check-prefix=DEV %s
+// RUN:   -target-cpu gfx906 | FileCheck -check-prefix=DEV %s
 
 #include "Inputs/cuda.h"
 
 // Device side kernel name.
 // HOST: @[[KERN_CAPTURE:[0-9]+]] = {{.*}} c"_Z1gIZ12test_capturevEUlvE_EvT_\00"
 // HOST: @[[KERN_RESOLVE:[0-9]+]] = {{.*}} c"_Z1gIZ12test_resolvevEUlvE_EvT_\00"
 
 // Check functions emitted for test_capture in host compilation.

clang/test/CodeGenCUDA/managed-var.cu

 // REQUIRES: x86-registered-target, amdgpu-registered-target
 
 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \
-// RUN:   -emit-llvm -o - -x hip %s | FileCheck \
+// RUN:   -target-cpu gfx906 -emit-llvm -o - -x hip %s | FileCheck \
 // RUN:   -check-prefixes=COMMON,DEV,NORDC-D %s
 
 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \
-// RUN:   -emit-llvm -fgpu-rdc -cuid=abc -o - -x hip %s > %t.dev
+// RUN:   -target-cpu gfx906 -emit-llvm -fgpu-rdc -cuid=abc -o - -x hip %s > %t.dev
 // RUN: cat %t.dev | FileCheck -check-prefixes=COMMON,DEV,RDC-D %s
 
 // RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \
 // RUN:   -emit-llvm -o - -x hip %s | FileCheck \
 // RUN:   -check-prefixes=COMMON,HOST,NORDC %s
 
 // RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \
 // RUN:   -emit-llvm -fgpu-rdc -cuid=abc -o - -x hip %s > %t.host

clang/test/CodeGenCUDA/norecurse.cu

 // REQUIRES: amdgpu-registered-target
 // REQUIRES: nvptx-registered-target
 
 // RUN: %clang_cc1 -triple nvptx-nvidia-cuda -fcuda-is-device \
 // RUN:     -emit-llvm -disable-llvm-passes -o - %s | FileCheck %s
 
 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
-// RUN:     -emit-llvm -disable-llvm-passes -o - -x hip %s | FileCheck %s
+// RUN:     -target-cpu gfx906 -emit-llvm -disable-llvm-passes \
+// RUN:     -o - -x hip %s | FileCheck %s
 
 #include "Inputs/cuda.h"
 
 __global__ void kernel1(int a) {}
 // CHECK: define{{.*}}@_Z7kernel1i{{.*}}#[[ATTR:[0-9]*]]
 
 // CHECK: attributes #[[ATTR]] = {{.*}}norecurse

clang/test/CodeGenCUDA/static-device-var-no-rdc.cu

 // REQUIRES: x86-registered-target
 // REQUIRES: amdgpu-registered-target
 
 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \
-// RUN:   -emit-llvm -o - -x hip %s | FileCheck \
+// RUN:   -target-cpu gfx906 -emit-llvm -o - -x hip %s | FileCheck \
 // RUN:   -check-prefixes=DEV %s
 
 // RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \
 // RUN:   -emit-llvm -o - -x hip %s | FileCheck \
 // RUN:   -check-prefixes=HOST %s
 
 #include "Inputs/cuda.h"
 

clang/test/CodeGenCUDA/static-device-var-rdc.cu

 // REQUIRES: x86-registered-target
 // REQUIRES: amdgpu-registered-target
 
 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
-// RUN:   -fgpu-rdc -emit-llvm -o - -x hip %s | FileCheck \
+// RUN:   -target-cpu gfx906 -fgpu-rdc -emit-llvm -o - -x hip %s | FileCheck \
 // RUN:   -check-prefixes=DEV,INT-DEV %s
 
 // RUN: %clang_cc1 -triple x86_64-gnu-linux \
 // RUN:   -fgpu-rdc -emit-llvm -o - -x hip %s | FileCheck \
 // RUN:   -check-prefixes=HOST,INT-HOST %s
 
 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -cuid=abc \
-// RUN:   -fgpu-rdc -emit-llvm -o - -x hip %s > %t.dev
+// RUN:   -target-cpu gfx906 -fgpu-rdc -emit-llvm -o - -x hip %s > %t.dev
 // RUN: cat %t.dev | FileCheck -check-prefixes=DEV,EXT-DEV %s
 
 // RUN: %clang_cc1 -triple x86_64-gnu-linux -cuid=abc \
 // RUN:   -fgpu-rdc -emit-llvm -o - -x hip %s > %t.host
 // RUN: cat %t.host | FileCheck -check-prefixes=HOST,EXT-HOST %s
 
 // Check host and device compilations use the same postfixes for static
 // variable names.

clang/test/CodeGenCUDA/unnamed-types.cu

 // RUN: %clang_cc1 -std=c++11 -x hip -triple x86_64-linux-gnu -aux-triple amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --check-prefix=HOST
 // RUN: %clang_cc1 -std=c++11 -x hip -triple x86_64-pc-windows-msvc -aux-triple amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --check-prefix=MSVC
-// RUN: %clang_cc1 -std=c++11 -x hip -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm %s -o - | FileCheck %s --check-prefix=DEVICE
+// RUN: %clang_cc1 -std=c++11 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx906 -fcuda-is-device -emit-llvm %s -o - | FileCheck %s --check-prefix=DEVICE
 
 #include "Inputs/cuda.h"
 
 // HOST: @0 = private unnamed_addr constant [43 x i8] c"_Z2k0IZZ2f1PfENKUlS0_E_clES0_EUlfE_EvS0_T_\00", align 1
 // HOST: @1 = private unnamed_addr constant [60 x i8] c"_Z2k1IZ2f1PfEUlfE_Z2f1S0_EUlffE_Z2f1S0_EUlfE0_EvS0_T_T0_T1_\00", align 1
 // Check that, on MSVC, the same device kernel mangling name is generated.
 // MSVC: @0 = private unnamed_addr constant [43 x i8] c"_Z2k0IZZ2f1PfENKUlS0_E_clES0_EUlfE_EvS0_T_\00", align 1
 // MSVC: @1 = private unnamed_addr constant [60 x i8] c"_Z2k1IZ2f1PfEUlfE_Z2f1S0_EUlffE_Z2f1S0_EUlfE0_EvS0_T_T0_T1_\00", align 1

clang/test/Driver/cuda-flush-denormals-to-zero.cu

 
 // RUN: %clang -x hip -no-canonical-prefixes -### -target x86_64-linux-gnu -c -march=haswell --cuda-gpu-arch=gfx900 -fgpu-flush-denormals-to-zero -nocudainc -nogpulib %s 2>&1 | FileCheck -check-prefix=FTZ %s
 // RUN: %clang -x hip -no-canonical-prefixes -### -target x86_64-linux-gnu -c -march=haswell --cuda-gpu-arch=gfx900 -fno-gpu-flush-denormals-to-zero -nocudainc -nogpulib %s 2>&1 | FileCheck -check-prefix=NOFTZ %s
 
 // Test the default changing with no argument based on the subtarget in HIP mode
 // RUN: %clang -x hip -no-canonical-prefixes -### -target x86_64-linux-gnu -c -march=haswell --cuda-gpu-arch=gfx803 -nocudainc -nogpulib %s 2>&1 | FileCheck -check-prefix=FTZ %s
 // RUN: %clang -x hip -no-canonical-prefixes -### -target x86_64-linux-gnu -c -march=haswell --cuda-gpu-arch=gfx900 -nocudainc -nogpulib %s 2>&1 | FileCheck -check-prefix=NOFTZ %s
 
-// Test no subtarget, which should get the denormal setting of the default gfx803
-// RUN: %clang -x hip -no-canonical-prefixes -### -target x86_64-linux-gnu -c -march=haswell -nocudainc -nogpulib %s 2>&1 | FileCheck -check-prefix=FTZ %s
+// Test no subtarget, which should get the denormal setting of the default
+// CPU of AMDGPU target, which is 'none'.
+// RUN: %clang -x hip -no-canonical-prefixes -### -target x86_64-linux-gnu -c \
+// RUN:   -march=haswell -nocudainc -nogpulib %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=NOFTZ %s
 
 // Test multiple offload archs with different defaults.
 // RUN: %clang -x hip -no-canonical-prefixes -### -target x86_64-linux-gnu -c -march=haswell --cuda-gpu-arch=gfx803 --cuda-gpu-arch=gfx900 -nocudainc -nogpulib %s 2>&1 | FileCheck -check-prefix=MIXED-DEFAULT-MODE %s
 // RUN: %clang -x hip -no-canonical-prefixes -### -target x86_64-linux-gnu -c -march=haswell -fgpu-flush-denormals-to-zero --cuda-gpu-arch=gfx803 --cuda-gpu-arch=gfx900 -nocudainc -nogpulib %s 2>&1 | FileCheck -check-prefix=FTZX2 %s
 // RUN: %clang -x hip -no-canonical-prefixes -### -target x86_64-linux-gnu -c -march=haswell -fno-gpu-flush-denormals-to-zero --cuda-gpu-arch=gfx803 --cuda-gpu-arch=gfx900 -nocudainc -nogpulib %s 2>&1 | FileCheck -check-prefix=NOFTZ %s
 
 
 // CPUFTZ-NOT: -fdenormal-fp-math

clang/test/Driver/hip-default-gpu-arch.hip

 // REQUIRES: clang-driver
 // REQUIRES: x86-registered-target
 // REQUIRES: amdgpu-registered-target
 
 // RUN: %clang -### -c %s 2>&1 | FileCheck %s
 
-// CHECK: {{.*}}clang{{.*}}"-target-cpu" "gfx803"
+// CHECK-NOT: {{.*}}clang{{.*}}"-triple" "amdgcn-amd-amdhsa"{{.*}}"-target-cpu"
+// CHECK: {{.*}}clang{{.*}}"-triple" "amdgcn-amd-amdhsa"

clang/test/SemaCUDA/kernel-no-gpu.cu

+// RUN: %clang_cc1 -fcuda-is-device -verify=hip -x hip %s
+// RUN: %clang_cc1 -fcuda-is-device -verify=cuda %s
+// cuda-no-diagnostics
+
+#include "Inputs/cuda.h"
+
+__global__ void kern1() {}
+// hip-error@-1 {{compiling a HIP kernel without specifying an offload arch is not allowed}}
+
+// Make sure the error is emitted once.
+__global__ void kern2() {}