linjamaki created this revision.
Herald added subscribers: Naghasan, dexonsmith, wenlei, Anastasia, arphaman,
hiraditya.
linjamaki requested review of this revision.
Herald added projects: clang, LLVM.
Herald added subscribers: llvm-commits, cfe-commits.
Add new architectures ‘spirv32’, ‘spirv64’ and ‘spirlogical’
in Triple, ported from LLVM SPIR-V Backend repository
(https://github.com/KhronosGroup/LLVM-SPIRV-Backend). Add basic and
code generation target info for ‘spirv32’ and ‘spirv64’ and, thus,
enabling clang (LLVM IR) code emission to SPIR-V target. The target
information for SPIR-V is mostly reused from the SPIR target info as
starter.
Enable testing for SPIR-V target in many existing code generation
related clang tests where the SPIR target is tested and added two new
tests. SYCL related tests were skipped - they are left for SYCL
developers to decide if they want to enable or switch to SPIR-V
target.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D109144
Files:
clang/include/clang/Basic/DiagnosticGroups.td
clang/lib/Basic/Targets.cpp
clang/lib/Basic/Targets/SPIR.cpp
clang/lib/Basic/Targets/SPIR.h
clang/lib/CodeGen/TargetInfo.cpp
clang/lib/Driver/ToolChains/Clang.cpp
clang/lib/Frontend/InitPreprocessor.cpp
clang/lib/Headers/opencl-c-base.h
clang/lib/Headers/opencl-c.h
clang/test/CodeGen/complex-math.c
clang/test/CodeGen/ext-int-cc.c
clang/test/CodeGen/spir-half-type.cpp
clang/test/CodeGen/target-data.c
clang/test/CodeGenCXX/builtin-calling-conv.cpp
clang/test/CodeGenOpenCL/as_type.cl
clang/test/CodeGenOpenCL/atomic-ops-libcall.cl
clang/test/CodeGenOpenCL/blocks.cl
clang/test/CodeGenOpenCL/cast_image.cl
clang/test/CodeGenOpenCL/cl20-device-side-enqueue.cl
clang/test/CodeGenOpenCL/constant-addr-space-globals.cl
clang/test/CodeGenOpenCL/enqueue-kernel-non-entry-block.cl
clang/test/CodeGenOpenCL/fdeclare-opencl-builtins.cl
clang/test/CodeGenOpenCL/fpmath.cl
clang/test/CodeGenOpenCL/half.cl
clang/test/CodeGenOpenCL/kernel-arg-info.cl
clang/test/CodeGenOpenCL/opencl_types.cl
clang/test/CodeGenOpenCL/overload.cl
clang/test/CodeGenOpenCL/partial_initializer.cl
clang/test/CodeGenOpenCL/preserve_vec3.cl
clang/test/CodeGenOpenCL/printf.cl
clang/test/CodeGenOpenCL/private-array-initialization.cl
clang/test/CodeGenOpenCL/sampler.cl
clang/test/CodeGenOpenCL/size_t.cl
clang/test/CodeGenOpenCL/spir-calling-conv.cl
clang/test/CodeGenOpenCL/spir-debug-info-pointer-address-space.cl
clang/test/CodeGenOpenCL/spirv32_target.cl
clang/test/CodeGenOpenCL/spirv64_target.cl
clang/test/CodeGenOpenCL/to_addr_builtin.cl
clang/test/CodeGenOpenCL/vectorLoadStore.cl
clang/test/CodeGenOpenCL/vla.cl
clang/test/CodeGenOpenCLCXX/address-space-castoperators.cpp
clang/test/CodeGenOpenCLCXX/address-space-deduction.clcpp
clang/test/CodeGenOpenCLCXX/address-space-deduction2.clcpp
clang/test/CodeGenOpenCLCXX/addrspace-constructors.clcpp
clang/test/CodeGenOpenCLCXX/addrspace-conversion.clcpp
clang/test/CodeGenOpenCLCXX/addrspace-derived-base.clcpp
clang/test/CodeGenOpenCLCXX/addrspace-new-delete.clcpp
clang/test/CodeGenOpenCLCXX/addrspace-of-this.clcpp
clang/test/CodeGenOpenCLCXX/addrspace-operators.clcpp
clang/test/CodeGenOpenCLCXX/addrspace-references.clcpp
clang/test/CodeGenOpenCLCXX/addrspace-with-class.clcpp
clang/test/CodeGenOpenCLCXX/addrspace_cast.clcpp
clang/test/CodeGenOpenCLCXX/atexit.clcpp
clang/test/CodeGenOpenCLCXX/constexpr.clcpp
clang/test/CodeGenOpenCLCXX/global_init.clcpp
clang/test/CodeGenOpenCLCXX/local_addrspace_init.clcpp
clang/test/CodeGenOpenCLCXX/method-overload-address-space.clcpp
clang/test/CodeGenOpenCLCXX/reinterpret_cast.clcpp
clang/test/CodeGenOpenCLCXX/template-address-spaces.clcpp
clang/test/Driver/opencl.cl
clang/test/Headers/opencl-builtins.cl
clang/test/Headers/opencl-c-header.cl
clang/test/Index/pipe-size.cl
clang/test/Preprocessor/predefined-macros.c
clang/test/Sema/Float16.c
clang/test/SemaCXX/Float16.cpp
llvm/include/llvm/ADT/Triple.h
llvm/lib/Support/Triple.cpp
llvm/unittests/ADT/TripleTest.cpp
Index: llvm/unittests/ADT/TripleTest.cpp
===================================================================
--- llvm/unittests/ADT/TripleTest.cpp
+++ 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());
Index: llvm/lib/Support/Triple.cpp
===================================================================
--- llvm/lib/Support/Triple.cpp
+++ 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:
Index: llvm/include/llvm/ADT/Triple.h
===================================================================
--- llvm/include/llvm/ADT/Triple.h
+++ 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;
Index: clang/test/SemaCXX/Float16.cpp
===================================================================
--- clang/test/SemaCXX/Float16.cpp
+++ 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
Index: clang/test/Sema/Float16.c
===================================================================
--- clang/test/Sema/Float16.c
+++ 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
Index: clang/test/Preprocessor/predefined-macros.c
===================================================================
--- clang/test/Preprocessor/predefined-macros.c
+++ 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
Index: clang/test/Index/pipe-size.cl
===================================================================
--- clang/test/Index/pipe-size.cl
+++ 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 )
{
Index: clang/test/Headers/opencl-c-header.cl
===================================================================
--- clang/test/Headers/opencl-c-header.cl
+++ 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)
Index: clang/test/Headers/opencl-builtins.cl
===================================================================
--- clang/test/Headers/opencl-builtins.cl
+++ 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
Index: clang/test/Driver/opencl.cl
===================================================================
--- clang/test/Driver/opencl.cl
+++ 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"
Index: clang/test/CodeGenOpenCLCXX/template-address-spaces.clcpp
===================================================================
--- clang/test/CodeGenOpenCLCXX/template-address-spaces.clcpp
+++ 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 <typename T>
struct S{
Index: clang/test/CodeGenOpenCLCXX/reinterpret_cast.clcpp
===================================================================
--- clang/test/CodeGenOpenCLCXX/reinterpret_cast.clcpp
+++ 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)));
Index: clang/test/CodeGenOpenCLCXX/method-overload-address-space.clcpp
===================================================================
--- clang/test/CodeGenOpenCLCXX/method-overload-address-space.clcpp
+++ 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;
Index: clang/test/CodeGenOpenCLCXX/local_addrspace_init.clcpp
===================================================================
--- clang/test/CodeGenOpenCLCXX/local_addrspace_init.clcpp
+++ 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
Index: clang/test/CodeGenOpenCLCXX/global_init.clcpp
===================================================================
--- clang/test/CodeGenOpenCLCXX/global_init.clcpp
+++ 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() {}
Index: clang/test/CodeGenOpenCLCXX/constexpr.clcpp
===================================================================
--- clang/test/CodeGenOpenCLCXX/constexpr.clcpp
+++ 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)));
Index: clang/test/CodeGenOpenCLCXX/atexit.clcpp
===================================================================
--- clang/test/CodeGenOpenCLCXX/atexit.clcpp
+++ 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(){};
Index: clang/test/CodeGenOpenCLCXX/addrspace_cast.clcpp
===================================================================
--- clang/test/CodeGenOpenCLCXX/addrspace_cast.clcpp
+++ 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) {
Index: clang/test/CodeGenOpenCLCXX/addrspace-with-class.clcpp
===================================================================
--- clang/test/CodeGenOpenCLCXX/addrspace-with-class.clcpp
+++ 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.
Index: clang/test/CodeGenOpenCLCXX/addrspace-references.clcpp
===================================================================
--- clang/test/CodeGenOpenCLCXX/addrspace-references.clcpp
+++ 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)));
Index: clang/test/CodeGenOpenCLCXX/addrspace-operators.clcpp
===================================================================
--- clang/test/CodeGenOpenCLCXX/addrspace-operators.clcpp
+++ 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,
Index: clang/test/CodeGenOpenCLCXX/addrspace-of-this.clcpp
===================================================================
--- clang/test/CodeGenOpenCLCXX/addrspace-of-this.clcpp
+++ 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.
Index: clang/test/CodeGenOpenCLCXX/addrspace-new-delete.clcpp
===================================================================
--- clang/test/CodeGenOpenCLCXX/addrspace-new-delete.clcpp
+++ 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;
Index: clang/test/CodeGenOpenCLCXX/addrspace-derived-base.clcpp
===================================================================
--- clang/test/CodeGenOpenCLCXX/addrspace-derived-base.clcpp
+++ 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;
Index: clang/test/CodeGenOpenCLCXX/addrspace-conversion.clcpp
===================================================================
--- clang/test/CodeGenOpenCLCXX/addrspace-conversion.clcpp
+++ 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)
{
Index: clang/test/CodeGenOpenCLCXX/addrspace-constructors.clcpp
===================================================================
--- clang/test/CodeGenOpenCLCXX/addrspace-constructors.clcpp
+++ 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 }
Index: clang/test/CodeGenOpenCLCXX/address-space-deduction2.clcpp
===================================================================
--- clang/test/CodeGenOpenCLCXX/address-space-deduction2.clcpp
+++ 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:
Index: clang/test/CodeGenOpenCLCXX/address-space-deduction.clcpp
===================================================================
--- clang/test/CodeGenOpenCLCXX/address-space-deduction.clcpp
+++ 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 &
Index: clang/test/CodeGenOpenCLCXX/address-space-castoperators.cpp
===================================================================
--- clang/test/CodeGenOpenCLCXX/address-space-castoperators.cpp
+++ 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;
Index: clang/test/CodeGenOpenCL/vla.cl
===================================================================
--- clang/test/CodeGenOpenCL/vla.cl
+++ 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;
Index: clang/test/CodeGenOpenCL/vectorLoadStore.cl
===================================================================
--- clang/test/CodeGenOpenCL/vectorLoadStore.cl
+++ 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)));
Index: clang/test/CodeGenOpenCL/to_addr_builtin.cl
===================================================================
--- clang/test/CodeGenOpenCL/to_addr_builtin.cl
+++ 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 {
Index: clang/test/CodeGenOpenCL/spirv64_target.cl
===================================================================
--- /dev/null
+++ 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)*
+}
Index: clang/test/CodeGenOpenCL/spirv32_target.cl
===================================================================
--- /dev/null
+++ 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)*
+}
Index: clang/test/CodeGenOpenCL/spir-debug-info-pointer-address-space.cl
===================================================================
--- clang/test/CodeGenOpenCL/spir-debug-info-pointer-address-space.cl
+++ 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)
Index: clang/test/CodeGenOpenCL/spir-calling-conv.cl
===================================================================
--- clang/test/CodeGenOpenCL/spir-calling-conv.cl
+++ 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);
Index: clang/test/CodeGenOpenCL/size_t.cl
===================================================================
--- clang/test/CodeGenOpenCL/size_t.cl
+++ 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
Index: clang/test/CodeGenOpenCL/sampler.cl
===================================================================
--- clang/test/CodeGenOpenCL/sampler.cl
+++ 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
Index: clang/test/CodeGenOpenCL/private-array-initialization.cl
===================================================================
--- clang/test/CodeGenOpenCL/private-array-initialization.cl
+++ 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
Index: clang/test/CodeGenOpenCL/printf.cl
===================================================================
--- clang/test/CodeGenOpenCL/printf.cl
+++ 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;
Index: clang/test/CodeGenOpenCL/preserve_vec3.cl
===================================================================
--- clang/test/CodeGenOpenCL/preserve_vec3.cl
+++ 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)));
Index: clang/test/CodeGenOpenCL/partial_initializer.cl
===================================================================
--- clang/test/CodeGenOpenCL/partial_initializer.cl
+++ 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;
Index: clang/test/CodeGenOpenCL/overload.cl
===================================================================
--- clang/test/CodeGenOpenCL/overload.cl
+++ 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)));
Index: clang/test/CodeGenOpenCL/opencl_types.cl
===================================================================
--- clang/test/CodeGenOpenCL/opencl_types.cl
+++ 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
Index: clang/test/CodeGenOpenCL/kernel-arg-info.cl
===================================================================
--- clang/test/CodeGenOpenCL/kernel-arg-info.cl
+++ 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,
Index: clang/test/CodeGenOpenCL/half.cl
===================================================================
--- clang/test/CodeGenOpenCL/half.cl
+++ 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
Index: clang/test/CodeGenOpenCL/fpmath.cl
===================================================================
--- clang/test/CodeGenOpenCL/fpmath.cl
+++ 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;
Index: clang/test/CodeGenOpenCL/fdeclare-opencl-builtins.cl
===================================================================
--- clang/test/CodeGenOpenCL/fdeclare-opencl-builtins.cl
+++ 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
Index: clang/test/CodeGenOpenCL/enqueue-kernel-non-entry-block.cl
===================================================================
--- clang/test/CodeGenOpenCL/enqueue-kernel-non-entry-block.cl
+++ 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: "<stdin>"
// 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]])
Index: clang/test/CodeGenOpenCL/constant-addr-space-globals.cl
===================================================================
--- clang/test/CodeGenOpenCL/constant-addr-space-globals.cl
+++ 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};
Index: clang/test/CodeGenOpenCL/cl20-device-side-enqueue.cl
===================================================================
--- clang/test/CodeGenOpenCL/cl20-device-side-enqueue.cl
+++ 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
Index: clang/test/CodeGenOpenCL/cast_image.cl
===================================================================
--- clang/test/CodeGenOpenCL/cast_image.cl
+++ 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__
Index: clang/test/CodeGenOpenCL/blocks.cl
===================================================================
--- clang/test/CodeGenOpenCL/blocks.cl
+++ 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)* }
Index: clang/test/CodeGenOpenCL/atomic-ops-libcall.cl
===================================================================
--- clang/test/CodeGenOpenCL/atomic-ops-libcall.cl
+++ 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,
Index: clang/test/CodeGenOpenCL/as_type.cl
===================================================================
--- clang/test/CodeGenOpenCL/as_type.cl
+++ 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;
Index: clang/test/CodeGenCXX/builtin-calling-conv.cpp
===================================================================
--- clang/test/CodeGenCXX/builtin-calling-conv.cpp
+++ 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
Index: clang/test/CodeGen/target-data.c
===================================================================
--- clang/test/CodeGen/target-data.c
+++ 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 | \
Index: clang/test/CodeGen/spir-half-type.cpp
===================================================================
--- clang/test/CodeGen/spir-half-type.cpp
+++ 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
Index: clang/test/CodeGen/ext-int-cc.c
===================================================================
--- clang/test/CodeGen/ext-int-cc.c
+++ 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
Index: clang/test/CodeGen/complex-math.c
===================================================================
--- clang/test/CodeGen/complex-math.c
+++ 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(
Index: clang/lib/Headers/opencl-c.h
===================================================================
--- clang/lib/Headers/opencl-c.h
+++ 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))
Index: clang/lib/Headers/opencl-c-base.h
===================================================================
--- clang/lib/Headers/opencl-c-base.h
+++ 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)
Index: clang/lib/Frontend/InitPreprocessor.cpp
===================================================================
--- clang/lib/Frontend/InitPreprocessor.cpp
+++ 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__");
}
Index: clang/lib/Driver/ToolChains/Clang.cpp
===================================================================
--- clang/lib/Driver/ToolChains/Clang.cpp
+++ 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.
Index: clang/lib/CodeGen/TargetInfo.cpp
===================================================================
--- clang/lib/CodeGen/TargetInfo.cpp
+++ 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));
Index: clang/lib/Basic/Targets/SPIR.h
===================================================================
--- clang/lib/Basic/Targets/SPIR.h
+++ 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
Index: clang/lib/Basic/Targets/SPIR.cpp
===================================================================
--- clang/lib/Basic/Targets/SPIR.cpp
+++ 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);
+}
Index: clang/lib/Basic/Targets.cpp
===================================================================
--- clang/lib/Basic/Targets.cpp
+++ 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 ||
Index: clang/include/clang/Basic/DiagnosticGroups.td
===================================================================
--- clang/include/clang/Basic/DiagnosticGroups.td
+++ 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">;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
