[SYCL] Remove sycldevice target triple environment component

clang/lib/Basic/Targets.cpp

@@ -595,45 +595,39 @@ TargetInfo *AllocateTarget(const llvm::Triple &Triple,
     }
 
   case llvm::Triple::spir: {
-    if (Triple.getEnvironment() == llvm::Triple::SYCLDevice) {
-      llvm::Triple HT(Opts.HostTriple);
-      switch (HT.getOS()) {
-      case llvm::Triple::Win32:
-        switch (HT.getEnvironment()) {
-        default: // Assume MSVC for unknown environments
-        case llvm::Triple::MSVC:
-          assert(HT.getArch() == llvm::Triple::x86 &&
-                 "Unsupported host architecture");
-          return new MicrosoftX86_32SPIRTargetInfo(Triple, Opts);
-        }
-      case llvm::Triple::Linux:
-        return new LinuxTargetInfo<SPIR32SYCLDeviceTargetInfo>(Triple, Opts);
-      default:
-        return new SPIR32SYCLDeviceTargetInfo(Triple, Opts);
+    llvm::Triple HT(Opts.HostTriple);
+    switch (HT.getOS()) {
+    case llvm::Triple::Win32:
+      switch (HT.getEnvironment()) {
+      default: // Assume MSVC for unknown environments
+      case llvm::Triple::MSVC:
+        assert(HT.getArch() == llvm::Triple::x86 &&
+               "Unsupported host architecture");
+        return new MicrosoftX86_32SPIRTargetInfo(Triple, Opts);
       }
+    case llvm::Triple::Linux:
+      return new LinuxTargetInfo<SPIR32TargetInfo>(Triple, Opts);
+    default:
+      return new SPIR32TargetInfo(Triple, Opts);
     }
-    return new SPIR32TargetInfo(Triple, Opts);
   }
 
   case llvm::Triple::spir64: {
-    if (Triple.getEnvironment() == llvm::Triple::SYCLDevice) {
-      llvm::Triple HT(Opts.HostTriple);
-      switch (HT.getOS()) {
-      case llvm::Triple::Win32:
-        switch (HT.getEnvironment()) {
-        default: // Assume MSVC for unknown environments
-        case llvm::Triple::MSVC:
-          assert(HT.getArch() == llvm::Triple::x86_64 &&
-                 "Unsupported host architecture");
-          return new MicrosoftX86_64_SPIR64TargetInfo(Triple, Opts);
-        }
-      case llvm::Triple::Linux:
-        return new LinuxTargetInfo<SPIR64SYCLDeviceTargetInfo>(Triple, Opts);
-      default:
-        return new SPIR64SYCLDeviceTargetInfo(Triple, Opts);
+    llvm::Triple HT(Opts.HostTriple);
+    switch (HT.getOS()) {
+    case llvm::Triple::Win32:
+      switch (HT.getEnvironment()) {
+      default: // Assume MSVC for unknown environments
+      case llvm::Triple::MSVC:
+        assert(HT.getArch() == llvm::Triple::x86_64 &&
+               "Unsupported host architecture");
+        return new MicrosoftX86_64_SPIR64TargetInfo(Triple, Opts);
       }
+    case llvm::Triple::Linux:
+      return new LinuxTargetInfo<SPIR64TargetInfo>(Triple, Opts);
+    default:
+      return new SPIR64TargetInfo(Triple, Opts);
     }
-    return new SPIR64TargetInfo(Triple, Opts);
   }
 
   case llvm::Triple::wasm32:

clang/lib/Basic/Targets/SPIR.h

@@ -192,40 +192,13 @@ class LLVM_LIBRARY_VISIBILITY SPIR64TargetInfo : public SPIRTargetInfo {
                         MacroBuilder &Builder) const override;
 };
 
-class LLVM_LIBRARY_VISIBILITY SPIR32SYCLDeviceTargetInfo
-    : public SPIR32TargetInfo {
-public:
-  SPIR32SYCLDeviceTargetInfo(const llvm::Triple &Triple,
-                             const TargetOptions &Opts)
-      : SPIR32TargetInfo(Triple, Opts) {
-    // This is workaround for exception_ptr class.
-    // Exceptions is not allowed in sycl device code but we should be able
-    // to parse host code. So we allow compilation of exception_ptr but
-    // if exceptions are used in device code we should emit a diagnostic.
-    MaxAtomicInlineWidth = 32;
-  }
-};
-
-class LLVM_LIBRARY_VISIBILITY SPIR64SYCLDeviceTargetInfo
-    : public SPIR64TargetInfo {
-public:
-  SPIR64SYCLDeviceTargetInfo(const llvm::Triple &Triple,
-                             const TargetOptions &Opts)
-      : SPIR64TargetInfo(Triple, Opts) {
-    // This is workaround for exception_ptr class.
-    // Exceptions is not allowed in sycl device code but we should be able
-    // to parse host code. So we allow compilation of exception_ptr but
-    // if exceptions are used in device code we should emit a diagnostic.
-    MaxAtomicInlineWidth = 64;
-  }
-};
-
 // x86-32 SPIR Windows target
 class LLVM_LIBRARY_VISIBILITY WindowsX86_32SPIRTargetInfo
-    : public WindowsTargetInfo<SPIR32SYCLDeviceTargetInfo> {
+    : public WindowsTargetInfo<SPIR32TargetInfo> {
 public:
-  WindowsX86_32SPIRTargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts)
-      : WindowsTargetInfo<SPIR32SYCLDeviceTargetInfo>(Triple, Opts) {
+  WindowsX86_32SPIRTargetInfo(const llvm::Triple &Triple,
+                              const TargetOptions &Opts)
+      : WindowsTargetInfo<SPIR32TargetInfo>(Triple, Opts) {
     DoubleAlign = LongLongAlign = 64;
     WCharType = UnsignedShort;
   }
@@ -265,10 +238,11 @@ class LLVM_LIBRARY_VISIBILITY MicrosoftX86_32SPIRTargetInfo
 
 // x86-64 SPIR64 Windows target
 class LLVM_LIBRARY_VISIBILITY WindowsX86_64_SPIR64TargetInfo
-    : public WindowsTargetInfo<SPIR64SYCLDeviceTargetInfo> {
+    : public WindowsTargetInfo<SPIR64TargetInfo> {
 public:
-  WindowsX86_64_SPIR64TargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts)
-      : WindowsTargetInfo<SPIR64SYCLDeviceTargetInfo>(Triple, Opts) {
+  WindowsX86_64_SPIR64TargetInfo(const llvm::Triple &Triple,
+                                 const TargetOptions &Opts)
+      : WindowsTargetInfo<SPIR64TargetInfo>(Triple, Opts) {
     LongWidth = LongAlign = 32;
     DoubleAlign = LongLongAlign = 64;
     IntMaxType = SignedLongLong;

clang/lib/Driver/Driver.cpp

@@ -923,6 +923,18 @@ void Driver::CreateOffloadingDeviceToolChains(Compilation &C,
             continue;
           }
 
+          // Warn about deprecated `sycldevice` environment component.
+          if (TT.getEnvironmentName() == "sycldevice") {
+            Diag(clang::diag::warn_drv_deprecated_arg)
+                << TT.str() << TT.getArchName();
+            // Drop environment component.
+            std::string EffectiveTriple =
+                Twine(TT.getArchName() + "-" + TT.getVendorName() + "-" +
+                      TT.getOSName())
+                    .str();
+            TT.setTriple(EffectiveTriple);
+          }
+
           // Store the current triple so that we can check for duplicates in
           // the following iterations.
           FoundNormalizedTriples[NormalizedName] = Val;
@@ -985,7 +997,7 @@ void Driver::CreateOffloadingDeviceToolChains(Compilation &C,
       else
         SYCLTargetArch = "spir64";
     else if (HasValidSYCLRuntime)
-      // Triple for -fintelfpga is spir64_fpga-unknown-unknown-sycldevice.
+      // Triple for -fintelfpga is spir64_fpga.
       SYCLTargetArch = SYCLfpga ? "spir64_fpga" : "spir64";
     if (!SYCLTargetArch.empty()) {
       UniqueSYCLTriplesVec.push_back(MakeSYCLDeviceTriple(SYCLTargetArch));
@@ -1877,7 +1889,6 @@ llvm::Triple Driver::MakeSYCLDeviceTriple(StringRef TargetArch) const {
     TT.setArchName(TargetArch);
     TT.setVendor(llvm::Triple::UnknownVendor);
     TT.setOS(llvm::Triple::UnknownOS);
-    TT.setEnvironment(llvm::Triple::SYCLDevice);
     return TT;
   }
   return llvm::Triple(TargetArch);
@@ -2759,7 +2770,6 @@ bool hasFPGABinary(Compilation &C, std::string Object, types::ID Type) {
   TT.setArchName(types::getTypeName(Type));
   TT.setVendorName("intel");
   TT.setOS(llvm::Triple::UnknownOS);
-  TT.setEnvironment(llvm::Triple::SYCLDevice);
 
   // Checking uses -check-section option with the input file, no output
   // file and the target triple being looked for.
@@ -7467,7 +7477,6 @@ const ToolChain &Driver::getToolChain(const ArgList &Args,
         break;
       case llvm::Triple::MSVC:
       case llvm::Triple::UnknownEnvironment:
-      case llvm::Triple::SYCLDevice:
         if (Args.getLastArgValue(options::OPT_fuse_ld_EQ)
                 .startswith_insensitive("bfd"))
           TC = std::make_unique<toolchains::CrossWindowsToolChain>(

clang/lib/Driver/ToolChains/Clang.cpp

@@ -7627,9 +7627,9 @@ void Clang::AddClangCLArgs(const ArgList &Args, types::ID InputType,
                            bool *EmitCodeView) const {
   unsigned RTOptionID = options::OPT__SLASH_MT;
   bool isNVPTX = getToolChain().getTriple().isNVPTX();
-  bool isSYCLDevice =
-      getToolChain().getTriple().getEnvironment() == llvm::Triple::SYCLDevice;
-  bool isSYCL = Args.hasArg(options::OPT_fsycl) || isSYCLDevice;
+  bool isSPIR = getToolChain().getTriple().isSPIR();
+  // FIXME: isSYCL should not be enabled by "isSPIR"
+  bool isSYCL = Args.hasArg(options::OPT_fsycl) || isSPIR;
   // For SYCL Windows, /MD is the default.
   if (isSYCL)
     RTOptionID = options::OPT__SLASH_MD;
@@ -7641,7 +7641,7 @@ void Clang::AddClangCLArgs(const ArgList &Args, types::ID InputType,
 
   if (Arg *A = Args.getLastArg(options::OPT__SLASH_M_Group)) {
     RTOptionID = A->getOption().getID();
-    if (isSYCL && !isSYCLDevice &&
+    if (isSYCL && !isSPIR &&
         (RTOptionID == options::OPT__SLASH_MT ||
          RTOptionID == options::OPT__SLASH_MTd))
       // Use of /MT or /MTd is not supported for SYCL.
@@ -7653,9 +7653,9 @@ void Clang::AddClangCLArgs(const ArgList &Args, types::ID InputType,
   auto addPreDefines = [&](unsigned Defines) {
     if (Defines & addDEBUG)
       CmdArgs.push_back("-D_DEBUG");
-    if (Defines & addMT && !isSYCLDevice)
+    if (Defines & addMT && !isSPIR)
       CmdArgs.push_back("-D_MT");
-    if (Defines & addDLL && !isSYCLDevice)
+    if (Defines & addDLL && !isSPIR)
       CmdArgs.push_back("-D_DLL");
   };
   StringRef FlagForCRT;
@@ -8346,7 +8346,6 @@ void OffloadBundler::ConstructJobMultipleOutputs(
           TT.setArchName(types::getTypeName(InputType));
           TT.setVendorName("intel");
           TT.setOS(getToolChain().getTriple().getOS());
-          TT.setEnvironment(llvm::Triple::SYCLDevice);
           Triples += "sycl-";
           Triples += TT.normalize();
           continue;
@@ -8462,7 +8461,7 @@ void OffloadWrapper::ConstructJob(Compilation &C, const JobAction &JA,
     llvm::Triple TT = getToolChain().getTriple();
     SmallString<128> TargetTripleOpt = TT.getArchName();
     // When wrapping an FPGA device binary, we need to be sure to apply the
-    // appropriate triple that corresponds (fpga_aoc[xr]-intel-<os>-sycldevice)
+    // appropriate triple that corresponds (fpga_aoc[xr]-intel-<os>)
     // to the target triple setting.
     if (TT.getSubArch() == llvm::Triple::SPIRSubArch_fpga &&
         TCArgs.hasArg(options::OPT_fsycl_link_EQ)) {
@@ -8474,7 +8473,6 @@ void OffloadWrapper::ConstructJob(Compilation &C, const JobAction &JA,
         FPGAArch += "_emu";
       TT.setArchName(FPGAArch);
       TT.setVendorName("intel");
-      TT.setEnvironment(llvm::Triple::SYCLDevice);
       TargetTripleOpt = TT.str();
       // When wrapping an FPGA aocx binary to archive, do not emit registration
       // functions
@@ -8724,7 +8722,7 @@ void SPIRVTranslator::ConstructJob(Compilation &C, const JobAction &JA,
 
   TranslatorArgs.push_back("-o");
   TranslatorArgs.push_back(Output.getFilename());
-  if (getToolChain().getTriple().isSYCLDeviceEnvironment()) {
+  if (JA.isDeviceOffloading(Action::OFK_SYCL)) {
     TranslatorArgs.push_back("-spirv-max-version=1.3");
     // TODO: align debug info for FPGA H/W when its SPIR-V consumer is ready
     if (C.getDriver().isFPGAEmulationMode())

clang/lib/Driver/ToolChains/Linux.cpp

@@ -177,8 +177,8 @@ static StringRef getOSLibDir(const llvm::Triple &Triple, const ArgList &Args) {
 
 Linux::Linux(const Driver &D, const llvm::Triple &Triple, const ArgList &Args)
     : Generic_ELF(D, Triple, Args) {
-  // for SYCL device targets, we rely on the host GCC for proper compilation
-  if (Triple.getEnvironment() == llvm::Triple::SYCLDevice) {
+  // For SPIR device target, we rely on the host GCC for proper compilation.
+  if (Triple.isSPIR()) {
     GCCInstallation.init(llvm::Triple(llvm::sys::getProcessTriple()), Args);
   } else {
     GCCInstallation.init(Triple, Args);

clang/test/CodeGenSYCL/accessor-readonly-invalid-lib.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -triple spir64-unknown-unknown -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
 //
 // Test which verifies that readonly attribute is generated for unexpected access mode value.
 

clang/test/CodeGenSYCL/accessor-readonly.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -fsycl-is-device -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -triple spir64-unknown-unknown -fsycl-is-device -emit-llvm %s -o - | FileCheck %s
 
 // Test to check that readonly attribute is applied to accessors with access mode read.
 

clang/test/CodeGenSYCL/accessor_inheritance.cpp

@@ -1,4 +1,4 @@
-// RUN:  %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+// RUN:  %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
 #include "Inputs/sycl.hpp"
 
 struct Base {

clang/test/CodeGenSYCL/accessor_no_alias_property.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -emit-llvm %s -o - | FileCheck %s
 
 // This test validates the behaviour of noalias parameter attribute.
 

clang/test/CodeGenSYCL/address-space-cond-op.cpp

@@ -1,5 +1,5 @@
 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
-// RUN: %clang_cc1 -x c++ -triple spir64-unknown-linux-sycldevice -disable-llvm-passes -fsycl-is-device -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -x c++ -triple spir64-unknown-linux -disable-llvm-passes -fsycl-is-device -emit-llvm %s -o - | FileCheck %s
 
 struct S {
   unsigned short x;

clang/test/CodeGenSYCL/address-space-initializer.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice \
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown \
 // RUN: -emit-llvm -disable-llvm-passes -o - %s | FileCheck %s
 
 // This test checks that data for big constant initializer lists is placed

clang/test/CodeGenSYCL/address-space-new.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
 
 struct SpaceWaster {
   int i, j;

clang/test/CodeGenSYCL/address-space-of-returns.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -fsycl-is-device -disable-llvm-passes -emit-llvm -x c++ %s -o - | FileCheck %s
+// RUN: %clang_cc1 -triple spir64-unknown-linux -fsycl-is-device -disable-llvm-passes -emit-llvm -x c++ %s -o - | FileCheck %s
 
 struct A {
   int B[42];

clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
 void bar(int & Data) {}
 // CHECK-DAG: define {{.*}}spir_func void @[[RAW_REF:[a-zA-Z0-9_]+]](i32 addrspace(4)* align 4 dereferenceable(4) %
 void bar2(int & Data) {}

clang/test/CodeGenSYCL/anonymous_integration_footer.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -triple spir64-unknown-unknown-sycldevice -fsycl-int-footer=%t.h %s -emit-llvm -o %t.ll
+// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -triple spir64-unknown-unknown -fsycl-int-footer=%t.h %s -emit-llvm -o %t.ll
 // RUN: FileCheck -input-file=%t.h %s
 // A test that validates the more complex cases of the specialization-constant
 // integration footer details, basically any situation we can come up with that

clang/test/CodeGenSYCL/anonymous_integration_footer2.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -triple spir64-unknown-unknown-sycldevice -fsycl-int-footer=%t.h %s -emit-llvm -o %t.ll
+// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -triple spir64-unknown-unknown -fsycl-int-footer=%t.h %s -emit-llvm -o %t.ll
 // RUN: FileCheck -input-file=%t.h %s
 // A test that validates the more complex cases of the specialization-constant
 // integration footer details, basically any situation we can come up with that

clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
 
 // This test checks that compiler generates correct kernel wrapper for basic
 // case.

clang/test/CodeGenSYCL/buffer_location.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -emit-llvm %s -o - | FileCheck %s
 
 // CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE15kernel_function{{.*}} !kernel_arg_buffer_location ![[MDBL:[0-9]+]]
 // CHECK: ![[MDBL]] = !{i32 3, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 2, i32 -1, i32 -1, i32 -1, i32 2, i32 -1, i32 -1, i32 -1, i32 -1}

clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -sycl-std=2020 -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -sycl-std=2020 -emit-llvm -o - %s | FileCheck %s
 
 // Tests for IR of [[intel::scheduler_target_fmax_mhz()]], [[intel::num_simd_work_items()]],
 // [[intel::no_global_work_offset()]], [[intel::max_global_work_dim()]], [[sycl::reqd_sub_group_size()]],

clang/test/CodeGenSYCL/const-wg-init.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
 
 #include "Inputs/sycl.hpp"
 

clang/test/CodeGenSYCL/device-functions.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
 
 template <typename T>
 T bar(T arg);

clang/test/CodeGenSYCL/device-indirectly-callable-attr.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
 
 void helper() {}
 

clang/test/CodeGenSYCL/device-variables.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
 
 enum class test_type { value1, value2, value3 };
 

clang/test/CodeGenSYCL/disable_loop_pipelining.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -sycl-std=2020 -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -sycl-std=2020 -emit-llvm -o - %s | FileCheck %s
 
 #include "sycl.hpp"
 

clang/test/CodeGenSYCL/emit-all-decls-crash.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -femit-all-decls -o - | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -femit-all-decls -o - | FileCheck %s
 
 // This should not crash and we should not emit this declaration, even though
 // we have 'emit-all-decls'.

clang/test/CodeGenSYCL/emit-kernel-in-virtual-func.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -emit-llvm %s -o - | FileCheck %s
 
 template <typename name, typename Func>
 __attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {

clang/test/CodeGenSYCL/emit-kernel-in-virtual-func2.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -emit-llvm %s -o - | FileCheck %s
 
 template <typename name, typename Func>
 __attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {