From 3b156d99f9b3ac9a39985631b868d5716e212dff Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Fri, 28 Aug 2026 11:13:15 -0700 Subject: [PATCH 1/4] [clang][SYCL] Remove the Microsoft-specific SPIR/SPIR-V TargetInfo types The six Microsoft* device targets in SPIR.h differ from their Windows* bases only in getTargetDefines(): the two x86-32 ones add _M_IX86, the two x86-64 ones add _M_X64/_M_AMD64, and the two AArch64 ones override getTargetDefines() solely to forward to the base. Those macros are already predefined from the host target via -aux-triple, so the device-side copies are redundant. Requested by Tom Honermann in https://github.com/intel/llvm/pull/22942#discussion_r3786123704 --- clang/lib/Basic/Targets.cpp | 17 ++-- clang/lib/Basic/Targets/SPIR.h | 99 ------------------- .../CodeGenSYCL/windows-msvc-spir64-sse2.cpp | 4 +- 3 files changed, 9 insertions(+), 111 deletions(-) diff --git a/clang/lib/Basic/Targets.cpp b/clang/lib/Basic/Targets.cpp index 7f5d558d166ec..ad49552774fa0 100644 --- a/clang/lib/Basic/Targets.cpp +++ b/clang/lib/Basic/Targets.cpp @@ -710,7 +710,7 @@ std::unique_ptr AllocateTarget(const llvm::Triple &Triple, case llvm::Triple::MSVC: assert(HT.getArch() == llvm::Triple::x86 && "Unsupported host architecture"); - return std::make_unique(Triple, Opts); + return std::make_unique(Triple, Opts); } case llvm::Triple::Linux: return std::make_unique>(Triple, Opts); @@ -730,11 +730,9 @@ std::unique_ptr AllocateTarget(const llvm::Triple &Triple, case llvm::Triple::MSVC: switch (HT.getArch()) { case llvm::Triple::aarch64: - return std::make_unique(Triple, - Opts); + return std::make_unique(Triple, Opts); case llvm::Triple::x86_64: - return std::make_unique(Triple, - Opts); + return std::make_unique(Triple, Opts); default: llvm::report_fatal_error( "Unsupported host architecture (not x86_64 or aarch64)"); @@ -766,7 +764,7 @@ std::unique_ptr AllocateTarget(const llvm::Triple &Triple, case llvm::Triple::MSVC: assert(HT.getArch() == llvm::Triple::x86 && "Unsupported host architecture"); - return std::make_unique(Triple, Opts); + return std::make_unique(Triple, Opts); } default: return std::make_unique(Triple, Opts); @@ -789,11 +787,10 @@ std::unique_ptr AllocateTarget(const llvm::Triple &Triple, case llvm::Triple::MSVC: switch (HT.getArch()) { case llvm::Triple::aarch64: - return std::make_unique(Triple, - Opts); + return std::make_unique(Triple, Opts); case llvm::Triple::x86_64: - return std::make_unique(Triple, - Opts); + return std::make_unique(Triple, + Opts); default: llvm::report_fatal_error( "Unsupported host architecture (not x86_64 or aarch64)"); diff --git a/clang/lib/Basic/Targets/SPIR.h b/clang/lib/Basic/Targets/SPIR.h index 47fd03948b6ac..f6eb56f385836 100644 --- a/clang/lib/Basic/Targets/SPIR.h +++ b/clang/lib/Basic/Targets/SPIR.h @@ -343,26 +343,6 @@ class LLVM_LIBRARY_VISIBILITY WindowsX86_32SPIRTargetInfo } }; -// x86-32 SPIR Windows Visual Studio target -class LLVM_LIBRARY_VISIBILITY MicrosoftX86_32SPIRTargetInfo - : public WindowsX86_32SPIRTargetInfo { -public: - MicrosoftX86_32SPIRTargetInfo(const llvm::Triple &Triple, - const TargetOptions &Opts) - : WindowsX86_32SPIRTargetInfo(Triple, Opts) { - } - - void getTargetDefines(const LangOptions &Opts, - MacroBuilder &Builder) const override { - WindowsX86_32SPIRTargetInfo::getTargetDefines(Opts, Builder); - // The value of the following reflects processor type. - // 300=386, 400=486, 500=Pentium, 600=Blend (default) - // We lost the original triple, so we use the default. - // TBD should we keep these lines? Copied from X86.h. - Builder.defineMacro("_M_IX86", "600"); - } -}; - // x86-64 SPIR64 Windows target class LLVM_LIBRARY_VISIBILITY WindowsX86_64_SPIR64TargetInfo : public WindowsTargetInfo { @@ -400,23 +380,6 @@ class LLVM_LIBRARY_VISIBILITY WindowsX86_64_SPIR64TargetInfo const std::vector &FeaturesVec) const override; }; -// x86-64 SPIR64 Windows Visual Studio target -class LLVM_LIBRARY_VISIBILITY MicrosoftX86_64_SPIR64TargetInfo - : public WindowsX86_64_SPIR64TargetInfo { -public: - MicrosoftX86_64_SPIR64TargetInfo(const llvm::Triple &Triple, - const TargetOptions &Opts) - : WindowsX86_64_SPIR64TargetInfo(Triple, Opts) { - } - - void getTargetDefines(const LangOptions &Opts, - MacroBuilder &Builder) const override { - WindowsX86_64_SPIR64TargetInfo::getTargetDefines(Opts, Builder); - Builder.defineMacro("_M_X64", "100"); - Builder.defineMacro("_M_AMD64", "100"); - } -}; - // ARM64 SPIR64 Windows target class LLVM_LIBRARY_VISIBILITY WindowsARM64_SPIR64TargetInfo : public WindowsTargetInfo { @@ -444,20 +407,6 @@ class LLVM_LIBRARY_VISIBILITY WindowsARM64_SPIR64TargetInfo } }; -// ARM64 SPIR64 Windows Visual Studio target -class LLVM_LIBRARY_VISIBILITY MicrosoftARM64_SPIR64TargetInfo - : public WindowsARM64_SPIR64TargetInfo { -public: - MicrosoftARM64_SPIR64TargetInfo(const llvm::Triple &Triple, - const TargetOptions &Opts) - : WindowsARM64_SPIR64TargetInfo(Triple, Opts) {} - - void getTargetDefines(const LangOptions &Opts, - MacroBuilder &Builder) const override { - WindowsARM64_SPIR64TargetInfo::getTargetDefines(Opts, Builder); - } -}; - class LLVM_LIBRARY_VISIBILITY BaseSPIRVTargetInfo : public BaseSPIRTargetInfo { public: BaseSPIRVTargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts) @@ -615,24 +564,6 @@ class LLVM_LIBRARY_VISIBILITY WindowsX86_32SPIRV32TargetInfo } }; -// x86-32 SPIRV32 Windows Visual Studio target -class LLVM_LIBRARY_VISIBILITY MicrosoftX86_32SPIRV32TargetInfo - : public WindowsX86_32SPIRV32TargetInfo { -public: - MicrosoftX86_32SPIRV32TargetInfo(const llvm::Triple &Triple, - const TargetOptions &Opts) - : WindowsX86_32SPIRV32TargetInfo(Triple, Opts) {} - - void getTargetDefines(const LangOptions &Opts, - MacroBuilder &Builder) const override { - WindowsX86_32SPIRV32TargetInfo::getTargetDefines(Opts, Builder); - // The value of the following reflects processor type. - // 300=386, 400=486, 500=Pentium, 600=Blend (default) - // We lost the original triple, so we use the default. - Builder.defineMacro("_M_IX86", "600"); - } -}; - // x86-64 SPIRV64 Windows target class LLVM_LIBRARY_VISIBILITY WindowsX86_64_SPIRV64TargetInfo : public WindowsTargetInfo { @@ -665,22 +596,6 @@ class LLVM_LIBRARY_VISIBILITY WindowsX86_64_SPIRV64TargetInfo } }; -// x86-64 SPIRV64 Windows Visual Studio target -class LLVM_LIBRARY_VISIBILITY MicrosoftX86_64_SPIRV64TargetInfo - : public WindowsX86_64_SPIRV64TargetInfo { -public: - MicrosoftX86_64_SPIRV64TargetInfo(const llvm::Triple &Triple, - const TargetOptions &Opts) - : WindowsX86_64_SPIRV64TargetInfo(Triple, Opts) {} - - void getTargetDefines(const LangOptions &Opts, - MacroBuilder &Builder) const override { - WindowsX86_64_SPIRV64TargetInfo::getTargetDefines(Opts, Builder); - Builder.defineMacro("_M_X64", "100"); - Builder.defineMacro("_M_AMD64", "100"); - } -}; - // ARM64 SPIRV64 Windows target class LLVM_LIBRARY_VISIBILITY WindowsARM64_SPIRV64TargetInfo : public WindowsTargetInfo { @@ -708,20 +623,6 @@ class LLVM_LIBRARY_VISIBILITY WindowsARM64_SPIRV64TargetInfo } }; -// ARM64 SPIRV64 Windows Visual Studio target -class LLVM_LIBRARY_VISIBILITY MicrosoftARM64_SPIRV64TargetInfo - : public WindowsARM64_SPIRV64TargetInfo { -public: - MicrosoftARM64_SPIRV64TargetInfo(const llvm::Triple &Triple, - const TargetOptions &Opts) - : WindowsARM64_SPIRV64TargetInfo(Triple, Opts) {} - - void getTargetDefines(const LangOptions &Opts, - MacroBuilder &Builder) const override { - WindowsARM64_SPIRV64TargetInfo::getTargetDefines(Opts, Builder); - } -}; - class LLVM_LIBRARY_VISIBILITY SPIRV64AMDGCNTargetInfo final : public BaseSPIRVTargetInfo { public: diff --git a/clang/test/CodeGenSYCL/windows-msvc-spir64-sse2.cpp b/clang/test/CodeGenSYCL/windows-msvc-spir64-sse2.cpp index 021b17a9cf904..742cdba9bd50d 100644 --- a/clang/test/CodeGenSYCL/windows-msvc-spir64-sse2.cpp +++ b/clang/test/CodeGenSYCL/windows-msvc-spir64-sse2.cpp @@ -1,8 +1,8 @@ // RUN: %clang_cc1 -triple spir64-unknown-unknown -aux-triple x86_64-pc-windows-msvc \ // RUN: -fsycl-is-device -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s -// When SYCL device code is compiled with a Windows-MSVC host, the device -// target (spir64) defines _M_X64 so that MSVC STL headers take the x86 +// When SYCL device code is compiled with a Windows-MSVC host, _M_X64 is +// predefined from the host target, so MSVC STL headers take the x86 // intrinsics path. The device target feature set must correspondingly carry // sse/sse2 so that function-level __target__ attributes (e.g. VS2026 // 's [[gnu::target("fma")]] on _Sqr_error_x86_x64_fma) don't strip From 6bad4763b4971b18dc79ad37849d41c48df6e4ac Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Fri, 28 Aug 2026 11:14:39 -0700 Subject: [PATCH 2/4] [clang][NFC] Collapse redundant host environment switches in AllocateTarget With the Microsoft-specific SPIR/SPIR-V TargetInfo types gone, the `default:` and `case llvm::Triple::MSVC:` labels select the same target, so the four `switch (HT.getEnvironment())` blocks no longer distinguish anything. --- clang/lib/Basic/Targets.cpp | 61 +++++++++++++------------------------ 1 file changed, 22 insertions(+), 39 deletions(-) diff --git a/clang/lib/Basic/Targets.cpp b/clang/lib/Basic/Targets.cpp index ad49552774fa0..fe0cf2f2833d8 100644 --- a/clang/lib/Basic/Targets.cpp +++ b/clang/lib/Basic/Targets.cpp @@ -705,13 +705,9 @@ std::unique_ptr AllocateTarget(const llvm::Triple &Triple, 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 std::make_unique(Triple, Opts); - } + assert(HT.getArch() == llvm::Triple::x86 && + "Unsupported host architecture"); + return std::make_unique(Triple, Opts); case llvm::Triple::Linux: return std::make_unique>(Triple, Opts); default: @@ -725,18 +721,14 @@ std::unique_ptr AllocateTarget(const llvm::Triple &Triple, switch (HT.getOS()) { case llvm::Triple::Win32: - switch (HT.getEnvironment()) { - default: // Assume MSVC for unknown environments - case llvm::Triple::MSVC: - switch (HT.getArch()) { - case llvm::Triple::aarch64: - return std::make_unique(Triple, Opts); - case llvm::Triple::x86_64: - return std::make_unique(Triple, Opts); - default: - llvm::report_fatal_error( - "Unsupported host architecture (not x86_64 or aarch64)"); - } + switch (HT.getArch()) { + case llvm::Triple::aarch64: + return std::make_unique(Triple, Opts); + case llvm::Triple::x86_64: + return std::make_unique(Triple, Opts); + default: + llvm::report_fatal_error( + "Unsupported host architecture (not x86_64 or aarch64)"); } case llvm::Triple::Linux: if (IsFPGASubArch) @@ -759,13 +751,9 @@ std::unique_ptr AllocateTarget(const llvm::Triple &Triple, 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 std::make_unique(Triple, Opts); - } + assert(HT.getArch() == llvm::Triple::x86 && + "Unsupported host architecture"); + return std::make_unique(Triple, Opts); default: return std::make_unique(Triple, Opts); } @@ -782,19 +770,14 @@ std::unique_ptr AllocateTarget(const llvm::Triple &Triple, 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: - switch (HT.getArch()) { - case llvm::Triple::aarch64: - return std::make_unique(Triple, Opts); - case llvm::Triple::x86_64: - return std::make_unique(Triple, - Opts); - default: - llvm::report_fatal_error( - "Unsupported host architecture (not x86_64 or aarch64)"); - } + switch (HT.getArch()) { + case llvm::Triple::aarch64: + return std::make_unique(Triple, Opts); + case llvm::Triple::x86_64: + return std::make_unique(Triple, Opts); + default: + llvm::report_fatal_error( + "Unsupported host architecture (not x86_64 or aarch64)"); } default: return std::make_unique(Triple, Opts); From 8e1eb32cbad44cb656ba1ab96c7cd59c3db6ae2a Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Sat, 29 Aug 2026 21:10:05 -0700 Subject: [PATCH 3/4] [clang][NFC] Drop SPIR/SPIR-V Windows ctor assignments duplicating the host BaseSPIRTargetInfo already copies LongWidth/LongAlign, DoubleAlign, LongLongAlign, IntMaxType, Int64Type and WCharType from the host target, and every MSVC host supplies exactly the values these constructors re-assigned. Only SizeType/PtrDiffType/IntPtrType are kept, since the base class does not copy those and the SPIR64/SPIRV64 defaults are wrong for a Windows host. --- clang/lib/Basic/Targets/SPIR.h | 30 ++---------------------------- 1 file changed, 2 insertions(+), 28 deletions(-) diff --git a/clang/lib/Basic/Targets/SPIR.h b/clang/lib/Basic/Targets/SPIR.h index f6eb56f385836..05b496b5fda41 100644 --- a/clang/lib/Basic/Targets/SPIR.h +++ b/clang/lib/Basic/Targets/SPIR.h @@ -325,10 +325,7 @@ class LLVM_LIBRARY_VISIBILITY WindowsX86_32SPIRTargetInfo public: WindowsX86_32SPIRTargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts) - : WindowsTargetInfo(Triple, Opts) { - DoubleAlign = LongLongAlign = 64; - WCharType = UnsignedShort; - } + : WindowsTargetInfo(Triple, Opts) {} BuiltinVaListKind getBuiltinVaListKind() const override { return TargetInfo::CharPtrBuiltinVaList; @@ -350,14 +347,9 @@ class LLVM_LIBRARY_VISIBILITY WindowsX86_64_SPIR64TargetInfo WindowsX86_64_SPIR64TargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts) : WindowsTargetInfo(Triple, Opts) { - LongWidth = LongAlign = 32; - DoubleAlign = LongLongAlign = 64; - IntMaxType = SignedLongLong; - Int64Type = SignedLongLong; SizeType = UnsignedLongLong; PtrDiffType = SignedLongLong; IntPtrType = SignedLongLong; - WCharType = UnsignedShort; } BuiltinVaListKind getBuiltinVaListKind() const override { @@ -387,14 +379,9 @@ class LLVM_LIBRARY_VISIBILITY WindowsARM64_SPIR64TargetInfo WindowsARM64_SPIR64TargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts) : WindowsTargetInfo(Triple, Opts) { - LongWidth = LongAlign = 32; - DoubleAlign = LongLongAlign = 64; - IntMaxType = SignedLongLong; - Int64Type = SignedLongLong; SizeType = UnsignedLongLong; PtrDiffType = SignedLongLong; IntPtrType = SignedLongLong; - WCharType = UnsignedShort; } BuiltinVaListKind getBuiltinVaListKind() const override { @@ -546,10 +533,7 @@ class LLVM_LIBRARY_VISIBILITY WindowsX86_32SPIRV32TargetInfo public: WindowsX86_32SPIRV32TargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts) - : WindowsTargetInfo(Triple, Opts) { - DoubleAlign = LongLongAlign = 64; - WCharType = UnsignedShort; - } + : WindowsTargetInfo(Triple, Opts) {} BuiltinVaListKind getBuiltinVaListKind() const override { return TargetInfo::CharPtrBuiltinVaList; @@ -571,14 +555,9 @@ class LLVM_LIBRARY_VISIBILITY WindowsX86_64_SPIRV64TargetInfo WindowsX86_64_SPIRV64TargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts) : WindowsTargetInfo(Triple, Opts) { - LongWidth = LongAlign = 32; - DoubleAlign = LongLongAlign = 64; - IntMaxType = SignedLongLong; - Int64Type = SignedLongLong; SizeType = UnsignedLongLong; PtrDiffType = SignedLongLong; IntPtrType = SignedLongLong; - WCharType = UnsignedShort; } BuiltinVaListKind getBuiltinVaListKind() const override { @@ -603,14 +582,9 @@ class LLVM_LIBRARY_VISIBILITY WindowsARM64_SPIRV64TargetInfo WindowsARM64_SPIRV64TargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts) : WindowsTargetInfo(Triple, Opts) { - LongWidth = LongAlign = 32; - DoubleAlign = LongLongAlign = 64; - IntMaxType = SignedLongLong; - Int64Type = SignedLongLong; SizeType = UnsignedLongLong; PtrDiffType = SignedLongLong; IntPtrType = SignedLongLong; - WCharType = UnsignedShort; } BuiltinVaListKind getBuiltinVaListKind() const override { From 7e2789352ca678fc6c1c8c2a834598d766962665 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Mon, 31 Aug 2026 08:09:19 -0700 Subject: [PATCH 4/4] [clang][SYCL] Remove the FPGA SPIR TargetInfo type FPGA support has already been removed from the driver: -fintelfpga and -fsycl-targets=spir64_fpga are rejected with err_drv_unsupported_opt_removed, so SPIR64FPGATargetInfo was only reachable through a direct -cc1 -triple spir64_fpga invocation. Follow up for https://github.com/intel/llvm/pull/22942 --- clang/lib/Basic/Targets.cpp | 6 ---- clang/lib/Basic/Targets/SPIR.h | 8 ----- .../sycl-host-intelfpga-bitint.cpp | 24 -------------- .../CodeGenSYCL/sycl-intelfpga-bitint.cpp | 31 ------------------- clang/test/SemaSYCL/sycl-intelfpga.cpp | 18 ----------- 5 files changed, 87 deletions(-) delete mode 100644 clang/test/CodeGenSYCL/sycl-host-intelfpga-bitint.cpp delete mode 100644 clang/test/CodeGenSYCL/sycl-intelfpga-bitint.cpp delete mode 100644 clang/test/SemaSYCL/sycl-intelfpga.cpp diff --git a/clang/lib/Basic/Targets.cpp b/clang/lib/Basic/Targets.cpp index fe0cf2f2833d8..2e10ef6d41dba 100644 --- a/clang/lib/Basic/Targets.cpp +++ b/clang/lib/Basic/Targets.cpp @@ -717,8 +717,6 @@ std::unique_ptr AllocateTarget(const llvm::Triple &Triple, case llvm::Triple::spir64: { llvm::Triple HT(Opts.HostTriple); - bool IsFPGASubArch = Triple.getSubArch() == llvm::Triple::SPIRSubArch_fpga; - switch (HT.getOS()) { case llvm::Triple::Win32: switch (HT.getArch()) { @@ -731,12 +729,8 @@ std::unique_ptr AllocateTarget(const llvm::Triple &Triple, "Unsupported host architecture (not x86_64 or aarch64)"); } case llvm::Triple::Linux: - if (IsFPGASubArch) - return std::make_unique>(Triple, Opts); return std::make_unique>(Triple, Opts); default: - if (IsFPGASubArch) - return std::make_unique(Triple, Opts); return std::make_unique(Triple, Opts); } } diff --git a/clang/lib/Basic/Targets/SPIR.h b/clang/lib/Basic/Targets/SPIR.h index 3275cfd78b308..574af9c7bf42f 100644 --- a/clang/lib/Basic/Targets/SPIR.h +++ b/clang/lib/Basic/Targets/SPIR.h @@ -279,14 +279,6 @@ class LLVM_LIBRARY_VISIBILITY SPIR64TargetInfo : public SPIRTargetInfo { MacroBuilder &Builder) const override; }; -// spir64_fpga target -class LLVM_LIBRARY_VISIBILITY SPIR64FPGATargetInfo : public SPIR64TargetInfo { -public: - SPIR64FPGATargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts) - : SPIR64TargetInfo(Triple, Opts) {} - virtual size_t getMaxBitIntWidth() const override { return 4096; } -}; - // x86-32 SPIR Windows target class LLVM_LIBRARY_VISIBILITY WindowsX86_32SPIRTargetInfo : public WindowsTargetInfo { diff --git a/clang/test/CodeGenSYCL/sycl-host-intelfpga-bitint.cpp b/clang/test/CodeGenSYCL/sycl-host-intelfpga-bitint.cpp deleted file mode 100644 index 521754bb9e402..0000000000000 --- a/clang/test/CodeGenSYCL/sycl-host-intelfpga-bitint.cpp +++ /dev/null @@ -1,24 +0,0 @@ -// RUN: %clang_cc1 -fsycl-is-host -fintelfpga -triple x86_64 -aux-triple spir64_fpga -emit-llvm %s -o - | FileCheck %s - -// This test checks that we generate appropriate code for division -// operations of _BitInts of size greater than 128 bits, since it -// is allowed when -fintelfpga is enabled. The test uses a value of -// 4096, the maximum bitsize that is currently supported. - -// CHECK: define{{.*}} void @_Z3fooDB4096_S_(ptr {{.*}} sret(i4096) align 8 %agg.result, ptr {{.*}} byval(i4096) align 8 %[[ARG1:[0-9]+]], ptr {{.*}} byval(i4096) align 8 %[[ARG2:[0-9]+]]) -signed _BitInt(4096) foo(signed _BitInt(4096) a, signed _BitInt(4096) b) { - // CHECK: %[[VAR_A:a]].addr = alloca i4096, align 8 - // CHECK: %[[VAR_B:b]].addr = alloca i4096, align 8 - // CHECK: %[[VAR_A]] = load i4096, ptr %[[ARG1]], align 8 - // CHECK: %[[VAR_B]] = load i4096, ptr %[[ARG2]], align 8 - // CHECK: store i4096 %[[VAR_A]], ptr %[[VAR_A]].addr, align 8 - // CHECK: store i4096 %[[VAR_B]], ptr %[[VAR_B]].addr, align 8 - // CHECK: %[[TEMP1:[0-9]+]] = load i4096, ptr %[[VAR_A]].addr, align 8 - // CHECK: %[[TEMP2:[0-9]+]] = load i4096, ptr %[[VAR_B]].addr, align 8 - // CHECK: %div = sdiv i4096 %[[TEMP1]], %[[TEMP2]] - // CHECK: store i4096 %div, ptr %agg.result, align 8 - // CHECK: %[[RES:[0-9+]]] = load i4096, ptr %agg.result, align 8 - // CHECK: store i4096 %[[RES]], ptr %agg.result, align 8 - // CHECK: ret void - return a / b; -} diff --git a/clang/test/CodeGenSYCL/sycl-intelfpga-bitint.cpp b/clang/test/CodeGenSYCL/sycl-intelfpga-bitint.cpp deleted file mode 100644 index 56f70767b8804..0000000000000 --- a/clang/test/CodeGenSYCL/sycl-intelfpga-bitint.cpp +++ /dev/null @@ -1,31 +0,0 @@ -// RUN: %clang_cc1 -fsycl-is-device -fintelfpga -triple spir64_fpga -aux-triple x86_64-unknown-linux-gnu -IInputs -emit-llvm %s -o - | FileCheck %s - -// This test checks that we generate appropriate code for division -// operations of _BitInts of size greater than 128 bits, since it -// is allowed when -fintelfpga is enabled. The test uses a value of -// 4096 for bitint size, the maximum that is currently supported. - -#include "Inputs/sycl.hpp" - -// CHECK: define{{.*}} void @_Z3fooDB4096_S_(ptr {{.*}} sret(i4096) align 8 %agg.result, ptr {{.*}} byval(i4096) align 8 %[[ARG1:[0-9]+]], ptr {{.*}} byval(i4096) align 8 %[[ARG2:[0-9]+]]) -signed _BitInt(4096) foo(signed _BitInt(4096) a, signed _BitInt(4096) b) { - // CHECK: %a.addr.ascast = addrspacecast ptr %a.addr to ptr addrspace(4) - // CHECK: %b.addr.ascast = addrspacecast ptr %b.addr to ptr addrspace(4) - // CHECK: %a = load i4096, ptr %[[ARG1]], align 8 - // CHECK: %b = load i4096, ptr %[[ARG2]], align 8 - // CHECK: store i4096 %a, ptr addrspace(4) %a.addr.ascast, align 8 - // CHECK: store i4096 %b, ptr addrspace(4) %b.addr.ascast, align 8 - // CHECK: %2 = load i4096, ptr addrspace(4) %a.addr.ascast, align 8 - // CHECK: %3 = load i4096, ptr addrspace(4) %b.addr.ascast, align 8 - // CHECK: %div = sdiv i4096 %2, %3 - return a / b; -} - -int main() { - sycl::handler h; - auto lambda = []() { - _BitInt(4096) a, b = 3, c = 4; - a = foo(b, c); - }; - h.single_task(lambda); -} diff --git a/clang/test/SemaSYCL/sycl-intelfpga.cpp b/clang/test/SemaSYCL/sycl-intelfpga.cpp deleted file mode 100644 index 36d9782e38a88..0000000000000 --- a/clang/test/SemaSYCL/sycl-intelfpga.cpp +++ /dev/null @@ -1,18 +0,0 @@ -// RUN: %clang_cc1 -fsycl-is-device -fintelfpga -verify=device-intelfpga -fsyntax-only %s -triple spir64_fpga -// RUN: %clang_cc1 -fsycl-is-device -verify=device -fsyntax-only -triple spir64 %s - -// Tests that we do not issue errors for _Bitints of size greater than 128 -// when -fintelfpga is enabled. The backend is expected to be able to handle -// this, upto a maximum size of 4096. When -fintelfpga is not passed, -// we continue to diagnose size greater than 128. - -// device-intelfpga-error@+2 3{{signed _BitInt of bit sizes greater than 4096 not supported}} -// device-error@+1 3{{signed _BitInt of bit sizes greater than 128 not supported}} -signed _BitInt(4097) foo(signed _BitInt(4097) a, signed _BitInt(4097) b) { - return a / b; -} -// device-error@+2 3{{signed _BitInt of bit sizes greater than 128 not supported}} -// device-intelfpga-no-diagnostic@+1 -signed _BitInt(215) foo(signed _BitInt(215) a, signed _BitInt(215) b) { - return a + b; -}