diff --git a/clang/include/clang/Basic/DiagnosticGroups.td b/clang/include/clang/Basic/DiagnosticGroups.td --- a/clang/include/clang/Basic/DiagnosticGroups.td +++ b/clang/include/clang/Basic/DiagnosticGroups.td @@ -1260,8 +1260,9 @@ def UnknownArgument : DiagGroup<"unknown-argument">; // A warning group for warnings about code that clang accepts when -// compiling OpenCL C/C++ but which is not compatible with the SPIR spec. +// compiling OpenCL C/C++ but which is not compatible with the SPIR(-V) spec. def SpirCompat : DiagGroup<"spir-compat">; +def : DiagGroup<"spirv-compat", [SpirCompat]>; // Alias. // Warning for the GlobalISel options. def GlobalISel : DiagGroup<"global-isel">; diff --git a/clang/lib/Basic/Targets.cpp b/clang/lib/Basic/Targets.cpp --- a/clang/lib/Basic/Targets.cpp +++ b/clang/lib/Basic/Targets.cpp @@ -606,6 +606,18 @@ return nullptr; return new SPIR64TargetInfo(Triple, Opts); } + case llvm::Triple::spirv32: { + if (os != llvm::Triple::UnknownOS || + Triple.getEnvironment() != llvm::Triple::UnknownEnvironment) + return nullptr; + return new SPIRV32TargetInfo(Triple, Opts); + } + case llvm::Triple::spirv64: { + if (os != llvm::Triple::UnknownOS || + Triple.getEnvironment() != llvm::Triple::UnknownEnvironment) + return nullptr; + return new SPIRV64TargetInfo(Triple, Opts); + } case llvm::Triple::wasm32: if (Triple.getSubArch() != llvm::Triple::NoSubArch || Triple.getVendor() != llvm::Triple::UnknownVendor || diff --git a/clang/lib/Basic/Targets/SPIR.h b/clang/lib/Basic/Targets/SPIR.h --- a/clang/lib/Basic/Targets/SPIR.h +++ b/clang/lib/Basic/Targets/SPIR.h @@ -1,4 +1,4 @@ -//===--- SPIR.h - Declare SPIR target feature support -----------*- C++ -*-===// +//===--- SPIR.h - Declare SPIR and SPIR-V target feature support *- C++ -*-===// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// // -// This file declares SPIR TargetInfo objects. +// This file declares SPIR and SPIR-V TargetInfo objects. // //===----------------------------------------------------------------------===// @@ -21,6 +21,7 @@ namespace clang { namespace targets { +// Used by both the SPIR and SPIR-V targets. static const unsigned SPIRDefIsPrivMap[] = { 0, // Default 1, // opencl_global @@ -44,6 +45,7 @@ 0 // ptr64 }; +// Used by both the SPIR and SPIR-V targets. static const unsigned SPIRDefIsGenMap[] = { 4, // Default // OpenCL address space values for this map are dummy and they can't be used @@ -67,14 +69,15 @@ 0 // ptr64 }; -class LLVM_LIBRARY_VISIBILITY SPIRTargetInfo : public TargetInfo { -public: - SPIRTargetInfo(const llvm::Triple &Triple, const TargetOptions &) +// Base class for SPIR and SPIR-V target info. +class LLVM_LIBRARY_VISIBILITY BaseSPIRTargetInfo : public TargetInfo { +protected: + BaseSPIRTargetInfo(const llvm::Triple &Triple, const TargetOptions &) : TargetInfo(Triple) { assert(getTriple().getOS() == llvm::Triple::UnknownOS && - "SPIR target must use unknown OS"); + "SPIR(-V) target must use unknown OS"); assert(getTriple().getEnvironment() == llvm::Triple::UnknownEnvironment && - "SPIR target must use unknown environment type"); + "SPIR(-V) target must use unknown environment type"); TLSSupported = false; VLASupported = false; LongWidth = LongAlign = 64; @@ -87,13 +90,7 @@ NoAsmVariants = true; } - void getTargetDefines(const LangOptions &Opts, - MacroBuilder &Builder) const override; - - bool hasFeature(StringRef Feature) const override { - return Feature == "spir"; - } - +public: // SPIR supports the half type and the only llvm intrinsic allowed in SPIR is // memcpy as per section 3 of the SPIR spec. bool useFP16ConversionIntrinsics() const override { return false; } @@ -149,7 +146,7 @@ void setSupportedOpenCLOpts() override { // Assume all OpenCL extensions and optional core features are supported - // for SPIR since it is a generic target. + // for SPIR and SPIR-V since they are generic targets. supportAllOpenCLOpts(); } @@ -158,6 +155,24 @@ bool hasInt128Type() const override { return false; } }; +class LLVM_LIBRARY_VISIBILITY SPIRTargetInfo : public BaseSPIRTargetInfo { +public: + SPIRTargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts) + : BaseSPIRTargetInfo(Triple, Opts) { + assert(getTriple().getOS() == llvm::Triple::UnknownOS && + "SPIR target must use unknown OS"); + assert(getTriple().getEnvironment() == llvm::Triple::UnknownEnvironment && + "SPIR target must use unknown environment type"); + } + + void getTargetDefines(const LangOptions &Opts, + MacroBuilder &Builder) const override; + + bool hasFeature(StringRef Feature) const override { + return Feature == "spir"; + } +}; + class LLVM_LIBRARY_VISIBILITY SPIR32TargetInfo : public SPIRTargetInfo { public: SPIR32TargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts) @@ -187,6 +202,55 @@ void getTargetDefines(const LangOptions &Opts, MacroBuilder &Builder) const override; }; + +class LLVM_LIBRARY_VISIBILITY SPIRVTargetInfo : public BaseSPIRTargetInfo { +public: + SPIRVTargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts) + : BaseSPIRTargetInfo(Triple, Opts) { + assert(getTriple().getOS() == llvm::Triple::UnknownOS && + "SPIR-V target must use unknown OS"); + assert(getTriple().getEnvironment() == llvm::Triple::UnknownEnvironment && + "SPIR-V target must use unknown environment type"); + } + + void getTargetDefines(const LangOptions &Opts, + MacroBuilder &Builder) const override; + + bool hasFeature(StringRef Feature) const override { + return Feature == "spirv"; + } +}; + +class LLVM_LIBRARY_VISIBILITY SPIRV32TargetInfo : public SPIRVTargetInfo { +public: + SPIRV32TargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts) + : SPIRVTargetInfo(Triple, Opts) { + PointerWidth = PointerAlign = 32; + SizeType = TargetInfo::UnsignedInt; + PtrDiffType = IntPtrType = TargetInfo::SignedInt; + resetDataLayout("e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-" + "v96:128-v192:256-v256:256-v512:512-v1024:1024"); + } + + void getTargetDefines(const LangOptions &Opts, + MacroBuilder &Builder) const override; +}; + +class LLVM_LIBRARY_VISIBILITY SPIRV64TargetInfo : public SPIRVTargetInfo { +public: + SPIRV64TargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts) + : SPIRVTargetInfo(Triple, Opts) { + PointerWidth = PointerAlign = 64; + SizeType = TargetInfo::UnsignedLong; + PtrDiffType = IntPtrType = TargetInfo::SignedLong; + resetDataLayout("e-i64:64-v16:16-v24:32-v32:32-v48:64-" + "v96:128-v192:256-v256:256-v512:512-v1024:1024"); + } + + void getTargetDefines(const LangOptions &Opts, + MacroBuilder &Builder) const override; +}; + } // namespace targets } // namespace clang #endif // LLVM_CLANG_LIB_BASIC_TARGETS_SPIR_H diff --git a/clang/lib/Basic/Targets/SPIR.cpp b/clang/lib/Basic/Targets/SPIR.cpp --- a/clang/lib/Basic/Targets/SPIR.cpp +++ b/clang/lib/Basic/Targets/SPIR.cpp @@ -1,4 +1,4 @@ -//===--- SPIR.cpp - Implement SPIR target feature support -----------------===// +//===--- SPIR.cpp - Implement SPIR and SPIR-V target feature support ------===// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// // -// This file implements SPIR TargetInfo objects. +// This file implements SPIR and SPIR-V TargetInfo objects. // //===----------------------------------------------------------------------===// @@ -32,3 +32,20 @@ SPIRTargetInfo::getTargetDefines(Opts, Builder); DefineStd(Builder, "SPIR64", Opts); } + +void SPIRVTargetInfo::getTargetDefines(const LangOptions &Opts, + MacroBuilder &Builder) const { + DefineStd(Builder, "SPIRV", Opts); +} + +void SPIRV32TargetInfo::getTargetDefines(const LangOptions &Opts, + MacroBuilder &Builder) const { + SPIRVTargetInfo::getTargetDefines(Opts, Builder); + DefineStd(Builder, "SPIRV32", Opts); +} + +void SPIRV64TargetInfo::getTargetDefines(const LangOptions &Opts, + MacroBuilder &Builder) const { + SPIRVTargetInfo::getTargetDefines(Opts, Builder); + DefineStd(Builder, "SPIRV64", Opts); +} diff --git a/clang/lib/CodeGen/TargetInfo.cpp b/clang/lib/CodeGen/TargetInfo.cpp --- a/clang/lib/CodeGen/TargetInfo.cpp +++ b/clang/lib/CodeGen/TargetInfo.cpp @@ -10179,7 +10179,7 @@ } } //===----------------------------------------------------------------------===// -// SPIR ABI Implementation +// SPIR and SPIR-V ABI Implementation //===----------------------------------------------------------------------===// namespace { @@ -11289,6 +11289,8 @@ return SetCGInfo(new ARCTargetCodeGenInfo(Types)); case llvm::Triple::spir: case llvm::Triple::spir64: + case llvm::Triple::spirv32: + case llvm::Triple::spirv64: return SetCGInfo(new SPIRTargetCodeGenInfo(Types)); case llvm::Triple::ve: return SetCGInfo(new VETargetCodeGenInfo(Types)); diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -4356,7 +4356,7 @@ TC.addClangWarningOptions(CmdArgs); // FIXME: Subclass ToolChain for SPIR and move this to addClangWarningOptions. - if (Triple.isSPIR()) + if (Triple.isSPIR() || Triple.isSPIRV()) CmdArgs.push_back("-Wspir-compat"); // Select the appropriate action. diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp --- a/clang/lib/Frontend/InitPreprocessor.cpp +++ b/clang/lib/Frontend/InitPreprocessor.cpp @@ -1178,7 +1178,7 @@ if (LangOpts.OpenCL) { InitializeOpenCLFeatureTestMacros(TI, LangOpts, Builder); - if (TI.getTriple().isSPIR()) + if (TI.getTriple().isSPIR() || TI.getTriple().isSPIRV()) Builder.defineMacro("__IMAGE_SUPPORT__"); } diff --git a/clang/lib/Headers/opencl-c-base.h b/clang/lib/Headers/opencl-c-base.h --- a/clang/lib/Headers/opencl-c-base.h +++ b/clang/lib/Headers/opencl-c-base.h @@ -12,8 +12,8 @@ // Define extension macros #if (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200) -// For SPIR all extensions are supported. -#if defined(__SPIR__) +// For SPIR and SPIR-V all extensions are supported. +#if defined(__SPIR__) || defined(__SPIRV__) #define cl_khr_subgroup_extended_types 1 #define cl_khr_subgroup_non_uniform_vote 1 #define cl_khr_subgroup_ballot 1 @@ -26,7 +26,7 @@ #define __opencl_c_integer_dot_product_input_4x8bit 1 #define __opencl_c_integer_dot_product_input_4x8bit_packed 1 -#endif // defined(__SPIR__) +#endif // defined(__SPIR__) || defined(__SPIRV__) #endif // (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200) // Define feature macros for OpenCL C 2.0 @@ -46,8 +46,8 @@ // Define header-only feature macros for OpenCL C 3.0. #if (__OPENCL_CPP_VERSION__ == 202100 || __OPENCL_C_VERSION__ == 300) -// For the SPIR target all features are supported. -#if defined(__SPIR__) +// For the SPIR and SPIR-V target all features are supported. +#if defined(__SPIR__) || defined(__SPIRV__) #define __opencl_c_atomic_scope_all_devices 1 #endif // defined(__SPIR__) #endif // (__OPENCL_CPP_VERSION__ == 202100 || __OPENCL_C_VERSION__ == 300) diff --git a/clang/lib/Headers/opencl-c.h b/clang/lib/Headers/opencl-c.h --- a/clang/lib/Headers/opencl-c.h +++ b/clang/lib/Headers/opencl-c.h @@ -24,10 +24,10 @@ #endif //__OPENCL_C_VERSION__ < CL_VERSION_2_0 -#if (defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_1_2)) && defined(__SPIR__) +#if (defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_1_2)) && (defined(__SPIR__) || defined(__SPIRV__)) #pragma OPENCL EXTENSION cl_intel_planar_yuv : begin #pragma OPENCL EXTENSION cl_intel_planar_yuv : end -#endif // (defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_1_2)) && defined(__SPIR__) +#endif // (defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_1_2)) && (defined(__SPIR__) || defined(__SPIRV__)) #define __ovld __attribute__((overloadable)) #define __conv __attribute__((convergent)) diff --git a/clang/test/CodeGen/complex-math.c b/clang/test/CodeGen/complex-math.c --- a/clang/test/CodeGen/complex-math.c +++ b/clang/test/CodeGen/complex-math.c @@ -7,6 +7,7 @@ // RUN: %clang_cc1 %s -O0 -fno-experimental-new-pass-manager -emit-llvm -triple thumbv7k-apple-watchos2.0 -o - -target-abi aapcs16 | FileCheck %s --check-prefix=ARM7K // RUN: %clang_cc1 %s -O0 -fno-experimental-new-pass-manager -emit-llvm -triple aarch64-unknown-unknown -ffast-math -ffp-contract=fast -o - | FileCheck %s --check-prefix=AARCH64-FASTMATH // RUN: %clang_cc1 %s -O0 -fno-experimental-new-pass-manager -emit-llvm -triple spir -o - | FileCheck %s --check-prefix=SPIR +// RUN: %clang_cc1 %s -O0 -fno-experimental-new-pass-manager -emit-llvm -triple spirv32 -o - | FileCheck %s --check-prefix=SPIR float _Complex add_float_rr(float a, float b) { // X86-LABEL: @add_float_rr( diff --git a/clang/test/CodeGen/ext-int-cc.c b/clang/test/CodeGen/ext-int-cc.c --- a/clang/test/CodeGen/ext-int-cc.c +++ b/clang/test/CodeGen/ext-int-cc.c @@ -10,6 +10,8 @@ // RUN: %clang_cc1 -triple mips -O3 -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefixes=MIPS // RUN: %clang_cc1 -triple spir64 -O3 -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefixes=SPIR64 // RUN: %clang_cc1 -triple spir -O3 -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefixes=SPIR +// RUN: %clang_cc1 -triple spirv32 -O3 -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefixes=SPIR +// RUN: %clang_cc1 -triple spirv64 -O3 -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefixes=SPIR64 // RUN: %clang_cc1 -triple hexagon -O3 -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefixes=HEX // RUN: %clang_cc1 -triple lanai -O3 -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefixes=LANAI // RUN: %clang_cc1 -triple r600 -O3 -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefixes=R600 diff --git a/clang/test/CodeGen/spir-half-type.cpp b/clang/test/CodeGen/spir-half-type.cpp --- a/clang/test/CodeGen/spir-half-type.cpp +++ b/clang/test/CodeGen/spir-half-type.cpp @@ -2,6 +2,10 @@ // RUN: %clang_cc1 -O0 -triple spir64 -emit-llvm %s -o - | FileCheck %s // RUN: %clang_cc1 -O0 -triple spir -fexperimental-new-pass-manager -emit-llvm %s -o - | FileCheck %s // RUN: %clang_cc1 -O0 -triple spir64 -fexperimental-new-pass-manager -emit-llvm %s -o - | FileCheck %s +// RUN: %clang_cc1 -O0 -triple spirv32 -emit-llvm %s -o - | FileCheck %s +// RUN: %clang_cc1 -O0 -triple spirv64 -emit-llvm %s -o - | FileCheck %s +// RUN: %clang_cc1 -O0 -triple spirv32 -fexperimental-new-pass-manager -emit-llvm %s -o - | FileCheck %s +// RUN: %clang_cc1 -O0 -triple spirv64 -fexperimental-new-pass-manager -emit-llvm %s -o - | FileCheck %s // This file tests that using the _Float16 type with the spir target will not // use the llvm intrinsics but instead will use the half arithmetic diff --git a/clang/test/CodeGen/target-data.c b/clang/test/CodeGen/target-data.c --- a/clang/test/CodeGen/target-data.c +++ b/clang/test/CodeGen/target-data.c @@ -267,10 +267,14 @@ // RUN: %clang_cc1 -triple spir-unknown -o - -emit-llvm %s | \ // RUN: FileCheck %s -check-prefix=SPIR +// RUN: %clang_cc1 -triple spirv32-unknown -o - -emit-llvm %s | \ +// RUN: FileCheck %s -check-prefix=SPIR // SPIR: target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" // RUN: %clang_cc1 -triple spir64-unknown -o - -emit-llvm %s | \ // RUN: FileCheck %s -check-prefix=SPIR64 +// RUN: %clang_cc1 -triple spirv64-unknown -o - -emit-llvm %s | \ +// RUN: FileCheck %s -check-prefix=SPIR64 // SPIR64: target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" // RUN: %clang_cc1 -triple bpfel -o - -emit-llvm %s | \ diff --git a/clang/test/CodeGenCXX/builtin-calling-conv.cpp b/clang/test/CodeGenCXX/builtin-calling-conv.cpp --- a/clang/test/CodeGenCXX/builtin-calling-conv.cpp +++ b/clang/test/CodeGenCXX/builtin-calling-conv.cpp @@ -1,7 +1,9 @@ // RUN: %clang_cc1 -triple x86_64-linux-pc -DREDECL -emit-llvm %s -o - | FileCheck %s -check-prefix LINUX // RUN: %clang_cc1 -triple spir-unknown-unknown -DREDECL -DSPIR -emit-llvm %s -o - | FileCheck %s -check-prefix SPIR +// RUN: %clang_cc1 -triple spirv32-unknown-unknown -DREDECL -DSPIR -emit-llvm %s -o - | FileCheck %s -check-prefix SPIR // RUN: %clang_cc1 -triple x86_64-linux-pc -emit-llvm %s -o - | FileCheck %s -check-prefix LINUX // RUN: %clang_cc1 -triple spir-unknown-unknown -DSPIR -emit-llvm %s -o - | FileCheck %s -check-prefix SPIR +// RUN: %clang_cc1 -triple spirv32-unknown-unknown -DSPIR -emit-llvm %s -o - | FileCheck %s -check-prefix SPIR // RUN: %clang_cc1 -triple i386-windows-pc -fdefault-calling-conv=stdcall -emit-llvm %s -o - | FileCheck %s -check-prefix WIN32 #ifdef REDECL diff --git a/clang/test/CodeGenOpenCL/as_type.cl b/clang/test/CodeGenOpenCL/as_type.cl --- a/clang/test/CodeGenOpenCL/as_type.cl +++ b/clang/test/CodeGenOpenCL/as_type.cl @@ -1,4 +1,5 @@ // RUN: %clang_cc1 %s -emit-llvm -triple spir-unknown-unknown -o - | FileCheck %s +// RUN: %clang_cc1 %s -emit-llvm -triple spirv32-unknown-unknown -o - | FileCheck %s typedef __attribute__(( ext_vector_type(3) )) char char3; typedef __attribute__(( ext_vector_type(4) )) char char4; diff --git a/clang/test/CodeGenOpenCL/atomic-ops-libcall.cl b/clang/test/CodeGenOpenCL/atomic-ops-libcall.cl --- a/clang/test/CodeGenOpenCL/atomic-ops-libcall.cl +++ b/clang/test/CodeGenOpenCL/atomic-ops-libcall.cl @@ -1,4 +1,5 @@ // RUN: %clang_cc1 < %s -cl-std=CL2.0 -triple spir64 -emit-llvm | FileCheck -check-prefix=SPIR %s +// RUN: %clang_cc1 < %s -cl-std=CL2.0 -triple spirv64 -emit-llvm | FileCheck -check-prefix=SPIR %s // RUN: %clang_cc1 < %s -cl-std=CL2.0 -triple armv5e-none-linux-gnueabi -emit-llvm | FileCheck -check-prefix=ARM %s typedef enum memory_order { memory_order_relaxed = __ATOMIC_RELAXED, diff --git a/clang/test/CodeGenOpenCL/blocks.cl b/clang/test/CodeGenOpenCL/blocks.cl --- a/clang/test/CodeGenOpenCL/blocks.cl +++ b/clang/test/CodeGenOpenCL/blocks.cl @@ -1,6 +1,8 @@ // RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -o - -O0 -triple spir-unknown-unknown | FileCheck -check-prefixes=COMMON,SPIR %s +// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -o - -O0 -triple spirv32-unknown-unknown | FileCheck -check-prefixes=COMMON,SPIR %s // RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -o - -O0 -triple amdgcn-amd-amdhsa | FileCheck -check-prefixes=COMMON,AMDGCN %s // RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -o - -O0 -debug-info-kind=limited -triple spir-unknown-unknown | FileCheck -check-prefixes=CHECK-DEBUG %s +// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -o - -O0 -debug-info-kind=limited -triple spirv32-unknown-unknown | FileCheck -check-prefixes=CHECK-DEBUG %s // RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -o - -O0 -debug-info-kind=limited -triple amdgcn-amd-amdhsa | FileCheck -check-prefixes=CHECK-DEBUG %s // SPIR: %struct.__opencl_block_literal_generic = type { i32, i32, i8 addrspace(4)* } diff --git a/clang/test/CodeGenOpenCL/cast_image.cl b/clang/test/CodeGenOpenCL/cast_image.cl --- a/clang/test/CodeGenOpenCL/cast_image.cl +++ b/clang/test/CodeGenOpenCL/cast_image.cl @@ -1,5 +1,6 @@ // RUN: %clang_cc1 -emit-llvm -o - -triple amdgcn--amdhsa %s | FileCheck --check-prefix=AMDGCN %s // RUN: %clang_cc1 -emit-llvm -o - -triple spir-unknown-unknown %s | FileCheck --check-prefix=SPIR %s +// RUN: %clang_cc1 -emit-llvm -o - -triple spirv32-unknown-unknown %s | FileCheck --check-prefix=SPIR %s #ifdef __AMDGCN__ diff --git a/clang/test/CodeGenOpenCL/cl20-device-side-enqueue.cl b/clang/test/CodeGenOpenCL/cl20-device-side-enqueue.cl --- a/clang/test/CodeGenOpenCL/cl20-device-side-enqueue.cl +++ b/clang/test/CodeGenOpenCL/cl20-device-side-enqueue.cl @@ -1,6 +1,9 @@ // RUN: %clang_cc1 %s -cl-std=CL2.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spir-unknown-unknown" | FileCheck %s --check-prefix=COMMON --check-prefix=B32 // RUN: %clang_cc1 %s -cl-std=CL2.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spir64-unknown-unknown" | FileCheck %s --check-prefix=COMMON --check-prefix=B64 // RUN: %clang_cc1 %s -cl-std=CL2.0 -ffake-address-space-map -O1 -emit-llvm -o - -triple "spir64-unknown-unknown" | FileCheck %s --check-prefix=CHECK-LIFETIMES +// RUN: %clang_cc1 %s -cl-std=CL2.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spirv32-unknown-unknown" | FileCheck %s --check-prefix=COMMON --check-prefix=B32 +// RUN: %clang_cc1 %s -cl-std=CL2.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spirv64-unknown-unknown" | FileCheck %s --check-prefix=COMMON --check-prefix=B64 +// RUN: %clang_cc1 %s -cl-std=CL2.0 -ffake-address-space-map -O1 -emit-llvm -o - -triple "spirv64-unknown-unknown" | FileCheck %s --check-prefix=CHECK-LIFETIMES #pragma OPENCL EXTENSION cl_khr_subgroups : enable diff --git a/clang/test/CodeGenOpenCL/constant-addr-space-globals.cl b/clang/test/CodeGenOpenCL/constant-addr-space-globals.cl --- a/clang/test/CodeGenOpenCL/constant-addr-space-globals.cl +++ b/clang/test/CodeGenOpenCL/constant-addr-space-globals.cl @@ -1,4 +1,5 @@ // RUN: %clang_cc1 %s -triple "spir64-unknown-unknown" -cl-opt-disable -ffake-address-space-map -emit-llvm -o - | FileCheck %s +// RUN: %clang_cc1 %s -triple "spirv64-unknown-unknown" -cl-opt-disable -ffake-address-space-map -emit-llvm -o - | FileCheck %s // CHECK: @array ={{.*}} addrspace({{[0-9]+}}) constant __constant float array[2] = {0.0f, 1.0f}; diff --git a/clang/test/CodeGenOpenCL/enqueue-kernel-non-entry-block.cl b/clang/test/CodeGenOpenCL/enqueue-kernel-non-entry-block.cl --- a/clang/test/CodeGenOpenCL/enqueue-kernel-non-entry-block.cl +++ b/clang/test/CodeGenOpenCL/enqueue-kernel-non-entry-block.cl @@ -1,6 +1,8 @@ // RUN: %clang_cc1 -cl-std=CL2.0 -O0 -emit-llvm -o - -triple amdgcn < %s | FileCheck %s --check-prefixes=COMMON,AMDGPU // RUN: %clang_cc1 -cl-std=CL2.0 -O0 -emit-llvm -o - -triple "spir-unknown-unknown" < %s | FileCheck %s --check-prefixes=COMMON,SPIR32 // RUN: %clang_cc1 -cl-std=CL2.0 -O0 -emit-llvm -o - -triple "spir64-unknown-unknown" < %s | FileCheck %s --check-prefixes=COMMON,SPIR64 +// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -emit-llvm -o - -triple "spirv32-unknown-unknown" < %s | FileCheck %s --check-prefixes=COMMON,SPIR32 +// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -emit-llvm -o - -triple "spirv64-unknown-unknown" < %s | FileCheck %s --check-prefixes=COMMON,SPIR64 // RUN: %clang_cc1 -cl-std=CL2.0 -O0 -debug-info-kind=limited -gno-column-info -emit-llvm -o - -triple amdgcn < %s | FileCheck %s --check-prefixes=CHECK-DEBUG // Check that the enqueue_kernel array temporary is in the entry block to avoid @@ -29,5 +31,5 @@ // CHECK-DEBUG: ![[TESTFILE:[0-9]+]] = !DIFile(filename: "" // CHECK-DEBUG: ![[TESTSCOPE:[0-9]+]] = distinct !DISubprogram(name: "test", {{.*}} file: ![[TESTFILE]] -// CHECK-DEBUG: ![[IFSCOPE:[0-9]+]] = distinct !DILexicalBlock(scope: ![[TESTSCOPE]], file: ![[TESTFILE]], line: 24) -// CHECK-DEBUG: ![[TEMPLOCATION]] = !DILocation(line: 25, scope: ![[IFSCOPE]]) +// CHECK-DEBUG: ![[IFSCOPE:[0-9]+]] = distinct !DILexicalBlock(scope: ![[TESTSCOPE]], file: ![[TESTFILE]], line: 26) +// CHECK-DEBUG: ![[TEMPLOCATION]] = !DILocation(line: 27, scope: ![[IFSCOPE]]) diff --git a/clang/test/CodeGenOpenCL/fdeclare-opencl-builtins.cl b/clang/test/CodeGenOpenCL/fdeclare-opencl-builtins.cl --- a/clang/test/CodeGenOpenCL/fdeclare-opencl-builtins.cl +++ b/clang/test/CodeGenOpenCL/fdeclare-opencl-builtins.cl @@ -1,5 +1,7 @@ // RUN: %clang_cc1 -emit-llvm -o - -O0 -triple spir-unknown-unknown -cl-std=CL1.2 -finclude-default-header %s | FileCheck %s // RUN: %clang_cc1 -emit-llvm -o - -O0 -triple spir-unknown-unknown -cl-std=CL1.2 -fdeclare-opencl-builtins -finclude-default-header %s | FileCheck %s +// RUN: %clang_cc1 -emit-llvm -o - -O0 -triple spirv32-unknown-unknown -cl-std=CL1.2 -finclude-default-header %s | FileCheck %s +// RUN: %clang_cc1 -emit-llvm -o - -O0 -triple spirv32-unknown-unknown -cl-std=CL1.2 -fdeclare-opencl-builtins -finclude-default-header %s | FileCheck %s // Test that mix is correctly defined. // CHECK-LABEL: @test_float diff --git a/clang/test/CodeGenOpenCL/fpmath.cl b/clang/test/CodeGenOpenCL/fpmath.cl --- a/clang/test/CodeGenOpenCL/fpmath.cl +++ b/clang/test/CodeGenOpenCL/fpmath.cl @@ -1,7 +1,10 @@ // RUN: %clang_cc1 %s -emit-llvm -o - -triple spir-unknown-unknown | FileCheck --check-prefix=CHECK --check-prefix=NODIVOPT %s // RUN: %clang_cc1 %s -emit-llvm -o - -triple spir-unknown-unknown -cl-fp32-correctly-rounded-divide-sqrt | FileCheck --check-prefix=CHECK --check-prefix=DIVOPT %s +// RUN: %clang_cc1 %s -emit-llvm -o - -triple spirv32-unknown-unknown | FileCheck --check-prefix=CHECK --check-prefix=NODIVOPT %s +// RUN: %clang_cc1 %s -emit-llvm -o - -triple spirv32-unknown-unknown -cl-fp32-correctly-rounded-divide-sqrt | FileCheck --check-prefix=CHECK --check-prefix=DIVOPT %s // RUN: %clang_cc1 %s -emit-llvm -o - -DNOFP64 -cl-std=CL1.2 -triple r600-unknown-unknown -target-cpu r600 -pedantic | FileCheck --check-prefix=CHECK-FLT %s // RUN: %clang_cc1 %s -emit-llvm -o - -DFP64 -cl-std=CL1.2 -triple spir-unknown-unknown -pedantic | FileCheck --check-prefix=CHECK-DBL %s +// RUN: %clang_cc1 %s -emit-llvm -o - -DFP64 -cl-std=CL1.2 -triple spirv32-unknown-unknown -pedantic | FileCheck --check-prefix=CHECK-DBL %s typedef __attribute__(( ext_vector_type(4) )) float float4; diff --git a/clang/test/CodeGenOpenCL/half.cl b/clang/test/CodeGenOpenCL/half.cl --- a/clang/test/CodeGenOpenCL/half.cl +++ b/clang/test/CodeGenOpenCL/half.cl @@ -1,4 +1,5 @@ // RUN: %clang_cc1 %s -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s +// RUN: %clang_cc1 %s -emit-llvm -o - -triple spirv32-unknown-unknown | FileCheck %s // RUN: %clang_cc1 %s -emit-llvm -o - -triple x86_64-pc-win32 | FileCheck %s #pragma OPENCL EXTENSION cl_khr_fp16 : enable diff --git a/clang/test/CodeGenOpenCL/kernel-arg-info.cl b/clang/test/CodeGenOpenCL/kernel-arg-info.cl --- a/clang/test/CodeGenOpenCL/kernel-arg-info.cl +++ b/clang/test/CodeGenOpenCL/kernel-arg-info.cl @@ -1,5 +1,7 @@ // RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -o - -triple spir-unknown-unknown | FileCheck %s // RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -o - -triple spir-unknown-unknown -cl-kernel-arg-info | FileCheck %s -check-prefix ARGINFO +// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -o - -triple spirv32-unknown-unknown | FileCheck %s +// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -o - -triple spirv32-unknown-unknown -cl-kernel-arg-info | FileCheck %s -check-prefix ARGINFO kernel void foo(global int * globalintp, global int * restrict globalintrestrictp, global const int * globalconstintp, diff --git a/clang/test/CodeGenOpenCL/opencl_types.cl b/clang/test/CodeGenOpenCL/opencl_types.cl --- a/clang/test/CodeGenOpenCL/opencl_types.cl +++ b/clang/test/CodeGenOpenCL/opencl_types.cl @@ -1,4 +1,5 @@ // RUN: %clang_cc1 -cl-std=CL2.0 %s -triple "spir-unknown-unknown" -emit-llvm -o - -O0 | FileCheck %s --check-prefixes=CHECK-COM,CHECK-SPIR +// RUN: %clang_cc1 -cl-std=CL2.0 %s -triple "spirv32-unknown-unknown" -emit-llvm -o - -O0 | FileCheck %s --check-prefixes=CHECK-COM,CHECK-SPIR // RUN: %clang_cc1 -cl-std=CL2.0 %s -triple "amdgcn--amdhsa" -emit-llvm -o - -O0 | FileCheck %s --check-prefixes=CHECK-COM,CHECK-AMDGCN #define CLK_ADDRESS_CLAMP_TO_EDGE 2 diff --git a/clang/test/CodeGenOpenCL/overload.cl b/clang/test/CodeGenOpenCL/overload.cl --- a/clang/test/CodeGenOpenCL/overload.cl +++ b/clang/test/CodeGenOpenCL/overload.cl @@ -1,5 +1,7 @@ // RUN: %clang_cc1 -cl-std=CL2.0 -emit-llvm -o - -triple spir-unknown-unknown %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -emit-llvm -o - -triple spirv32-unknown-unknown %s | FileCheck %s // RUN: %clang_cc1 -cl-std=CL3.0 -cl-ext=+__opencl_c_generic_address_space -emit-llvm -o - -triple spir-unknown-unknown %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL3.0 -cl-ext=+__opencl_c_generic_address_space -emit-llvm -o - -triple spirv32-unknown-unknown %s | FileCheck %s typedef short short4 __attribute__((ext_vector_type(4))); diff --git a/clang/test/CodeGenOpenCL/partial_initializer.cl b/clang/test/CodeGenOpenCL/partial_initializer.cl --- a/clang/test/CodeGenOpenCL/partial_initializer.cl +++ b/clang/test/CodeGenOpenCL/partial_initializer.cl @@ -1,4 +1,5 @@ // RUN: %clang_cc1 -triple spir-unknown-unknown -cl-std=CL2.0 -emit-llvm %s -O0 -o - | FileCheck %s +// RUN: %clang_cc1 -triple spirv32-unknown-unknown -cl-std=CL2.0 -emit-llvm %s -O0 -o - | FileCheck %s typedef __attribute__(( ext_vector_type(2) )) int int2; typedef __attribute__(( ext_vector_type(4) )) int int4; diff --git a/clang/test/CodeGenOpenCL/preserve_vec3.cl b/clang/test/CodeGenOpenCL/preserve_vec3.cl --- a/clang/test/CodeGenOpenCL/preserve_vec3.cl +++ b/clang/test/CodeGenOpenCL/preserve_vec3.cl @@ -1,4 +1,5 @@ // RUN: %clang_cc1 %s -emit-llvm -o - -triple spir-unknown-unknown -fpreserve-vec3-type | FileCheck %s +// RUN: %clang_cc1 %s -emit-llvm -o - -triple spirv32-unknown-unknown -fpreserve-vec3-type | FileCheck %s typedef char char3 __attribute__((ext_vector_type(3))); typedef short short3 __attribute__((ext_vector_type(3))); diff --git a/clang/test/CodeGenOpenCL/printf.cl b/clang/test/CodeGenOpenCL/printf.cl --- a/clang/test/CodeGenOpenCL/printf.cl +++ b/clang/test/CodeGenOpenCL/printf.cl @@ -2,6 +2,10 @@ // RUN: %clang_cc1 -cl-std=CL1.2 -cl-ext=-cl_khr_fp64 -triple spir-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck -check-prefixes=NOFP64,ALL %s // RUN: %clang_cc1 -cl-std=CL3.0 -cl-ext=+__opencl_c_fp64,+cl_khr_fp64 -triple spir-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck -check-prefixes=FP64,ALL %s // RUN: %clang_cc1 -cl-std=CL3.0 -cl-ext=-__opencl_c_fp64,-cl_khr_fp64 -triple spir-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck -check-prefixes=NOFP64,ALL %s +// RUN: %clang_cc1 -cl-std=CL1.2 -cl-ext=-+cl_khr_fp64 -triple spirv32-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck -check-prefixes=FP64,ALL %s +// RUN: %clang_cc1 -cl-std=CL1.2 -cl-ext=-cl_khr_fp64 -triple spirv32-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck -check-prefixes=NOFP64,ALL %s +// RUN: %clang_cc1 -cl-std=CL3.0 -cl-ext=+__opencl_c_fp64,+cl_khr_fp64 -triple spirv32-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck -check-prefixes=FP64,ALL %s +// RUN: %clang_cc1 -cl-std=CL3.0 -cl-ext=-__opencl_c_fp64,-cl_khr_fp64 -triple spirv32-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck -check-prefixes=NOFP64,ALL %s typedef __attribute__((ext_vector_type(2))) float float2; typedef __attribute__((ext_vector_type(2))) half half2; diff --git a/clang/test/CodeGenOpenCL/private-array-initialization.cl b/clang/test/CodeGenOpenCL/private-array-initialization.cl --- a/clang/test/CodeGenOpenCL/private-array-initialization.cl +++ b/clang/test/CodeGenOpenCL/private-array-initialization.cl @@ -1,4 +1,5 @@ // RUN: %clang_cc1 %s -triple spir-unknown-unknown -O0 -emit-llvm -o - | FileCheck -check-prefix=PRIVATE0 %s +// RUN: %clang_cc1 %s -triple spirv32-unknown-unknown -O0 -emit-llvm -o - | FileCheck -check-prefix=PRIVATE0 %s // RUN: %clang_cc1 %s -triple amdgcn-amd-amdhsa-unknown -O0 -emit-llvm -o - | FileCheck -check-prefix=PRIVATE5 %s // CHECK: @test.arr = private unnamed_addr addrspace(2) constant [3 x i32] [i32 1, i32 2, i32 3], align 4 diff --git a/clang/test/CodeGenOpenCL/sampler.cl b/clang/test/CodeGenOpenCL/sampler.cl --- a/clang/test/CodeGenOpenCL/sampler.cl +++ b/clang/test/CodeGenOpenCL/sampler.cl @@ -1,6 +1,9 @@ // RUN: %clang_cc1 %s -emit-llvm -triple spir-unknown-unknown -o - -O0 | FileCheck %s // RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -triple spir-unknown-unknown -o - -O0 | FileCheck %s // RUN: %clang_cc1 %s -cl-std=clc++ -emit-llvm -triple spir-unknown-unknown -o - -O0 | FileCheck %s +// RUN: %clang_cc1 %s -emit-llvm -triple spirv32-unknown-unknown -o - -O0 | FileCheck %s +// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -triple spirv32-unknown-unknown -o - -O0 | FileCheck %s +// RUN: %clang_cc1 %s -cl-std=clc++ -emit-llvm -triple spirv32-unknown-unknown -o - -O0 | FileCheck %s // // This test covers 5 cases of sampler initialzation: // 1. function argument passing diff --git a/clang/test/CodeGenOpenCL/size_t.cl b/clang/test/CodeGenOpenCL/size_t.cl --- a/clang/test/CodeGenOpenCL/size_t.cl +++ b/clang/test/CodeGenOpenCL/size_t.cl @@ -1,5 +1,7 @@ // RUN: %clang_cc1 %s -cl-std=CL2.0 -finclude-default-header -fdeclare-opencl-builtins -emit-llvm -O0 -triple spir-unknown-unknown -o - | FileCheck --check-prefix=SZ32 %s // RUN: %clang_cc1 %s -cl-std=CL2.0 -finclude-default-header -fdeclare-opencl-builtins -emit-llvm -O0 -triple spir64-unknown-unknown -o - | FileCheck --check-prefix=SZ64 --check-prefix=SZ64ONLY %s +// RUN: %clang_cc1 %s -cl-std=CL2.0 -finclude-default-header -fdeclare-opencl-builtins -emit-llvm -O0 -triple spirv32-unknown-unknown -o - | FileCheck --check-prefix=SZ32 %s +// RUN: %clang_cc1 %s -cl-std=CL2.0 -finclude-default-header -fdeclare-opencl-builtins -emit-llvm -O0 -triple spirv64-unknown-unknown -o - | FileCheck --check-prefix=SZ64 --check-prefix=SZ64ONLY %s // RUN: %clang_cc1 %s -cl-std=CL2.0 -finclude-default-header -fdeclare-opencl-builtins -emit-llvm -O0 -triple amdgcn -o - | FileCheck --check-prefix=SZ64 --check-prefix=AMDGCN %s // RUN: %clang_cc1 %s -cl-std=CL2.0 -finclude-default-header -fdeclare-opencl-builtins -emit-llvm -O0 -triple amdgcn---opencl -o - | FileCheck --check-prefix=SZ64 --check-prefix=AMDGCN %s diff --git a/clang/test/CodeGenOpenCL/spir-calling-conv.cl b/clang/test/CodeGenOpenCL/spir-calling-conv.cl --- a/clang/test/CodeGenOpenCL/spir-calling-conv.cl +++ b/clang/test/CodeGenOpenCL/spir-calling-conv.cl @@ -1,4 +1,5 @@ // RUN: %clang_cc1 %s -triple "spir-unknown-unknown" -emit-llvm -o - | FileCheck %s +// RUN: %clang_cc1 %s -triple "spirv32-unknown-unknown" -emit-llvm -o - | FileCheck %s int get_dummy_id(int D); diff --git a/clang/test/CodeGenOpenCL/spir-debug-info-pointer-address-space.cl b/clang/test/CodeGenOpenCL/spir-debug-info-pointer-address-space.cl --- a/clang/test/CodeGenOpenCL/spir-debug-info-pointer-address-space.cl +++ b/clang/test/CodeGenOpenCL/spir-debug-info-pointer-address-space.cl @@ -1,5 +1,7 @@ // RUN: %clang_cc1 -cl-std=CL2.0 -debug-info-kind=limited -dwarf-version=5 -emit-llvm -O0 -triple spir-unknown-unknown -o - %s | FileCheck %s // RUN: %clang_cc1 -cl-std=CL2.0 -debug-info-kind=limited -dwarf-version=5 -emit-llvm -O0 -triple spir64-unknown-unknown -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -debug-info-kind=limited -dwarf-version=5 -emit-llvm -O0 -triple spirv32-unknown-unknown -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -debug-info-kind=limited -dwarf-version=5 -emit-llvm -O0 -triple spirv64-unknown-unknown -o - %s | FileCheck %s // CHECK-DAG: ![[DWARF_ADDRESS_SPACE_GLOBAL:[0-9]+]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !{{[0-9]+}}, size: {{[0-9]+}}, dwarfAddressSpace: 1) // CHECK-DAG: ![[DWARF_ADDRESS_SPACE_CONSTANT:[0-9]+]] = !DIDerivedType(tag: DW_TAG_pointer_type, baseType: !{{[0-9]+}}, size: {{[0-9]+}}, dwarfAddressSpace: 2) diff --git a/clang/test/CodeGenOpenCL/spirv32_target.cl b/clang/test/CodeGenOpenCL/spirv32_target.cl new file mode 100644 --- /dev/null +++ b/clang/test/CodeGenOpenCL/spirv32_target.cl @@ -0,0 +1,22 @@ +// RUN: %clang_cc1 %s -triple "spirv32-unknown-unknown" -emit-llvm -o - | FileCheck %s + +// CHECK: target triple = "spirv32-unknown-unknown" + +typedef struct { + char c; + void *v; + void *v2; +} my_st; + +kernel void foo(global long *arg) { + int res1[sizeof(my_st) == 12 ? 1 : -1]; + int res2[sizeof(void *) == 4 ? 1 : -1]; + int res3[sizeof(arg) == 4 ? 1 : -1]; + + my_st *tmp = 0; + + arg[0] = (long)(&tmp->v); +//CHECK: store i64 4, i64 addrspace(1)* + arg[1] = (long)(&tmp->v2); +//CHECK: store i64 8, i64 addrspace(1)* +} diff --git a/clang/test/CodeGenOpenCL/spirv64_target.cl b/clang/test/CodeGenOpenCL/spirv64_target.cl new file mode 100644 --- /dev/null +++ b/clang/test/CodeGenOpenCL/spirv64_target.cl @@ -0,0 +1,21 @@ +// RUN: %clang_cc1 %s -triple "spirv64-unknown-unknown" -emit-llvm -o - | FileCheck %s + +// CHECK: target triple = "spirv64-unknown-unknown" + +typedef struct { + char c; + void *v; + void *v2; +} my_st; + +kernel void foo(global long *arg) { + int res1[sizeof(my_st) == 24 ? 1 : -1]; + int res2[sizeof(void *) == 8 ? 1 : -1]; + int res3[sizeof(arg) == 8 ? 1 : -1]; + + my_st *tmp = 0; + arg[3] = (long)(&tmp->v); +//CHECK: store i64 8, i64 addrspace(1)* + arg[4] = (long)(&tmp->v2); +//CHECK: store i64 16, i64 addrspace(1)* +} diff --git a/clang/test/CodeGenOpenCL/to_addr_builtin.cl b/clang/test/CodeGenOpenCL/to_addr_builtin.cl --- a/clang/test/CodeGenOpenCL/to_addr_builtin.cl +++ b/clang/test/CodeGenOpenCL/to_addr_builtin.cl @@ -1,4 +1,5 @@ // RUN: %clang_cc1 -triple spir-unknown-unknown -emit-llvm -O0 -cl-std=clc++ -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple spirv32-unknown-unknown -emit-llvm -O0 -cl-std=clc++ -o - %s | FileCheck %s // CHECK: %[[A:.*]] = type { float, float, float } typedef struct { diff --git a/clang/test/CodeGenOpenCL/vectorLoadStore.cl b/clang/test/CodeGenOpenCL/vectorLoadStore.cl --- a/clang/test/CodeGenOpenCL/vectorLoadStore.cl +++ b/clang/test/CodeGenOpenCL/vectorLoadStore.cl @@ -1,4 +1,5 @@ // RUN: %clang_cc1 -cl-std=CL2.0 -triple "spir-unknown-unknown" %s -emit-llvm -O0 -o - | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -triple "spirv32-unknown-unknown" %s -emit-llvm -O0 -o - | FileCheck %s typedef char char2 __attribute((ext_vector_type(2))); typedef char char3 __attribute((ext_vector_type(3))); diff --git a/clang/test/CodeGenOpenCL/vla.cl b/clang/test/CodeGenOpenCL/vla.cl --- a/clang/test/CodeGenOpenCL/vla.cl +++ b/clang/test/CodeGenOpenCL/vla.cl @@ -1,4 +1,5 @@ // RUN: %clang_cc1 -emit-llvm -triple "spir-unknown-unknown" -O0 -cl-std=CL2.0 -o - %s | FileCheck -check-prefixes=CHECK,SPIR %s +// RUN: %clang_cc1 -emit-llvm -triple "spirv32-unknown-unknown" -O0 -cl-std=CL2.0 -o - %s | FileCheck -check-prefixes=CHECK,SPIR %s // RUN: %clang_cc1 -emit-llvm -triple amdgcn-amd-amdhsa -O0 -cl-std=CL2.0 -o - %s | FileCheck -check-prefixes=CHECK,AMDGCN %s constant int sz0 = 5; diff --git a/clang/test/CodeGenOpenCLCXX/address-space-castoperators.cpp b/clang/test/CodeGenOpenCLCXX/address-space-castoperators.cpp --- a/clang/test/CodeGenOpenCLCXX/address-space-castoperators.cpp +++ b/clang/test/CodeGenOpenCLCXX/address-space-castoperators.cpp @@ -1,4 +1,5 @@ // RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=clc++ -emit-llvm -O0 -o - | FileCheck %s +// RUN: %clang_cc1 %s -triple spirv32-unknown-unknown -cl-std=clc++ -emit-llvm -O0 -o - | FileCheck %s void test_reinterpret_cast(){ __private float x; diff --git a/clang/test/CodeGenOpenCLCXX/address-space-deduction.clcpp b/clang/test/CodeGenOpenCLCXX/address-space-deduction.clcpp --- a/clang/test/CodeGenOpenCLCXX/address-space-deduction.clcpp +++ b/clang/test/CodeGenOpenCLCXX/address-space-deduction.clcpp @@ -1,5 +1,7 @@ // RUN: %clang_cc1 %s -triple spir-unknown-unknown -O0 -emit-llvm -o - | FileCheck %s -check-prefixes=COMMON,PTR // RUN: %clang_cc1 %s -triple spir-unknown-unknown -O0 -emit-llvm -o - -DREF | FileCheck %s -check-prefixes=COMMON,REF +// RUN: %clang_cc1 %s -triple spirv32-unknown-unknown -O0 -emit-llvm -o - | FileCheck %s -check-prefixes=COMMON,PTR +// RUN: %clang_cc1 %s -triple spirv32-unknown-unknown -O0 -emit-llvm -o - -DREF | FileCheck %s -check-prefixes=COMMON,REF #ifdef REF #define PTR & diff --git a/clang/test/CodeGenOpenCLCXX/address-space-deduction2.clcpp b/clang/test/CodeGenOpenCLCXX/address-space-deduction2.clcpp --- a/clang/test/CodeGenOpenCLCXX/address-space-deduction2.clcpp +++ b/clang/test/CodeGenOpenCLCXX/address-space-deduction2.clcpp @@ -1,4 +1,5 @@ // RUN: %clang_cc1 %s -triple spir-unknown-unknown -O0 -emit-llvm -o - | FileCheck %s +// RUN: %clang_cc1 %s -triple spirv32-unknown-unknown -O0 -emit-llvm -o - | FileCheck %s class P { public: diff --git a/clang/test/CodeGenOpenCLCXX/addrspace-constructors.clcpp b/clang/test/CodeGenOpenCLCXX/addrspace-constructors.clcpp --- a/clang/test/CodeGenOpenCLCXX/addrspace-constructors.clcpp +++ b/clang/test/CodeGenOpenCLCXX/addrspace-constructors.clcpp @@ -1,4 +1,5 @@ // RUN: %clang_cc1 %s -emit-llvm -o - -O0 -triple spir-unknown-unknown | FileCheck %s +// RUN: %clang_cc1 %s -emit-llvm -o - -O0 -triple spirv32-unknown-unknown | FileCheck %s // CHECK: %struct.X = type { i32 } diff --git a/clang/test/CodeGenOpenCLCXX/addrspace-conversion.clcpp b/clang/test/CodeGenOpenCLCXX/addrspace-conversion.clcpp --- a/clang/test/CodeGenOpenCLCXX/addrspace-conversion.clcpp +++ b/clang/test/CodeGenOpenCLCXX/addrspace-conversion.clcpp @@ -1,4 +1,5 @@ // RUN: %clang_cc1 %s -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s +// RUN: %clang_cc1 %s -triple spirv32-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s void bar(__generic volatile unsigned int* ptr) { diff --git a/clang/test/CodeGenOpenCLCXX/addrspace-derived-base.clcpp b/clang/test/CodeGenOpenCLCXX/addrspace-derived-base.clcpp --- a/clang/test/CodeGenOpenCLCXX/addrspace-derived-base.clcpp +++ b/clang/test/CodeGenOpenCLCXX/addrspace-derived-base.clcpp @@ -1,4 +1,5 @@ // RUN: %clang_cc1 %s -triple spir -emit-llvm -O0 -o - | FileCheck %s +// RUN: %clang_cc1 %s -triple spirv32 -emit-llvm -O0 -o - | FileCheck %s struct B { int mb; diff --git a/clang/test/CodeGenOpenCLCXX/addrspace-new-delete.clcpp b/clang/test/CodeGenOpenCLCXX/addrspace-new-delete.clcpp --- a/clang/test/CodeGenOpenCLCXX/addrspace-new-delete.clcpp +++ b/clang/test/CodeGenOpenCLCXX/addrspace-new-delete.clcpp @@ -1,4 +1,5 @@ // RUN: %clang_cc1 %s -triple spir -emit-llvm -O0 -o - | FileCheck %s +// RUN: %clang_cc1 %s -triple spirv32 -emit-llvm -O0 -o - | FileCheck %s typedef __SIZE_TYPE__ size_t; diff --git a/clang/test/CodeGenOpenCLCXX/addrspace-of-this.clcpp b/clang/test/CodeGenOpenCLCXX/addrspace-of-this.clcpp --- a/clang/test/CodeGenOpenCLCXX/addrspace-of-this.clcpp +++ b/clang/test/CodeGenOpenCLCXX/addrspace-of-this.clcpp @@ -1,6 +1,9 @@ // RUN: %clang_cc1 %s -triple spir-unknown-unknown -emit-llvm -pedantic -verify -O0 -o - -DDECL | FileCheck %s --check-prefixes="COMMON,EXPL" // RUN: %clang_cc1 %s -triple spir-unknown-unknown -emit-llvm -pedantic -verify -O0 -o - -DDECL -DUSE_DEFLT | FileCheck %s --check-prefixes="COMMON,IMPL" // RUN: %clang_cc1 %s -triple spir-unknown-unknown -emit-llvm -pedantic -verify -O0 -o - | FileCheck %s --check-prefixes="COMMON,IMPL" +// RUN: %clang_cc1 %s -triple spirv32-unknown-unknown -emit-llvm -pedantic -verify -O0 -o - -DDECL | FileCheck %s --check-prefixes="COMMON,EXPL" +// RUN: %clang_cc1 %s -triple spirv32-unknown-unknown -emit-llvm -pedantic -verify -O0 -o - -DDECL -DUSE_DEFLT | FileCheck %s --check-prefixes="COMMON,IMPL" +// RUN: %clang_cc1 %s -triple spirv32-unknown-unknown -emit-llvm -pedantic -verify -O0 -o - | FileCheck %s --check-prefixes="COMMON,IMPL" // expected-no-diagnostics // Test that the 'this' pointer is in the __generic address space. diff --git a/clang/test/CodeGenOpenCLCXX/addrspace-operators.clcpp b/clang/test/CodeGenOpenCLCXX/addrspace-operators.clcpp --- a/clang/test/CodeGenOpenCLCXX/addrspace-operators.clcpp +++ b/clang/test/CodeGenOpenCLCXX/addrspace-operators.clcpp @@ -1,4 +1,5 @@ //RUN: %clang_cc1 %s -triple spir -emit-llvm -O0 -o - | FileCheck %s +//RUN: %clang_cc1 %s -triple spirv32 -emit-llvm -O0 -o - | FileCheck %s enum E { a, diff --git a/clang/test/CodeGenOpenCLCXX/addrspace-references.clcpp b/clang/test/CodeGenOpenCLCXX/addrspace-references.clcpp --- a/clang/test/CodeGenOpenCLCXX/addrspace-references.clcpp +++ b/clang/test/CodeGenOpenCLCXX/addrspace-references.clcpp @@ -1,4 +1,5 @@ //RUN: %clang_cc1 %s -triple spir -emit-llvm -o - -O0 | FileCheck %s +//RUN: %clang_cc1 %s -triple spirv32 -emit-llvm -o - -O0 | FileCheck %s typedef short short2 __attribute__((ext_vector_type(2))); diff --git a/clang/test/CodeGenOpenCLCXX/addrspace-with-class.clcpp b/clang/test/CodeGenOpenCLCXX/addrspace-with-class.clcpp --- a/clang/test/CodeGenOpenCLCXX/addrspace-with-class.clcpp +++ b/clang/test/CodeGenOpenCLCXX/addrspace-with-class.clcpp @@ -1,5 +1,7 @@ // RUN: %clang_cc1 %s -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s // RUN: %clang_cc1 %s -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s --check-prefix=CHECK-DEFINITIONS +// RUN: %clang_cc1 %s -triple spirv32-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s +// RUN: %clang_cc1 %s -triple spirv32-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s --check-prefix=CHECK-DEFINITIONS // This test ensures the proper address spaces and address space cast are used // for constructors, member functions and destructors. diff --git a/clang/test/CodeGenOpenCLCXX/addrspace_cast.clcpp b/clang/test/CodeGenOpenCLCXX/addrspace_cast.clcpp --- a/clang/test/CodeGenOpenCLCXX/addrspace_cast.clcpp +++ b/clang/test/CodeGenOpenCLCXX/addrspace_cast.clcpp @@ -1,4 +1,5 @@ //RUN: %clang_cc1 %s -triple spir -emit-llvm -O0 -o - | FileCheck %s +//RUN: %clang_cc1 %s -triple spirv32 -emit-llvm -O0 -o - | FileCheck %s //CHECK-LABEL: define{{.*}} spir_func void @_Z3barPU3AS1i void bar(global int *gl) { diff --git a/clang/test/CodeGenOpenCLCXX/atexit.clcpp b/clang/test/CodeGenOpenCLCXX/atexit.clcpp --- a/clang/test/CodeGenOpenCLCXX/atexit.clcpp +++ b/clang/test/CodeGenOpenCLCXX/atexit.clcpp @@ -1,4 +1,5 @@ //RUN: %clang_cc1 %s -triple spir -emit-llvm -O0 -o - | FileCheck %s +//RUN: %clang_cc1 %s -triple spirv32 -emit-llvm -O0 -o - | FileCheck %s struct S { ~S(){}; diff --git a/clang/test/CodeGenOpenCLCXX/constexpr.clcpp b/clang/test/CodeGenOpenCLCXX/constexpr.clcpp --- a/clang/test/CodeGenOpenCLCXX/constexpr.clcpp +++ b/clang/test/CodeGenOpenCLCXX/constexpr.clcpp @@ -1,4 +1,5 @@ // RUN: %clang_cc1 %s -triple spir-unknown-unknown -O0 -emit-llvm -o - | FileCheck %s +// RUN: %clang_cc1 %s -triple spirv32-unknown-unknown -O0 -emit-llvm -o - | FileCheck %s typedef int int2 __attribute__((ext_vector_type(2))); typedef int int4 __attribute__((ext_vector_type(4))); diff --git a/clang/test/CodeGenOpenCLCXX/global_init.clcpp b/clang/test/CodeGenOpenCLCXX/global_init.clcpp --- a/clang/test/CodeGenOpenCLCXX/global_init.clcpp +++ b/clang/test/CodeGenOpenCLCXX/global_init.clcpp @@ -1,4 +1,5 @@ // RUN: %clang_cc1 %s -triple spir -emit-llvm -O0 -o - | FileCheck %s +// RUN: %clang_cc1 %s -triple spirv32 -emit-llvm -O0 -o - | FileCheck %s struct S { S() {} diff --git a/clang/test/CodeGenOpenCLCXX/local_addrspace_init.clcpp b/clang/test/CodeGenOpenCLCXX/local_addrspace_init.clcpp --- a/clang/test/CodeGenOpenCLCXX/local_addrspace_init.clcpp +++ b/clang/test/CodeGenOpenCLCXX/local_addrspace_init.clcpp @@ -1,4 +1,5 @@ // RUN: %clang_cc1 %s -triple spir -emit-llvm -O0 -o - | FileCheck %s +// RUN: %clang_cc1 %s -triple spirv32 -emit-llvm -O0 -o - | FileCheck %s // Test that we don't initialize local address space objects. //CHECK: @_ZZ4testE1i = internal addrspace(3) global i32 undef diff --git a/clang/test/CodeGenOpenCLCXX/method-overload-address-space.clcpp b/clang/test/CodeGenOpenCLCXX/method-overload-address-space.clcpp --- a/clang/test/CodeGenOpenCLCXX/method-overload-address-space.clcpp +++ b/clang/test/CodeGenOpenCLCXX/method-overload-address-space.clcpp @@ -1,4 +1,5 @@ //RUN: %clang_cc1 %s -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s +//RUN: %clang_cc1 %s -triple spirv32-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s struct C { void foo() __local; diff --git a/clang/test/CodeGenOpenCLCXX/reinterpret_cast.clcpp b/clang/test/CodeGenOpenCLCXX/reinterpret_cast.clcpp --- a/clang/test/CodeGenOpenCLCXX/reinterpret_cast.clcpp +++ b/clang/test/CodeGenOpenCLCXX/reinterpret_cast.clcpp @@ -1,4 +1,5 @@ //RUN: %clang_cc1 %s -triple spir -emit-llvm -O0 -o - | FileCheck %s +//RUN: %clang_cc1 %s -triple spirv32 -emit-llvm -O0 -o - | FileCheck %s typedef int int2 __attribute__((ext_vector_type(2))); typedef int int4 __attribute__((ext_vector_type(4))); diff --git a/clang/test/CodeGenOpenCLCXX/template-address-spaces.clcpp b/clang/test/CodeGenOpenCLCXX/template-address-spaces.clcpp --- a/clang/test/CodeGenOpenCLCXX/template-address-spaces.clcpp +++ b/clang/test/CodeGenOpenCLCXX/template-address-spaces.clcpp @@ -1,4 +1,5 @@ // RUN: %clang_cc1 %s -emit-llvm -o - -O0 -triple spir-unknown-unknown | FileCheck %s +// RUN: %clang_cc1 %s -emit-llvm -o - -O0 -triple spirv32-unknown-unknown | FileCheck %s template struct S{ diff --git a/clang/test/Driver/opencl.cl b/clang/test/Driver/opencl.cl --- a/clang/test/Driver/opencl.cl +++ b/clang/test/Driver/opencl.cl @@ -20,6 +20,7 @@ // RUN: not %clang -cl-std=c99 -DOPENCL %s 2>&1 | FileCheck --check-prefix=CHECK-C99 %s // RUN: not %clang -cl-std=invalid -DOPENCL %s 2>&1 | FileCheck --check-prefix=CHECK-INVALID %s // RUN: %clang -S -### -target spir-unknown-unknown %s 2>&1 | FileCheck --check-prefix=CHECK-W-SPIR-COMPAT %s +// RUN: %clang -S -### -target spirv32-unknown-unknown %s 2>&1 | FileCheck --check-prefix=CHECK-W-SPIR-COMPAT %s // RUN: %clang -S -### -target amdgcn-amd-amdhsa-opencl %s 2>&1 | FileCheck --check-prefix=CHECK-NO-W-SPIR-COMPAT %s // CHECK-CL: "-cc1" {{.*}} "-cl-std=CL" diff --git a/clang/test/Headers/opencl-builtins.cl b/clang/test/Headers/opencl-builtins.cl --- a/clang/test/Headers/opencl-builtins.cl +++ b/clang/test/Headers/opencl-builtins.cl @@ -1,6 +1,8 @@ // RUN: clang-tblgen -gen-clang-opencl-builtin-tests %clang_src_sema_dir/OpenCLBuiltins.td -o %t.cl // RUN: %clang_cc1 -include %s %t.cl -triple spir -verify -fsyntax-only -cl-std=CL2.0 -finclude-default-header // RUN: %clang_cc1 -include %s %t.cl -triple spir -verify -fsyntax-only -cl-std=CL2.0 -fdeclare-opencl-builtins -finclude-default-header +// RUN: %clang_cc1 -include %s %t.cl -triple spirv32 -verify -fsyntax-only -cl-std=CL2.0 -finclude-default-header +// RUN: %clang_cc1 -include %s %t.cl -triple spirv32 -verify -fsyntax-only -cl-std=CL2.0 -fdeclare-opencl-builtins -finclude-default-header // Generate an OpenCL source file containing a call to each builtin from // OpenCLBuiltins.td and then run that generated source file through the diff --git a/clang/test/Headers/opencl-c-header.cl b/clang/test/Headers/opencl-c-header.cl --- a/clang/test/Headers/opencl-c-header.cl +++ b/clang/test/Headers/opencl-c-header.cl @@ -4,6 +4,9 @@ // RUN: %clang_cc1 -O0 -triple spir-unknown-unknown -internal-isystem ../../lib/Headers -include opencl-c.h -emit-llvm -o - %s -verify -cl-std=clc++ | FileCheck %s --check-prefix=CHECK20 // RUN: %clang_cc1 -O0 -triple spir-unknown-unknown -internal-isystem ../../lib/Headers -include opencl-c.h -emit-llvm -o - %s -verify -cl-std=CL3.0 | FileCheck %s +// RUN: %clang_cc1 -O0 -triple spirv32-unknown-unknown -internal-isystem ../../lib/Headers -include opencl-c.h -emit-llvm -o - %s -verify | FileCheck %s + + // Test including the default header as a module. // The module should be compiled only once and loaded from cache afterwards. // Change the directory mode to read only to make sure no new modules are created. @@ -90,7 +93,7 @@ // Check that extension macros are defined correctly. // For SPIR all extensions are supported. -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) // Verify that cl_intel_planar_yuv extension is defined from OpenCL 1.2 onwards. #if defined(__OPENCL_CPP_VERSION__) || (__OPENCL_C_VERSION__ >= CL_VERSION_1_2) diff --git a/clang/test/Index/pipe-size.cl b/clang/test/Index/pipe-size.cl --- a/clang/test/Index/pipe-size.cl +++ b/clang/test/Index/pipe-size.cl @@ -1,6 +1,8 @@ // RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple x86_64-unknown-linux-gnu %s -o - | FileCheck %s --check-prefix=X86 // RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple spir-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPIR // RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple spir64-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPIR64 +// RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPIR +// RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPIR64 // RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple amdgcn-amd-amdhsa %s -o - | FileCheck %s --check-prefix=AMDGCN __kernel void testPipe( pipe int test ) { diff --git a/clang/test/Preprocessor/predefined-macros.c b/clang/test/Preprocessor/predefined-macros.c --- a/clang/test/Preprocessor/predefined-macros.c +++ b/clang/test/Preprocessor/predefined-macros.c @@ -200,7 +200,21 @@ // CHECK-SPIR64-DAG: #define __SPIR64__ 1 // CHECK-SPIR64-NOT: #define __SPIR32__ 1 -// RUN: %clang_cc1 %s -E -dM -o - -x hip -triple amdgcn-amd-amdhsa \ +// RUN: %clang_cc1 %s -E -dM -o - -x cl -triple spirv32-unknown-unknown \ +// RUN: | FileCheck -match-full-lines %s --check-prefix=CHECK-SPIRV32 +// CHECK-SPIRV32-DAG: #define __IMAGE_SUPPORT__ 1 +// CHECK-SPIRV32-DAG: #define __SPIRV__ 1 +// CHECK-SPIRV32-DAG: #define __SPIRV32__ 1 +// CHECK-SPIRV32-NOT: #define __SPIRV64__ 1 + +// RUN: %clang_cc1 %s -E -dM -o - -x cl -triple spirv64-unknown-unknown \ +// RUN: | FileCheck -match-full-lines %s --check-prefix=CHECK-SPIRV64 +// CHECK-SPIRV64-DAG: #define __IMAGE_SUPPORT__ 1 +// CHECK-SPIRV64-DAG: #define __SPIRV__ 1 +// CHECK-SPIRV64-DAG: #define __SPIRV64__ 1 +// CHECK-SPIRV64-NOT: #define __SPIRV32__ 1 + +// RUN: %clang_cc1 %s -E -dM -o - -x hip -triple amdgcn-amd-amdhsa \ // RUN: | FileCheck -match-full-lines %s --check-prefix=CHECK-HIP // CHECK-HIP-NOT: #define __CUDA_ARCH__ // CHECK-HIP: #define __HIPCC__ 1 diff --git a/clang/test/Sema/Float16.c b/clang/test/Sema/Float16.c --- a/clang/test/Sema/Float16.c +++ b/clang/test/Sema/Float16.c @@ -1,6 +1,7 @@ // RUN: %clang_cc1 -fsyntax-only -verify -triple x86_64-linux-pc %s // RUN: %clang_cc1 -fsyntax-only -verify -triple x86_64-linux-pc -target-feature +avx512fp16 %s -DHAVE // RUN: %clang_cc1 -fsyntax-only -verify -triple spir-unknown-unknown %s -DHAVE +// RUN: %clang_cc1 -fsyntax-only -verify -triple spirv32-unknown-unknown %s -DHAVE // RUN: %clang_cc1 -fsyntax-only -verify -triple armv7a-linux-gnu %s -DHAVE // RUN: %clang_cc1 -fsyntax-only -verify -triple aarch64-linux-gnu %s -DHAVE diff --git a/clang/test/SemaCXX/Float16.cpp b/clang/test/SemaCXX/Float16.cpp --- a/clang/test/SemaCXX/Float16.cpp +++ b/clang/test/SemaCXX/Float16.cpp @@ -1,5 +1,6 @@ // RUN: %clang_cc1 -fsyntax-only -verify -triple x86_64-linux-pc %s // RUN: %clang_cc1 -fsyntax-only -verify -triple spir-unknown-unknown %s -DHAVE +// RUN: %clang_cc1 -fsyntax-only -verify -triple spirv32-unknown-unknown %s -DHAVE // RUN: %clang_cc1 -fsyntax-only -verify -triple armv7a-linux-gnu %s -DHAVE // RUN: %clang_cc1 -fsyntax-only -verify -triple aarch64-linux-gnu %s -DHAVE diff --git a/llvm/include/llvm/ADT/Triple.h b/llvm/include/llvm/ADT/Triple.h --- a/llvm/include/llvm/ADT/Triple.h +++ b/llvm/include/llvm/ADT/Triple.h @@ -93,6 +93,9 @@ hsail64, // AMD HSAIL with 64-bit pointers spir, // SPIR: standard portable IR for OpenCL 32-bit version spir64, // SPIR: standard portable IR for OpenCL 64-bit version + spirv32, // SPIR-V with 32-bit pointers + spirv64, // SPIR-V with 64-bit pointers + spirvlogical, // SPIR-V with logical addressing kalimba, // Kalimba: generic kalimba shave, // SHAVE: Movidius vector VLIW processors lanai, // Lanai: Lanai 32-bit @@ -698,6 +701,17 @@ return getArch() == Triple::spir || getArch() == Triple::spir64; } + /// Tests whether the target is SPIR-V (32- or 64-bit). + bool isSPIRVPhysical() const { + return getArch() == Triple::spirv32 || getArch() == Triple::spirv64; + } + + /// Tests whether the target is SPIR-V with logical addressing + bool isSPIRVLogical() const { return getArch() == spirvlogical; } + + /// Tests whether the target is SPIR-V (32/64-bit or logical). + bool isSPIRV() const { return isSPIRVPhysical() || isSPIRVLogical(); } + /// Tests whether the target is NVPTX (32- or 64-bit). bool isNVPTX() const { return getArch() == Triple::nvptx || getArch() == Triple::nvptx64; diff --git a/llvm/lib/Support/Triple.cpp b/llvm/lib/Support/Triple.cpp --- a/llvm/lib/Support/Triple.cpp +++ b/llvm/lib/Support/Triple.cpp @@ -67,6 +67,9 @@ case sparcv9: return "sparcv9"; case spir64: return "spir64"; case spir: return "spir"; + case spirv32: return "spirv32"; + case spirv64: return "spirv64"; + case spirvlogical: return "spirvlogical"; case systemz: return "s390x"; case tce: return "tce"; case tcele: return "tcele"; @@ -147,6 +150,11 @@ case spir: case spir64: return "spir"; + + case spirv32: + case spirv64: + case spirvlogical:return "spirv"; + case kalimba: return "kalimba"; case lanai: return "lanai"; case shave: return "shave"; @@ -323,6 +331,9 @@ .Case("hsail64", hsail64) .Case("spir", spir) .Case("spir64", spir64) + .Case("spirv32", spirv32) + .Case("spirv64", spirv64) + .Case("spirvlogical", spirvlogical) .Case("kalimba", kalimba) .Case("lanai", lanai) .Case("shave", shave) @@ -456,6 +467,9 @@ .Case("hsail64", Triple::hsail64) .Case("spir", Triple::spir) .Case("spir64", Triple::spir64) + .Case("spirv32", Triple::spirv32) + .Case("spirv64", Triple::spirv64) + .Case("spirvlogical", Triple::spirvlogical) .StartsWith("kalimba", Triple::kalimba) .Case("lanai", Triple::lanai) .Case("renderscript32", Triple::renderscript32) @@ -753,6 +767,12 @@ case Triple::wasm32: case Triple::wasm64: return Triple::Wasm; + + case Triple::spirv32: + case Triple::spirv64: + case Triple::spirvlogical: + // TODO: In future this will be Triple::SPIRV. + return Triple::UnknownObjectFormat; } llvm_unreachable("unknown architecture"); } @@ -1268,6 +1288,7 @@ static unsigned getArchPointerBitWidth(llvm::Triple::ArchType Arch) { switch (Arch) { case llvm::Triple::UnknownArch: + case llvm::Triple::spirvlogical: return 0; case llvm::Triple::avr: @@ -1298,6 +1319,7 @@ case llvm::Triple::sparc: case llvm::Triple::sparcel: case llvm::Triple::spir: + case llvm::Triple::spirv32: case llvm::Triple::tce: case llvm::Triple::tcele: case llvm::Triple::thumb: @@ -1324,6 +1346,7 @@ case llvm::Triple::riscv64: case llvm::Triple::sparcv9: case llvm::Triple::spir64: + case llvm::Triple::spirv64: case llvm::Triple::systemz: case llvm::Triple::ve: case llvm::Triple::wasm64: @@ -1383,6 +1406,7 @@ case Triple::sparc: case Triple::sparcel: case Triple::spir: + case Triple::spirv32: case Triple::tce: case Triple::tcele: case Triple::thumb: @@ -1407,6 +1431,8 @@ case Triple::riscv64: T.setArch(Triple::riscv32); break; case Triple::sparcv9: T.setArch(Triple::sparc); break; case Triple::spir64: T.setArch(Triple::spir); break; + case Triple::spirv64: T.setArch(Triple::spirv32); break; + case Triple::spirvlogical: T.setArch(Triple::spirv32); break; case Triple::wasm64: T.setArch(Triple::wasm32); break; case Triple::x86_64: T.setArch(Triple::x86); break; } @@ -1451,6 +1477,7 @@ case Triple::riscv64: case Triple::sparcv9: case Triple::spir64: + case Triple::spirv64: case Triple::systemz: case Triple::ve: case Triple::wasm64: @@ -1473,6 +1500,8 @@ case Triple::riscv32: T.setArch(Triple::riscv64); break; case Triple::sparc: T.setArch(Triple::sparcv9); break; case Triple::spir: T.setArch(Triple::spir64); break; + case Triple::spirv32: T.setArch(Triple::spirv64); break; + case Triple::spirvlogical: T.setArch(Triple::spirv64); break; case Triple::thumb: T.setArch(Triple::aarch64); break; case Triple::thumbeb: T.setArch(Triple::aarch64_be); break; case Triple::wasm32: T.setArch(Triple::wasm64); break; @@ -1509,6 +1538,9 @@ case Triple::shave: case Triple::spir64: case Triple::spir: + case Triple::spirv32: + case Triple::spirv64: + case Triple::spirvlogical: case Triple::wasm32: case Triple::wasm64: case Triple::x86: @@ -1604,6 +1636,9 @@ case Triple::sparcel: case Triple::spir64: case Triple::spir: + case Triple::spirv32: + case Triple::spirv64: + case Triple::spirvlogical: case Triple::tcele: case Triple::thumb: case Triple::ve: diff --git a/llvm/unittests/ADT/TripleTest.cpp b/llvm/unittests/ADT/TripleTest.cpp --- a/llvm/unittests/ADT/TripleTest.cpp +++ b/llvm/unittests/ADT/TripleTest.cpp @@ -224,6 +224,21 @@ EXPECT_EQ(Triple::UnknownVendor, T.getVendor()); EXPECT_EQ(Triple::UnknownOS, T.getOS()); + T = Triple("spirv32-unknown-unknown"); + EXPECT_EQ(Triple::spirv32, T.getArch()); + EXPECT_EQ(Triple::UnknownVendor, T.getVendor()); + EXPECT_EQ(Triple::UnknownOS, T.getOS()); + + T = Triple("spirv64-unknown-unknown"); + EXPECT_EQ(Triple::spirv64, T.getArch()); + EXPECT_EQ(Triple::UnknownVendor, T.getVendor()); + EXPECT_EQ(Triple::UnknownOS, T.getOS()); + + T = Triple("spirvlogical-unknown-unknown"); + EXPECT_EQ(Triple::spirvlogical, T.getArch()); + EXPECT_EQ(Triple::UnknownVendor, T.getVendor()); + EXPECT_EQ(Triple::UnknownOS, T.getOS()); + T = Triple("x86_64-unknown-ananas"); EXPECT_EQ(Triple::x86_64, T.getArch()); EXPECT_EQ(Triple::UnknownVendor, T.getVendor()); @@ -865,6 +880,21 @@ EXPECT_FALSE(T.isArch32Bit()); EXPECT_TRUE(T.isArch64Bit()); + T.setArch(Triple::spirv32); + EXPECT_FALSE(T.isArch16Bit()); + EXPECT_TRUE(T.isArch32Bit()); + EXPECT_FALSE(T.isArch64Bit()); + + T.setArch(Triple::spirv64); + EXPECT_FALSE(T.isArch16Bit()); + EXPECT_FALSE(T.isArch32Bit()); + EXPECT_TRUE(T.isArch64Bit()); + + T.setArch(Triple::spirvlogical); + EXPECT_FALSE(T.isArch16Bit()); + EXPECT_FALSE(T.isArch32Bit()); + EXPECT_FALSE(T.isArch64Bit()); + T.setArch(Triple::sparc); EXPECT_FALSE(T.isArch16Bit()); EXPECT_TRUE(T.isArch32Bit()); @@ -1000,6 +1030,18 @@ EXPECT_EQ(Triple::spir, T.get32BitArchVariant().getArch()); EXPECT_EQ(Triple::spir64, T.get64BitArchVariant().getArch()); + T.setArch(Triple::spirv32); + EXPECT_EQ(Triple::spirv32, T.get32BitArchVariant().getArch()); + EXPECT_EQ(Triple::spirv64, T.get64BitArchVariant().getArch()); + + T.setArch(Triple::spirv64); + EXPECT_EQ(Triple::spirv32, T.get32BitArchVariant().getArch()); + EXPECT_EQ(Triple::spirv64, T.get64BitArchVariant().getArch()); + + T.setArch(Triple::spirvlogical); + EXPECT_EQ(Triple::spirv32, T.get32BitArchVariant().getArch()); + EXPECT_EQ(Triple::spirv64, T.get64BitArchVariant().getArch()); + T.setArch(Triple::wasm32); EXPECT_EQ(Triple::wasm32, T.get32BitArchVariant().getArch()); EXPECT_EQ(Triple::wasm64, T.get64BitArchVariant().getArch());