[clang] ebd9753 - [CodeGenOpenCLCXX] Convert tests to opaque pointers (NFC)
Nikita Popov via cfe-commits
cfe-commits at lists.llvm.org
Thu Jan 5 02:03:23 PST 2023
Author: Nikita Popov
Date: 2023-01-05T11:03:15+01:00
New Revision: ebd97534e71aafaa637e466d700f66bbfa63d56b
URL: https://github.com/llvm/llvm-project/commit/ebd97534e71aafaa637e466d700f66bbfa63d56b
DIFF: https://github.com/llvm/llvm-project/commit/ebd97534e71aafaa637e466d700f66bbfa63d56b.diff
LOG: [CodeGenOpenCLCXX] Convert tests to opaque pointers (NFC)
Added:
Modified:
clang/test/CodeGenOpenCLCXX/address-space-deduction.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/method-overload-address-space.clcpp
clang/test/CodeGenOpenCLCXX/reinterpret_cast.clcpp
clang/test/CodeGenOpenCLCXX/template-address-spaces.clcpp
Removed:
################################################################################
diff --git a/clang/test/CodeGenOpenCLCXX/address-space-deduction.clcpp b/clang/test/CodeGenOpenCLCXX/address-space-deduction.clcpp
index 40bac730fd00b..5d7360f406dc3 100644
--- a/clang/test/CodeGenOpenCLCXX/address-space-deduction.clcpp
+++ b/clang/test/CodeGenOpenCLCXX/address-space-deduction.clcpp
@@ -1,5 +1,5 @@
-// RUN: %clang_cc1 -no-opaque-pointers %s -triple spir-unknown-unknown -O0 -emit-llvm -o - | FileCheck %s -check-prefixes=COMMON,PTR
-// RUN: %clang_cc1 -no-opaque-pointers %s -triple spir-unknown-unknown -O0 -emit-llvm -o - -DREF | FileCheck %s -check-prefixes=COMMON,REF
+// 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
#ifdef REF
#define PTR &
@@ -11,26 +11,26 @@
//COMMON: @glob ={{.*}} addrspace(1) global i32
int glob;
-//PTR: @glob_p ={{.*}} addrspace(1) global i32 addrspace(4)* addrspacecast (i32 addrspace(1)* @glob to i32 addrspace(4)*)
-//REF: @glob_p ={{.*}} addrspace(1) constant i32 addrspace(4)* addrspacecast (i32 addrspace(1)* @glob to i32 addrspace(4)*)
+//PTR: @glob_p ={{.*}} addrspace(1) global ptr addrspace(4) addrspacecast (ptr addrspace(1) @glob to ptr addrspace(4))
+//REF: @glob_p ={{.*}} addrspace(1) constant ptr addrspace(4) addrspacecast (ptr addrspace(1) @glob to ptr addrspace(4))
int PTR glob_p = ADR(glob);
//COMMON: @_ZZ3fooi{{P|R}}U3AS4iE6loc_st = internal addrspace(1) global i32
-//PTR: @_ZZ3fooiPU3AS4iE8loc_st_p = internal addrspace(1) global i32 addrspace(4)* addrspacecast (i32 addrspace(1)* @_ZZ3fooiPU3AS4iE6loc_st to i32 addrspace(4)*)
-//REF: @_ZZ3fooiRU3AS4iE8loc_st_p = internal addrspace(1) constant i32 addrspace(4)* addrspacecast (i32 addrspace(1)* @_ZZ3fooiRU3AS4iE6loc_st to i32 addrspace(4)*)
-//COMMON: @loc_ext_p = external addrspace(1) {{global|constant}} i32 addrspace(4)*
+//PTR: @_ZZ3fooiPU3AS4iE8loc_st_p = internal addrspace(1) global ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZZ3fooiPU3AS4iE6loc_st to ptr addrspace(4))
+//REF: @_ZZ3fooiRU3AS4iE8loc_st_p = internal addrspace(1) constant ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZZ3fooiRU3AS4iE6loc_st to ptr addrspace(4))
+//COMMON: @loc_ext_p = external addrspace(1) {{global|constant}} ptr addrspace(4)
//COMMON: @loc_ext = external addrspace(1) global i32
-//COMMON: define{{.*}} spir_func noundef i32 @_Z3fooi{{P|R}}U3AS4i(i32 noundef %par, i32 addrspace(4)*{{.*}} %par_p)
+//COMMON: define{{.*}} spir_func noundef i32 @_Z3fooi{{P|R}}U3AS4i(i32 noundef %par, ptr addrspace(4){{.*}} %par_p)
int foo(int par, int PTR par_p){
//COMMON: %loc = alloca i32
int loc;
- //COMMON: %loc_p = alloca i32 addrspace(4)*
- //COMMON: %loc_p_const = alloca i32*
- //COMMON: [[GAS:%[._a-z0-9]*]] ={{.*}} addrspacecast i32* %loc to i32 addrspace(4)*
- //COMMON: store i32 addrspace(4)* [[GAS]], i32 addrspace(4)** %loc_p
+ //COMMON: %loc_p = alloca ptr addrspace(4)
+ //COMMON: %loc_p_const = alloca ptr
+ //COMMON: [[GAS:%[._a-z0-9]*]] ={{.*}} addrspacecast ptr %loc to ptr addrspace(4)
+ //COMMON: store ptr addrspace(4) [[GAS]], ptr %loc_p
int PTR loc_p = ADR(loc);
- //COMMON: store i32* %loc, i32** %loc_p_const
+ //COMMON: store ptr %loc, ptr %loc_p_const
const __private int PTR loc_p_const = ADR(loc);
// CHECK directives for the following code are located above.
diff --git a/clang/test/CodeGenOpenCLCXX/addrspace-constructors.clcpp b/clang/test/CodeGenOpenCLCXX/addrspace-constructors.clcpp
index 3a87d47acc839..62258cae90647 100644
--- a/clang/test/CodeGenOpenCLCXX/addrspace-constructors.clcpp
+++ b/clang/test/CodeGenOpenCLCXX/addrspace-constructors.clcpp
@@ -1,6 +1,6 @@
-// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=clc++1.0 -DGENERIC -emit-llvm -o - -O0 -triple spir-unknown-unknown | FileCheck %s
-// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=clc++2021 -DGENERIC -emit-llvm -o - -O0 -triple spir-unknown-unknown | FileCheck %s
-// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=clc++2021 -cl-ext=-all,+__opencl_c_program_scope_global_variables -emit-llvm -o - -O0 -triple spir-unknown-unknown | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=clc++1.0 -DGENERIC -emit-llvm -o - -O0 -triple spir-unknown-unknown | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -DGENERIC -emit-llvm -o - -O0 -triple spir-unknown-unknown | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -cl-ext=-all,+__opencl_c_program_scope_global_variables -emit-llvm -o - -O0 -triple spir-unknown-unknown | FileCheck %s
// CHECK: %struct.X = type { i32 }
@@ -35,7 +35,7 @@ __global X gx;
kernel void k() {
// Check that the constructor for px calls the __private constructor.
// CHECK: %px = alloca %struct.X
- // CHECK-NEXT: call spir_func void @_ZN1XC1Ev(%struct.X* {{.*}}%px)
+ // CHECK-NEXT: call spir_func void @_ZN1XC1Ev(ptr {{.*}}%px)
__private X px;
__constant X cx1;
diff --git a/clang/test/CodeGenOpenCLCXX/addrspace-conversion.clcpp b/clang/test/CodeGenOpenCLCXX/addrspace-conversion.clcpp
index 347df195cb00d..2940afb092800 100644
--- a/clang/test/CodeGenOpenCLCXX/addrspace-conversion.clcpp
+++ b/clang/test/CodeGenOpenCLCXX/addrspace-conversion.clcpp
@@ -1,7 +1,7 @@
-// RUN: %clang_cc1 -no-opaque-pointers %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
void bar(__generic volatile unsigned int* ptr)
{
- //CHECK: addrspacecast i32 addrspace(4)* %{{.+}} to i32 addrspace(1)*
+ //CHECK: addrspacecast ptr addrspace(4) %{{.+}} to ptr addrspace(1)
auto gptr = (__global volatile unsigned int*)ptr;
}
diff --git a/clang/test/CodeGenOpenCLCXX/addrspace-derived-base.clcpp b/clang/test/CodeGenOpenCLCXX/addrspace-derived-base.clcpp
index de3621c29434a..808d9e31292aa 100644
--- a/clang/test/CodeGenOpenCLCXX/addrspace-derived-base.clcpp
+++ b/clang/test/CodeGenOpenCLCXX/addrspace-derived-base.clcpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -no-opaque-pointers %s -triple spir -emit-llvm -O0 -o - | FileCheck %s
+// RUN: %clang_cc1 %s -triple spir -emit-llvm -O0 -o - | FileCheck %s
struct B {
int mb;
@@ -12,15 +12,14 @@ public:
void foo() {
D d;
//CHECK-LABEL: foo
- //CHECK: addrspacecast %class.D* %d to %class.D addrspace(4)*
- //CHECK: call spir_func noundef i32 @_ZNU3AS41D5getmbEv(%class.D addrspace(4)*
+ //CHECK: addrspacecast ptr %d to ptr addrspace(4)
+ //CHECK: call spir_func noundef i32 @_ZNU3AS41D5getmbEv(ptr addrspace(4)
d.getmb();
}
//Derived and Base are in the same address space.
-//CHECK: define linkonce_odr spir_func noundef i32 @_ZNU3AS41D5getmbEv(%class.D addrspace(4)* {{[^,]*}} %this)
-//CHECK: bitcast %class.D addrspace(4)* %this1 to %struct.B addrspace(4)*
+//CHECK: define linkonce_odr spir_func noundef i32 @_ZNU3AS41D5getmbEv(ptr addrspace(4) {{[^,]*}} %this)
// Calling base method through multiple inheritance.
@@ -36,7 +35,6 @@ class Derived : public B, public B2 {
public:
void work() const { baseMethod(); }
// CHECK-LABEL: work
- // CHECK: bitcast i8 addrspace(4)* %add.ptr to %class.B2 addrspace(4)*
};
void pr43145(const Derived *argDerived) {
@@ -48,7 +46,6 @@ void pr43145(const Derived *argDerived) {
void pr43145_2(B *argB) {
Derived *x = (Derived*)argB;
// CHECK-LABEL: @_Z9pr43145_2
- // CHECK: bitcast %struct.B addrspace(4)* %0 to %class.Derived addrspace(4)*
}
// Assigning to reference returned by base class method through derived class.
@@ -58,15 +55,13 @@ void pr43145_3(int n) {
d.getRef() = n;
// CHECK-LABEL: @_Z9pr43145_3
- // CHECK: addrspacecast %class.Derived* %d to %class.Derived addrspace(4)*
- // CHECK: bitcast i8 addrspace(4)* %add.ptr to %class.B2 addrspace(4)*
+ // CHECK: addrspacecast ptr %d to ptr addrspace(4)
// CHECK: call {{.*}} @_ZNU3AS42B26getRefEv
private Derived *pd = &d;
pd->getRef() = n;
- // CHECK: addrspacecast %class.Derived* %4 to %class.Derived addrspace(4)*
- // CHECK: bitcast i8 addrspace(4)* %add.ptr1 to %class.B2 addrspace(4)*
+ // CHECK: addrspacecast ptr %2 to ptr addrspace(4)
// CHECK: call {{.*}} @_ZNU3AS42B26getRefEv
}
diff --git a/clang/test/CodeGenOpenCLCXX/addrspace-new-delete.clcpp b/clang/test/CodeGenOpenCLCXX/addrspace-new-delete.clcpp
index 0316df455eb56..ad5948a0b4e67 100644
--- a/clang/test/CodeGenOpenCLCXX/addrspace-new-delete.clcpp
+++ b/clang/test/CodeGenOpenCLCXX/addrspace-new-delete.clcpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -no-opaque-pointers %s -triple spir -emit-llvm -O0 -o - | FileCheck %s
+// RUN: %clang_cc1 %s -triple spir -emit-llvm -O0 -o - | FileCheck %s
typedef __SIZE_TYPE__ size_t;
@@ -9,8 +9,8 @@ public:
};
void test_new_delete(A **a) {
-// CHECK: %{{.*}} = call spir_func noundef i8 addrspace(4)* @_ZNU3AS41AnwEj(i32 {{.*}})
+// CHECK: %{{.*}} = call spir_func noundef ptr addrspace(4) @_ZNU3AS41AnwEj(i32 {{.*}})
*a = new A;
-// CHECK: call spir_func void @_ZNU3AS41AdlEPU3AS4v(i8 addrspace(4)* {{.*}})
+// CHECK: call spir_func void @_ZNU3AS41AdlEPU3AS4v(ptr addrspace(4) {{.*}})
delete *a;
}
diff --git a/clang/test/CodeGenOpenCLCXX/addrspace-of-this.clcpp b/clang/test/CodeGenOpenCLCXX/addrspace-of-this.clcpp
index a759d4bd5d565..1c09743d3c8db 100644
--- a/clang/test/CodeGenOpenCLCXX/addrspace-of-this.clcpp
+++ b/clang/test/CodeGenOpenCLCXX/addrspace-of-this.clcpp
@@ -1,6 +1,6 @@
-// RUN: %clang_cc1 -no-opaque-pointers %s -triple spir-unknown-unknown -emit-llvm -pedantic -verify -O0 -o - -DDECL | FileCheck %s --check-prefixes="COMMON,EXPL"
-// RUN: %clang_cc1 -no-opaque-pointers %s -triple spir-unknown-unknown -emit-llvm -pedantic -verify -O0 -o - -DDECL -DUSE_DEFLT | FileCheck %s --check-prefixes="COMMON,IMPL"
-// RUN: %clang_cc1 -no-opaque-pointers %s -triple spir-unknown-unknown -emit-llvm -pedantic -verify -O0 -o - | FileCheck %s --check-prefixes="COMMON,IMPL"
+// 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"
// expected-no-diagnostics
// Test that the 'this' pointer is in the __generic address space.
@@ -70,71 +70,64 @@ __kernel void test__global() {
}
// Test that the address space is __generic for all members
-// EXPL: @_ZNU3AS41CC2Ev(%class.C addrspace(4)* {{[^,]*}} %this)
-// EXPL: @_ZNU3AS41CC1Ev(%class.C addrspace(4)* {{[^,]*}} %this)
-// EXPL: @_ZNU3AS41CC2EOU3AS4S_(%class.C addrspace(4)* {{[^,]*}} %this
-// EXPL: @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* {{[^,]*}} %this
-// EXPL: @_ZNU3AS41CC2ERU3AS4KS_(%class.C addrspace(4)* {{[^,]*}} %this
-// EXPL: @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* {{[^,]*}} %this
-// EXPL: @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* {{[^,]*}} %this
-// EXPL: @_ZNU3AS4R1CaSEOU3AS4S_(%class.C addrspace(4)* {{[^,]*}} %this
-// COMMON: @_ZNU3AS41C7outsideEv(%class.C addrspace(4)* {{[^,]*}} %this)
+// EXPL: @_ZNU3AS41CC2Ev(ptr addrspace(4) {{[^,]*}} %this)
+// EXPL: @_ZNU3AS41CC1Ev(ptr addrspace(4) {{[^,]*}} %this)
+// EXPL: @_ZNU3AS41CC2EOU3AS4S_(ptr addrspace(4) {{[^,]*}} %this
+// EXPL: @_ZNU3AS41CC1EOU3AS4S_(ptr addrspace(4) {{[^,]*}} %this
+// EXPL: @_ZNU3AS41CC2ERU3AS4KS_(ptr addrspace(4) {{[^,]*}} %this
+// EXPL: @_ZNU3AS41CC1ERU3AS4KS_(ptr addrspace(4) {{[^,]*}} %this
+// EXPL: @_ZNU3AS41CaSERU3AS4KS_(ptr addrspace(4) {{[^,]*}} %this
+// EXPL: @_ZNU3AS4R1CaSEOU3AS4S_(ptr addrspace(4) {{[^,]*}} %this
+// COMMON: @_ZNU3AS41C7outsideEv(ptr addrspace(4) {{[^,]*}} %this)
// EXPL-LABEL: @__cxx_global_var_init()
-// EXPL: call spir_func void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* {{[^,]*}} addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
+// EXPL: call spir_func void @_ZNU3AS41CC1Ev(ptr addrspace(4) {{[^,]*}} addrspacecast (ptr addrspace(1) @c to ptr addrspace(4)))
// COMMON-LABEL: @test__global()
// Test the address space of 'this' when invoking a method.
-// COMMON: call spir_func noundef i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* {{[^,]*}} addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
+// COMMON: call spir_func noundef i32 @_ZNU3AS41C3getEv(ptr addrspace(4) {{[^,]*}} addrspacecast (ptr addrspace(1) @c to ptr addrspace(4)))
// Test the address space of 'this' when invoking a method using a pointer to the object.
-// COMMON: call spir_func noundef i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* {{[^,]*}} addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
+// COMMON: call spir_func noundef i32 @_ZNU3AS41C3getEv(ptr addrspace(4) {{[^,]*}} addrspacecast (ptr addrspace(1) @c to ptr addrspace(4)))
// Test the address space of 'this' when invoking a method that is declared in the file contex.
-// COMMON: call spir_func noundef i32 @_ZNU3AS41C7outsideEv(%class.C addrspace(4)* {{[^,]*}} addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
+// COMMON: call spir_func noundef i32 @_ZNU3AS41C7outsideEv(ptr addrspace(4) {{[^,]*}} addrspacecast (ptr addrspace(1) @c to ptr addrspace(4)))
// Test the address space of 'this' when invoking copy-constructor.
-// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
-// IMPL: [[C1VOID:%[0-9]+]] = bitcast %class.C* %c1 to i8*
-// IMPL: call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C1VOID]], i8 addrspace(4)* {{.*}}addrspacecast (i8 addrspace(1)* bitcast (%class.C addrspace(1)* @c to i8 addrspace(1)*) to i8 addrspace(4)*)
-// EXPL: call spir_func void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* {{[^,]*}} [[C1GEN]], %class.C addrspace(4)* noundef align 4 dereferenceable(4) addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
+// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast ptr %c1 to ptr addrspace(4)
+// IMPL: call void @llvm.memcpy.p0.p4.i32(ptr {{.*}}%c1, ptr addrspace(4) {{.*}}addrspacecast (ptr addrspace(1) @c to ptr addrspace(4))
+// EXPL: call spir_func void @_ZNU3AS41CC1ERU3AS4KS_(ptr addrspace(4) {{[^,]*}} [[C1GEN]], ptr addrspace(4) noundef align 4 dereferenceable(4) addrspacecast (ptr addrspace(1) @c to ptr addrspace(4)))
// Test the address space of 'this' when invoking a constructor.
-// EXPL: [[C2GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
-// EXPL: call spir_func void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* {{[^,]*}} [[C2GEN]])
+// EXPL: [[C2GEN:%[.a-z0-9]+]] = addrspacecast ptr %c2 to ptr addrspace(4)
+// EXPL: call spir_func void @_ZNU3AS41CC1Ev(ptr addrspace(4) {{[^,]*}} [[C2GEN]])
// Test the address space of 'this' when invoking assignment operator.
-// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
-// COMMON: [[C2GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
-// EXPL: call spir_func noundef align 4 dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* {{[^,]*}} [[C2GEN]], %class.C addrspace(4)* noundef align 4 dereferenceable(4) [[C1GEN]])
-// IMPL: [[C2GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C2GEN]] to i8 addrspace(4)*
-// IMPL: [[C1GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C1GEN]] to i8 addrspace(4)*
-// IMPL: call void @llvm.memcpy.p4i8.p4i8.i32(i8 addrspace(4)* {{.*}}[[C2GENVOID]], i8 addrspace(4)* {{.*}}[[C1GENVOID]]
+// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast ptr %c1 to ptr addrspace(4)
+// COMMON: [[C2GEN:%[.a-z0-9]+]] = addrspacecast ptr %c2 to ptr addrspace(4)
+// EXPL: call spir_func noundef align 4 dereferenceable(4) ptr addrspace(4) @_ZNU3AS41CaSERU3AS4KS_(ptr addrspace(4) {{[^,]*}} [[C2GEN]], ptr addrspace(4) noundef align 4 dereferenceable(4) [[C1GEN]])
+// IMPL: call void @llvm.memcpy.p4.p4.i32(ptr addrspace(4) {{.*}}[[C2GEN]], ptr addrspace(4) {{.*}}[[C1GEN]]
// Test the address space of 'this' when invoking the operator+
-// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
-// COMMON: [[C2GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
-// COMMON: call spir_func void @_ZNU3AS41CplERU3AS4KS_(%class.C* sret(%class.C) align 4 %c3, %class.C addrspace(4)* {{[^,]*}} [[C1GEN]], %class.C addrspace(4)* noundef align 4 dereferenceable(4) [[C2GEN]])
+// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast ptr %c1 to ptr addrspace(4)
+// COMMON: [[C2GEN:%[.a-z0-9]+]] = addrspacecast ptr %c2 to ptr addrspace(4)
+// COMMON: call spir_func void @_ZNU3AS41CplERU3AS4KS_(ptr sret(%class.C) align 4 %c3, ptr addrspace(4) {{[^,]*}} [[C1GEN]], ptr addrspace(4) noundef align 4 dereferenceable(4) [[C2GEN]])
// Test the address space of 'this' when invoking the move constructor
-// COMMON: [[C4GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c4 to %class.C addrspace(4)*
-// COMMON: [[CALL:%call[0-9]+]] = call spir_func noundef align 4 dereferenceable(4) %class.C addrspace(4)* @_Z3foov()
-// EXPL: call spir_func void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* {{[^,]*}} [[C4GEN]], %class.C addrspace(4)* noundef align 4 dereferenceable(4) [[CALL]])
-// IMPL: [[C4VOID:%[0-9]+]] = bitcast %class.C* %c4 to i8*
-// IMPL: [[CALLVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[CALL]] to i8 addrspace(4)*
-// IMPL: call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C4VOID]], i8 addrspace(4)* {{.*}}[[CALLVOID]]
+// COMMON: [[C4GEN:%[.a-z0-9]+]] = addrspacecast ptr %c4 to ptr addrspace(4)
+// COMMON: [[CALL:%call[0-9]+]] = call spir_func noundef align 4 dereferenceable(4) ptr addrspace(4) @_Z3foov()
+// EXPL: call spir_func void @_ZNU3AS41CC1EOU3AS4S_(ptr addrspace(4) {{[^,]*}} [[C4GEN]], ptr addrspace(4) noundef align 4 dereferenceable(4) [[CALL]])
+// IMPL: call void @llvm.memcpy.p0.p4.i32(ptr {{.*}}%c4, ptr addrspace(4) {{.*}}[[CALL]]
// Test the address space of 'this' when invoking the move assignment
-// COMMON: [[C5GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c5 to %class.C addrspace(4)*
-// COMMON: [[CALL:%call[0-9]+]] = call spir_func noundef align 4 dereferenceable(4) %class.C addrspace(4)* @_Z3foov()
-// EXPL: call spir_func void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* {{[^,]*}} [[C5GEN]], %class.C addrspace(4)* noundef align 4 dereferenceable(4) [[CALL]])
-// IMPL: [[C5VOID:%[0-9]+]] = bitcast %class.C* %c5 to i8*
-// IMPL: [[CALLVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[CALL]] to i8 addrspace(4)*
-// IMPL: call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C5VOID]], i8 addrspace(4)* {{.*}}[[CALLVOID]]
+// COMMON: [[C5GEN:%[.a-z0-9]+]] = addrspacecast ptr %c5 to ptr addrspace(4)
+// COMMON: [[CALL:%call[0-9]+]] = call spir_func noundef align 4 dereferenceable(4) ptr addrspace(4) @_Z3foov()
+// EXPL: call spir_func void @_ZNU3AS41CC1EOU3AS4S_(ptr addrspace(4) {{[^,]*}} [[C5GEN]], ptr addrspace(4) noundef align 4 dereferenceable(4) [[CALL]])
+// IMPL: call void @llvm.memcpy.p0.p4.i32(ptr {{.*}}%c5, ptr addrspace(4) {{.*}}[[CALL]]
// Tests address space of inline members
-//COMMON: @_ZNU3AS41C3getEv(%class.C addrspace(4)* {{[^,]*}} %this)
-//COMMON: @_ZNU3AS41CplERU3AS4KS_(%class.C* noalias sret(%class.C) align 4 %agg.result, %class.C addrspace(4)* {{[^,]*}} %this
+//COMMON: @_ZNU3AS41C3getEv(ptr addrspace(4) {{[^,]*}} %this)
+//COMMON: @_ZNU3AS41CplERU3AS4KS_(ptr noalias sret(%class.C) align 4 %agg.result, ptr addrspace(4) {{[^,]*}} %this
#define TEST(AS) \
__kernel void test##AS() { \
AS C c; \
@@ -149,60 +142,53 @@ TEST(__local)
// COMMON-LABEL: @test__local
// Test that we don't initialize an object in local address space.
-// EXPL-NOT: call spir_func void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localE1c to %class.C addrspace(4)*))
+// EXPL-NOT: call spir_func void @_ZNU3AS41CC1Ev(ptr addrspace(4) addrspacecast (ptr addrspace(3) @_ZZ11test__localE1c to ptr addrspace(4)))
// Test the address space of 'this' when invoking a method.
-// COMMON: call spir_func noundef i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* {{[^,]*}} addrspacecast (%class.C addrspace(3)* @_ZZ11test__localE1c to %class.C addrspace(4)*))
+// COMMON: call spir_func noundef i32 @_ZNU3AS41C3getEv(ptr addrspace(4) {{[^,]*}} addrspacecast (ptr addrspace(3) @_ZZ11test__localE1c to ptr addrspace(4)))
// Test the address space of 'this' when invoking copy-constructor.
-// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
-// EXPL: call spir_func void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* {{[^,]*}} [[C1GEN]], %class.C addrspace(4)* noundef align 4 dereferenceable(4) addrspacecast (%class.C addrspace(3)* @_ZZ11test__localE1c to %class.C addrspace(4)*))
-// IMPL: [[C1VOID:%[0-9]+]] = bitcast %class.C* %c1 to i8*
-// IMPL: call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C1VOID]], i8 addrspace(4)* {{.*}}addrspacecast (i8 addrspace(3)* bitcast (%class.C addrspace(3)* @_ZZ11test__localE1c to i8 addrspace(3)*) to i8 addrspace(4)*), i32 4, i1 false)
+// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast ptr %c1 to ptr addrspace(4)
+// EXPL: call spir_func void @_ZNU3AS41CC1ERU3AS4KS_(ptr addrspace(4) {{[^,]*}} [[C1GEN]], ptr addrspace(4) noundef align 4 dereferenceable(4) addrspacecast (ptr addrspace(3) @_ZZ11test__localE1c to ptr addrspace(4)))
+// IMPL: call void @llvm.memcpy.p0.p4.i32(ptr {{.*}}%c1, ptr addrspace(4) {{.*}}addrspacecast (ptr addrspace(3) @_ZZ11test__localE1c to ptr addrspace(4)), i32 4, i1 false)
// Test the address space of 'this' when invoking a constructor.
-// EXPL: [[C2GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
-// EXPL: call spir_func void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* {{[^,]*}} [[C2GEN]])
+// EXPL: [[C2GEN:%[.a-z0-9]+]] = addrspacecast ptr %c2 to ptr addrspace(4)
+// EXPL: call spir_func void @_ZNU3AS41CC1Ev(ptr addrspace(4) {{[^,]*}} [[C2GEN]])
// Test the address space of 'this' when invoking assignment operator.
-// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
-// COMMON: [[C2GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
-// EXPL: call spir_func noundef align 4 dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* {{[^,]*}} [[C2GEN]], %class.C addrspace(4)* noundef align 4 dereferenceable(4) [[C1GEN]])
-// IMPL: [[C2GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C2GEN]] to i8 addrspace(4)*
-// IMPL: [[C1GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C1GEN]] to i8 addrspace(4)*
-// IMPL: call void @llvm.memcpy.p4i8.p4i8.i32(i8 addrspace(4)* {{.*}}[[C2GENVOID]], i8 addrspace(4)* {{.*}}[[C1GENVOID]]
+// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast ptr %c1 to ptr addrspace(4)
+// COMMON: [[C2GEN:%[.a-z0-9]+]] = addrspacecast ptr %c2 to ptr addrspace(4)
+// EXPL: call spir_func noundef align 4 dereferenceable(4) ptr addrspace(4) @_ZNU3AS41CaSERU3AS4KS_(ptr addrspace(4) {{[^,]*}} [[C2GEN]], ptr addrspace(4) noundef align 4 dereferenceable(4) [[C1GEN]])
+// IMPL: call void @llvm.memcpy.p4.p4.i32(ptr addrspace(4) {{.*}}[[C2GEN]], ptr addrspace(4) {{.*}}[[C1GEN]]
TEST(__private)
// COMMON-LABEL: @test__private
// Test the address space of 'this' when invoking a constructor for an object in non-default address space
-// EXPL: [[CGEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
-// EXPL: call spir_func void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* {{[^,]*}} [[CGEN]])
+// EXPL: [[CGEN:%[.a-z0-9]+]] = addrspacecast ptr %c to ptr addrspace(4)
+// EXPL: call spir_func void @_ZNU3AS41CC1Ev(ptr addrspace(4) {{[^,]*}} [[CGEN]])
// Test the address space of 'this' when invoking a method.
-// COMMON: [[CGEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
-// COMMON: call spir_func noundef i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* {{[^,]*}} [[CGEN]])
+// COMMON: [[CGEN:%[.a-z0-9]+]] = addrspacecast ptr %c to ptr addrspace(4)
+// COMMON: call spir_func noundef i32 @_ZNU3AS41C3getEv(ptr addrspace(4) {{[^,]*}} [[CGEN]])
// Test the address space of 'this' when invoking a copy-constructor.
-// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
-// COMMON: [[CGEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
-// EXPL: call spir_func void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* {{[^,]*}} [[C1GEN]], %class.C addrspace(4)* noundef align 4 dereferenceable(4) [[CGEN]])
-// IMPL: [[C1VOID:%[0-9]+]] = bitcast %class.C* %c1 to i8*
-// IMPL: [[CGENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[CGEN]] to i8 addrspace(4)*
-// IMPL: call void @llvm.memcpy.p0i8.p4i8.i32(i8* {{.*}}[[C1VOID]], i8 addrspace(4)* {{.*}}[[CGENVOID]]
+// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast ptr %c1 to ptr addrspace(4)
+// COMMON: [[CGEN:%[.a-z0-9]+]] = addrspacecast ptr %c to ptr addrspace(4)
+// EXPL: call spir_func void @_ZNU3AS41CC1ERU3AS4KS_(ptr addrspace(4) {{[^,]*}} [[C1GEN]], ptr addrspace(4) noundef align 4 dereferenceable(4) [[CGEN]])
+// IMPL: call void @llvm.memcpy.p0.p4.i32(ptr {{.*}}%c1, ptr addrspace(4) {{.*}}[[CGEN]]
// Test the address space of 'this' when invoking a constructor.
-// EXPL: [[C2GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
-// EXPL: call spir_func void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* {{[^,]*}} [[C2GEN]])
+// EXPL: [[C2GEN:%[.a-z0-9]+]] = addrspacecast ptr %c2 to ptr addrspace(4)
+// EXPL: call spir_func void @_ZNU3AS41CC1Ev(ptr addrspace(4) {{[^,]*}} [[C2GEN]])
// Test the address space of 'this' when invoking a copy-assignment.
-// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
-// COMMON: [[C2GEN:%[.a-z0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
-// EXPL: call spir_func noundef align 4 dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* {{[^,]*}} [[C2GEN]], %class.C addrspace(4)* noundef align 4 dereferenceable(4) [[C1GEN]])
-// IMPL: [[C2GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C2GEN]] to i8 addrspace(4)*
-// IMPL: [[C1GENVOID:%[0-9]+]] = bitcast %class.C addrspace(4)* [[C1GEN]] to i8 addrspace(4)*
-// IMPL: call void @llvm.memcpy.p4i8.p4i8.i32(i8 addrspace(4)* {{.*}}[[C2GENVOID]], i8 addrspace(4)* {{.*}}[[C1GENVOID]]
+// COMMON: [[C1GEN:%[.a-z0-9]+]] = addrspacecast ptr %c1 to ptr addrspace(4)
+// COMMON: [[C2GEN:%[.a-z0-9]+]] = addrspacecast ptr %c2 to ptr addrspace(4)
+// EXPL: call spir_func noundef align 4 dereferenceable(4) ptr addrspace(4) @_ZNU3AS41CaSERU3AS4KS_(ptr addrspace(4) {{[^,]*}} [[C2GEN]], ptr addrspace(4) noundef align 4 dereferenceable(4) [[C1GEN]])
+// IMPL: call void @llvm.memcpy.p4.p4.i32(ptr addrspace(4) {{.*}}[[C2GEN]], ptr addrspace(4) {{.*}}[[C1GEN]]
// Test that calling a const method from a non-const method does not crash Clang.
class ConstAndNonConstMethod {
diff --git a/clang/test/CodeGenOpenCLCXX/addrspace-operators.clcpp b/clang/test/CodeGenOpenCLCXX/addrspace-operators.clcpp
index 93358b749a95d..1f9f849468661 100644
--- a/clang/test/CodeGenOpenCLCXX/addrspace-operators.clcpp
+++ b/clang/test/CodeGenOpenCLCXX/addrspace-operators.clcpp
@@ -1,4 +1,4 @@
-//RUN: %clang_cc1 -no-opaque-pointers %s -triple spir -emit-llvm -O0 -o - | FileCheck %s
+//RUN: %clang_cc1 %s -triple spir -emit-llvm -O0 -o - | FileCheck %s
enum E {
a,
@@ -19,49 +19,49 @@ __global int globI;
//CHECK-LABEL: define{{.*}} spir_func void @_Z3barv()
void bar() {
C c;
- //CHECK: [[A1:%[.a-z0-9]+]] ={{.*}} addrspacecast %class.C* [[C:%[a-z0-9]+]] to %class.C addrspace(4)*
- //CHECK: call spir_func void @_ZNU3AS41C6AssignE1E(%class.C addrspace(4)* {{[^,]*}} [[A1]], i32 noundef 0)
+ //CHECK: [[A1:%[.a-z0-9]+]] ={{.*}} addrspacecast ptr [[C:%[a-z0-9]+]] to ptr addrspace(4)
+ //CHECK: call spir_func void @_ZNU3AS41C6AssignE1E(ptr addrspace(4) {{[^,]*}} [[A1]], i32 noundef 0)
c.Assign(a);
- //CHECK: [[A2:%[.a-z0-9]+]] ={{.*}} addrspacecast %class.C* [[C]] to %class.C addrspace(4)*
- //CHECK: call spir_func void @_ZNU3AS41C8OrAssignE1E(%class.C addrspace(4)* {{[^,]*}} [[A2]], i32 noundef 0)
+ //CHECK: [[A2:%[.a-z0-9]+]] ={{.*}} addrspacecast ptr [[C]] to ptr addrspace(4)
+ //CHECK: call spir_func void @_ZNU3AS41C8OrAssignE1E(ptr addrspace(4) {{[^,]*}} [[A2]], i32 noundef 0)
c.OrAssign(a);
E e;
- //CHECK: store i32 1, i32* %e
+ //CHECK: store i32 1, ptr %e
e = b;
- //CHECK: store i32 0, i32 addrspace(1)* @globE
+ //CHECK: store i32 0, ptr addrspace(1) @globE
globE = a;
- //CHECK: store i32 %or, i32 addrspace(1)* @globI
+ //CHECK: store i32 %or, ptr addrspace(1) @globI
globI |= b;
- //CHECK: store i32 %add, i32 addrspace(1)* @globI
+ //CHECK: store i32 %add, ptr addrspace(1) @globI
globI += a;
- //CHECK: [[GVIV1:%[0-9]+]] = load volatile i32, i32 addrspace(1)* @globVI
+ //CHECK: [[GVIV1:%[0-9]+]] = load volatile i32, ptr addrspace(1) @globVI
//CHECK: [[AND:%[a-z0-9]+]] = and i32 [[GVIV1]], 1
- //CHECK: store volatile i32 [[AND]], i32 addrspace(1)* @globVI
+ //CHECK: store volatile i32 [[AND]], ptr addrspace(1) @globVI
globVI &= b;
- //CHECK: [[GVIV2:%[0-9]+]] = load volatile i32, i32 addrspace(1)* @globVI
+ //CHECK: [[GVIV2:%[0-9]+]] = load volatile i32, ptr addrspace(1) @globVI
//CHECK: [[SUB:%[a-z0-9]+]] = sub nsw i32 [[GVIV2]], 0
- //CHECK: store volatile i32 [[SUB]], i32 addrspace(1)* @globVI
+ //CHECK: store volatile i32 [[SUB]], ptr addrspace(1) @globVI
globVI -= a;
}
-//CHECK: define linkonce_odr spir_func void @_ZNU3AS41C6AssignE1E(%class.C addrspace(4)* {{[^,]*}} {{[ %a-z0-9]*}}, i32{{[ %a-z0-9]*}})
-//CHECK: [[THIS_ADDR:%[.a-z0-9]+]] = alloca %class.C addrspace(4)
+//CHECK: define linkonce_odr spir_func void @_ZNU3AS41C6AssignE1E(ptr addrspace(4) {{[^,]*}} {{[ %a-z0-9]*}}, i32{{[ %a-z0-9]*}})
+//CHECK: [[THIS_ADDR:%[.a-z0-9]+]] = alloca ptr addrspace(4)
//CHECK: [[E_ADDR:%[.a-z0-9]+]] = alloca i32
-//CHECK: store %class.C addrspace(4)* {{%[a-z0-9]+}}, %class.C addrspace(4)** [[THIS_ADDR]]
-//CHECK: store i32 {{%[a-z0-9]+}}, i32* [[E_ADDR]]
-//CHECK: [[THIS1:%[.a-z0-9]+]] = load %class.C addrspace(4)*, %class.C addrspace(4)** [[THIS_ADDR]]
-//CHECK: [[E:%[0-9]+]] = load i32, i32* [[E_ADDR]]
-//CHECK: [[ME:%[a-z0-9]+]] = getelementptr inbounds %class.C, %class.C addrspace(4)* [[THIS1]], i32 0, i32 0
-//CHECK: store i32 [[E]], i32 addrspace(4)* [[ME]]
+//CHECK: store ptr addrspace(4) {{%[a-z0-9]+}}, ptr [[THIS_ADDR]]
+//CHECK: store i32 {{%[a-z0-9]+}}, ptr [[E_ADDR]]
+//CHECK: [[THIS1:%[.a-z0-9]+]] = load ptr addrspace(4), ptr [[THIS_ADDR]]
+//CHECK: [[E:%[0-9]+]] = load i32, ptr [[E_ADDR]]
+//CHECK: [[ME:%[a-z0-9]+]] = getelementptr inbounds %class.C, ptr addrspace(4) [[THIS1]], i32 0, i32 0
+//CHECK: store i32 [[E]], ptr addrspace(4) [[ME]]
-//CHECK: define linkonce_odr spir_func void @_ZNU3AS41C8OrAssignE1E(%class.C addrspace(4)* {{[^,]*}} {{[ %a-z0-9]*}}, i32{{[ %a-z0-9]*}})
-//CHECK: [[THIS_ADDR:%[.a-z0-9]+]] = alloca %class.C addrspace(4)
+//CHECK: define linkonce_odr spir_func void @_ZNU3AS41C8OrAssignE1E(ptr addrspace(4) {{[^,]*}} {{[ %a-z0-9]*}}, i32{{[ %a-z0-9]*}})
+//CHECK: [[THIS_ADDR:%[.a-z0-9]+]] = alloca ptr addrspace(4)
//CHECK: [[E_ADDR:%[.a-z0-9]+]] = alloca i32
-//CHECK: store %class.C addrspace(4)* {{%[a-z0-9]+}}, %class.C addrspace(4)** [[THIS_ADDR]]
-//CHECK: store i32 {{%[a-z0-9]+}}, i32* [[E_ADDR]]
-//CHECK: [[THIS1:%[.a-z0-9]+]] = load %class.C addrspace(4)*, %class.C addrspace(4)** [[THIS_ADDR]]
-//CHECK: [[E:%[0-9]+]] = load i32, i32* [[E_ADDR]]
-//CHECK: [[MI_GEP:%[a-z0-9]+]] = getelementptr inbounds %class.C, %class.C addrspace(4)* [[THIS1]], i32 0, i32 1
-//CHECK: [[MI:%[0-9]+]] = load i32, i32 addrspace(4)* [[MI_GEP]]
+//CHECK: store ptr addrspace(4) {{%[a-z0-9]+}}, ptr [[THIS_ADDR]]
+//CHECK: store i32 {{%[a-z0-9]+}}, ptr [[E_ADDR]]
+//CHECK: [[THIS1:%[.a-z0-9]+]] = load ptr addrspace(4), ptr [[THIS_ADDR]]
+//CHECK: [[E:%[0-9]+]] = load i32, ptr [[E_ADDR]]
+//CHECK: [[MI_GEP:%[a-z0-9]+]] = getelementptr inbounds %class.C, ptr addrspace(4) [[THIS1]], i32 0, i32 1
+//CHECK: [[MI:%[0-9]+]] = load i32, ptr addrspace(4) [[MI_GEP]]
//CHECK: %or = or i32 [[MI]], [[E]]
diff --git a/clang/test/CodeGenOpenCLCXX/addrspace-references.clcpp b/clang/test/CodeGenOpenCLCXX/addrspace-references.clcpp
index 7a092e3e3d1c5..de840654bf74f 100644
--- a/clang/test/CodeGenOpenCLCXX/addrspace-references.clcpp
+++ b/clang/test/CodeGenOpenCLCXX/addrspace-references.clcpp
@@ -1,4 +1,4 @@
-//RUN: %clang_cc1 -no-opaque-pointers %s -triple spir -emit-llvm -o - -O0 | FileCheck %s
+//RUN: %clang_cc1 %s -triple spir -emit-llvm -o - -O0 | FileCheck %s
typedef short short2 __attribute__((ext_vector_type(2)));
@@ -15,9 +15,9 @@ void scalar() {
// to a temporary value allocated in private addr space. We need an
// addrspacecast before passing the value to the function.
// CHECK: [[REF:%.*]] = alloca i32
- // CHECK: store i32 1, i32* [[REF]]
- // CHECK: [[REG:%[.a-z0-9]+]] ={{.*}} addrspacecast i32* [[REF]] to i32 addrspace(4)*
- // CHECK: call spir_func noundef i32 @_Z3barRU3AS4Kj(i32 addrspace(4)* noundef align 4 dereferenceable(4) [[REG]])
+ // CHECK: store i32 1, ptr [[REF]]
+ // CHECK: [[REG:%[.a-z0-9]+]] ={{.*}} addrspacecast ptr [[REF]] to ptr addrspace(4)
+ // CHECK: call spir_func noundef i32 @_Z3barRU3AS4Kj(ptr addrspace(4) noundef align 4 dereferenceable(4) [[REG]])
bar(1);
}
@@ -26,8 +26,8 @@ void scalar() {
void list() {
C c1;
// CHECK: [[REF:%.*]] = alloca <2 x i16>
-// CHECK: store <2 x i16> <i16 1, i16 2>, <2 x i16>* [[REF]]
-// CHECK: [[REG:%[.a-z0-9]+]] = addrspacecast <2 x i16>* [[REF]] to <2 x i16> addrspace(4)*
-// CHECK: call {{.*}}void @_ZNU3AS41C3barERU3AS4KDv2_s(%class.C addrspace(4)* {{.*}}, <2 x i16> addrspace(4)*{{.*}} [[REG]])
+// CHECK: store <2 x i16> <i16 1, i16 2>, ptr [[REF]]
+// CHECK: [[REG:%[.a-z0-9]+]] = addrspacecast ptr [[REF]] to ptr addrspace(4)
+// CHECK: call {{.*}}void @_ZNU3AS41C3barERU3AS4KDv2_s(ptr addrspace(4) {{.*}}, ptr addrspace(4){{.*}} [[REG]])
c1.bar({1, 2});
}
diff --git a/clang/test/CodeGenOpenCLCXX/addrspace-with-class.clcpp b/clang/test/CodeGenOpenCLCXX/addrspace-with-class.clcpp
index 8b79bf53b53ff..18d97a989a436 100644
--- a/clang/test/CodeGenOpenCLCXX/addrspace-with-class.clcpp
+++ b/clang/test/CodeGenOpenCLCXX/addrspace-with-class.clcpp
@@ -1,5 +1,5 @@
-// RUN: %clang_cc1 -no-opaque-pointers %s -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s
-// RUN: %clang_cc1 -no-opaque-pointers %s -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s --check-prefix=CHECK-DEFINITIONS
+// 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
// This test ensures the proper address spaces and address space cast are used
// for constructors, member functions and destructors.
@@ -23,37 +23,37 @@ __constant MyType const2(2);
// CHECK: @glob ={{.*}} addrspace(1) global %struct.MyType zeroinitializer
MyType glob(1);
-// CHECK: call spir_func void @_ZNU3AS26MyTypeC1Ei(%struct.MyType addrspace(2)* {{[^,]*}} @const1, i32 noundef 1)
-// CHECK: call spir_func void @_ZNU3AS26MyTypeC1Ei(%struct.MyType addrspace(2)* {{[^,]*}} @const2, i32 noundef 2)
-// CHECK: call spir_func void @_ZNU3AS46MyTypeC1Ei(%struct.MyType addrspace(4)* {{[^,]*}} addrspacecast (%struct.MyType addrspace(1)* @glob to %struct.MyType addrspace(4)*), i32 noundef 1)
+// CHECK: call spir_func void @_ZNU3AS26MyTypeC1Ei(ptr addrspace(2) {{[^,]*}} @const1, i32 noundef 1)
+// CHECK: call spir_func void @_ZNU3AS26MyTypeC1Ei(ptr addrspace(2) {{[^,]*}} @const2, i32 noundef 2)
+// CHECK: call spir_func void @_ZNU3AS46MyTypeC1Ei(ptr addrspace(4) {{[^,]*}} addrspacecast (ptr addrspace(1) @glob to ptr addrspace(4)), i32 noundef 1)
// CHECK-LABEL: define{{.*}} spir_kernel void @fooGlobal()
kernel void fooGlobal() {
- // CHECK: call spir_func noundef i32 @_ZNU3AS46MyType3barEv(%struct.MyType addrspace(4)* {{[^,]*}} addrspacecast (%struct.MyType addrspace(1)* @glob to %struct.MyType addrspace(4)*))
+ // CHECK: call spir_func noundef i32 @_ZNU3AS46MyType3barEv(ptr addrspace(4) {{[^,]*}} addrspacecast (ptr addrspace(1) @glob to ptr addrspace(4)))
glob.bar();
- // CHECK: call spir_func noundef i32 @_ZNU3AS26MyType3barEv(%struct.MyType addrspace(2)* {{[^,]*}} @const1)
+ // CHECK: call spir_func noundef i32 @_ZNU3AS26MyType3barEv(ptr addrspace(2) {{[^,]*}} @const1)
const1.bar();
- // CHECK: call spir_func void @_ZNU3AS26MyTypeD1Ev(%struct.MyType addrspace(2)* {{[^,]*}} @const1)
+ // CHECK: call spir_func void @_ZNU3AS26MyTypeD1Ev(ptr addrspace(2) {{[^,]*}} @const1)
const1.~MyType();
}
// CHECK-LABEL: define{{.*}} spir_kernel void @fooLocal()
kernel void fooLocal() {
// CHECK: [[VAR:%.*]] = alloca %struct.MyType
- // CHECK: [[REG:%.*]] ={{.*}} addrspacecast %struct.MyType* [[VAR]] to %struct.MyType addrspace(4)*
- // CHECK: call spir_func void @_ZNU3AS46MyTypeC1Ei(%struct.MyType addrspace(4)* {{[^,]*}} [[REG]], i32 noundef 3)
+ // CHECK: [[REG:%.*]] ={{.*}} addrspacecast ptr [[VAR]] to ptr addrspace(4)
+ // CHECK: call spir_func void @_ZNU3AS46MyTypeC1Ei(ptr addrspace(4) {{[^,]*}} [[REG]], i32 noundef 3)
MyType myLocal(3);
- // CHECK: [[REG:%.*]] ={{.*}} addrspacecast %struct.MyType* [[VAR]] to %struct.MyType addrspace(4)*
- // CHECK: call spir_func noundef i32 @_ZNU3AS46MyType3barEv(%struct.MyType addrspace(4)* {{[^,]*}} [[REG]])
+ // CHECK: [[REG:%.*]] ={{.*}} addrspacecast ptr [[VAR]] to ptr addrspace(4)
+ // CHECK: call spir_func noundef i32 @_ZNU3AS46MyType3barEv(ptr addrspace(4) {{[^,]*}} [[REG]])
myLocal.bar();
- // CHECK: [[REG:%.*]] ={{.*}} addrspacecast %struct.MyType* [[VAR]] to %struct.MyType addrspace(4)*
- // CHECK: call spir_func void @_ZNU3AS46MyTypeD1Ev(%struct.MyType addrspace(4)* {{[^,]*}} [[REG]])
+ // CHECK: [[REG:%.*]] ={{.*}} addrspacecast ptr [[VAR]] to ptr addrspace(4)
+ // CHECK: call spir_func void @_ZNU3AS46MyTypeD1Ev(ptr addrspace(4) {{[^,]*}} [[REG]])
}
// Ensure all members are defined for all the required address spaces.
-// CHECK-DEFINITIONS-DAG: define linkonce_odr spir_func void @_ZNU3AS26MyTypeC1Ei(%struct.MyType addrspace(2)* {{[^,]*}} %this, i32 noundef %i)
-// CHECK-DEFINITIONS-DAG: define linkonce_odr spir_func void @_ZNU3AS46MyTypeC1Ei(%struct.MyType addrspace(4)* {{[^,]*}} %this, i32 noundef %i)
-// CHECK-DEFINITIONS-DAG: define linkonce_odr spir_func void @_ZNU3AS26MyTypeD1Ev(%struct.MyType addrspace(2)* {{[^,]*}} %this)
-// CHECK-DEFINITIONS-DAG: define linkonce_odr spir_func void @_ZNU3AS46MyTypeD1Ev(%struct.MyType addrspace(4)* {{[^,]*}} %this)
-// CHECK-DEFINITIONS-DAG: define linkonce_odr spir_func noundef i32 @_ZNU3AS26MyType3barEv(%struct.MyType addrspace(2)* {{[^,]*}} %this)
-// CHECK-DEFINITIONS-DAG: define linkonce_odr spir_func noundef i32 @_ZNU3AS46MyType3barEv(%struct.MyType addrspace(4)* {{[^,]*}} %this)
+// CHECK-DEFINITIONS-DAG: define linkonce_odr spir_func void @_ZNU3AS26MyTypeC1Ei(ptr addrspace(2) {{[^,]*}} %this, i32 noundef %i)
+// CHECK-DEFINITIONS-DAG: define linkonce_odr spir_func void @_ZNU3AS46MyTypeC1Ei(ptr addrspace(4) {{[^,]*}} %this, i32 noundef %i)
+// CHECK-DEFINITIONS-DAG: define linkonce_odr spir_func void @_ZNU3AS26MyTypeD1Ev(ptr addrspace(2) {{[^,]*}} %this)
+// CHECK-DEFINITIONS-DAG: define linkonce_odr spir_func void @_ZNU3AS46MyTypeD1Ev(ptr addrspace(4) {{[^,]*}} %this)
+// CHECK-DEFINITIONS-DAG: define linkonce_odr spir_func noundef i32 @_ZNU3AS26MyType3barEv(ptr addrspace(2) {{[^,]*}} %this)
+// CHECK-DEFINITIONS-DAG: define linkonce_odr spir_func noundef i32 @_ZNU3AS46MyType3barEv(ptr addrspace(4) {{[^,]*}} %this)
diff --git a/clang/test/CodeGenOpenCLCXX/addrspace_cast.clcpp b/clang/test/CodeGenOpenCLCXX/addrspace_cast.clcpp
index 2bba6edb79aed..44e4e0f356f9b 100644
--- a/clang/test/CodeGenOpenCLCXX/addrspace_cast.clcpp
+++ b/clang/test/CodeGenOpenCLCXX/addrspace_cast.clcpp
@@ -1,7 +1,7 @@
-//RUN: %clang_cc1 -no-opaque-pointers %s -triple spir -emit-llvm -O0 -o - | FileCheck %s
+//RUN: %clang_cc1 %s -triple spir -emit-llvm -O0 -o - | FileCheck %s
//CHECK-LABEL: define{{.*}} spir_func void @_Z3barPU3AS1i
void bar(global int *gl) {
- //CHECK: addrspacecast i32 addrspace(1)* %{{[0-9]+}} to i32 addrspace(4)*
+ //CHECK: addrspacecast ptr addrspace(1) %{{[0-9]+}} to ptr addrspace(4)
int *gen = addrspace_cast<int *>(gl);
}
diff --git a/clang/test/CodeGenOpenCLCXX/atexit.clcpp b/clang/test/CodeGenOpenCLCXX/atexit.clcpp
index 19209ac706f29..e33296a12ef2a 100644
--- a/clang/test/CodeGenOpenCLCXX/atexit.clcpp
+++ b/clang/test/CodeGenOpenCLCXX/atexit.clcpp
@@ -1,4 +1,4 @@
-//RUN: %clang_cc1 -no-opaque-pointers %s -triple spir -emit-llvm -O0 -o - | FileCheck %s
+//RUN: %clang_cc1 %s -triple spir -emit-llvm -O0 -o - | FileCheck %s
struct S {
~S(){};
@@ -6,6 +6,6 @@ struct S {
S s;
//CHECK-LABEL: define internal spir_func void @__cxx_global_var_init()
-//CHECK: call spir_func i32 @__cxa_atexit(void (i8*)* bitcast (void (%struct.S addrspace(4)*)* @_ZNU3AS41SD1Ev to void (i8*)*), i8* null, i8 addrspace(1)* @__dso_handle)
+//CHECK: call spir_func i32 @__cxa_atexit(ptr @_ZNU3AS41SD1Ev, ptr null, ptr addrspace(1) @__dso_handle)
-//CHECK: declare spir_func i32 @__cxa_atexit(void (i8*)*, i8*, i8 addrspace(1)*)
+//CHECK: declare spir_func i32 @__cxa_atexit(ptr, ptr, ptr addrspace(1))
diff --git a/clang/test/CodeGenOpenCLCXX/constexpr.clcpp b/clang/test/CodeGenOpenCLCXX/constexpr.clcpp
index 572604b98d040..0576418c944ba 100644
--- a/clang/test/CodeGenOpenCLCXX/constexpr.clcpp
+++ b/clang/test/CodeGenOpenCLCXX/constexpr.clcpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -no-opaque-pointers %s -triple spir-unknown-unknown -O0 -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 %s -triple spir-unknown-unknown -O0 -emit-llvm -o - | FileCheck %s
typedef int int2 __attribute__((ext_vector_type(2)));
typedef int int4 __attribute__((ext_vector_type(4)));
@@ -31,7 +31,7 @@ kernel void foo(global float *x) {
// Test evaluation of constant vectors.
// CHECK-LABEL: define{{.*}} spir_kernel void @vecEval
// CHECK: store i32 3
-// CHECK: store <2 x i32> <i32 22, i32 33>, <2 x i32>
+// CHECK: store <2 x i32> <i32 22, i32 33>, ptr
const int oneElt = int4(3).x;
const int2 twoElts = (int4)(11, 22, 33, 44).yz;
diff --git a/clang/test/CodeGenOpenCLCXX/method-overload-address-space.clcpp b/clang/test/CodeGenOpenCLCXX/method-overload-address-space.clcpp
index b165e5f65d30a..b10e2d790b75f 100644
--- a/clang/test/CodeGenOpenCLCXX/method-overload-address-space.clcpp
+++ b/clang/test/CodeGenOpenCLCXX/method-overload-address-space.clcpp
@@ -1,4 +1,4 @@
-//RUN: %clang_cc1 -no-opaque-pointers %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
struct C {
void foo() __local;
@@ -15,21 +15,21 @@ __kernel void k() {
__global C &c_ref = c1;
__global C *c_ptr;
- // CHECK: call spir_func void @_ZNU3AS11C3fooEv(%struct.C addrspace(1)*
+ // CHECK: call spir_func void @_ZNU3AS11C3fooEv(ptr addrspace(1)
c1.foo();
- // CHECK: call spir_func void @_ZNU3AS31C3fooEv(%struct.C addrspace(3)*
+ // CHECK: call spir_func void @_ZNU3AS31C3fooEv(ptr addrspace(3)
c2.foo();
- // CHECK: call spir_func void @_ZNU3AS41C3fooEv(%struct.C addrspace(4)*
+ // CHECK: call spir_func void @_ZNU3AS41C3fooEv(ptr addrspace(4)
c3.foo();
- // CHECK: call spir_func void @_ZNU3AS11C3fooEv(%struct.C addrspace(1)*
+ // CHECK: call spir_func void @_ZNU3AS11C3fooEv(ptr addrspace(1)
c_ptr->foo();
- // CHECK: spir_func void @_ZNU3AS11C3fooEv(%struct.C addrspace(1)*
+ // CHECK: spir_func void @_ZNU3AS11C3fooEv(ptr addrspace(1)
c_ref.foo();
- // CHECK: call spir_func void @_ZNU3AS41C3barEv(%struct.C addrspace(4)* {{[^,]*}} addrspacecast (%struct.C addrspace(1)* @c1 to %struct.C addrspace(4)*))
+ // CHECK: call spir_func void @_ZNU3AS41C3barEv(ptr addrspace(4) {{[^,]*}} addrspacecast (ptr addrspace(1) @c1 to ptr addrspace(4)))
c1.bar();
//FIXME: Doesn't compile yet
//c_ptr->bar();
- // CHECK: call spir_func void @_ZNU3AS41C3barEv(%struct.C addrspace(4)* {{[^,]*}} addrspacecast (%struct.C addrspace(1)* @c1 to %struct.C addrspace(4)*))
+ // CHECK: call spir_func void @_ZNU3AS41C3barEv(ptr addrspace(4) {{[^,]*}} addrspacecast (ptr addrspace(1) @c1 to ptr addrspace(4)))
c_ref.bar();
}
diff --git a/clang/test/CodeGenOpenCLCXX/reinterpret_cast.clcpp b/clang/test/CodeGenOpenCLCXX/reinterpret_cast.clcpp
index 98e2613608dd5..c8a58fea2bbd1 100644
--- a/clang/test/CodeGenOpenCLCXX/reinterpret_cast.clcpp
+++ b/clang/test/CodeGenOpenCLCXX/reinterpret_cast.clcpp
@@ -1,4 +1,4 @@
-//RUN: %clang_cc1 -no-opaque-pointers %s -triple spir -emit-llvm -O0 -o - | FileCheck %s
+//RUN: %clang_cc1 %s -triple spir -emit-llvm -O0 -o - | FileCheck %s
typedef int int2 __attribute__((ext_vector_type(2)));
typedef int int4 __attribute__((ext_vector_type(4)));
@@ -12,14 +12,14 @@ void bar(global int2 *in) {
auto i2 = reinterpret_cast<int2>(l);
__private short s1;
- // CHECK: %{{[0-9]+}} = load i16, i16* %s1, align 2
- // CHECK-NEXT: store i16 %{{[0-9]+}}, i16* %s2, align 2
+ // CHECK: %{{[0-9]+}} = load i16, ptr %s1, align 2
+ // CHECK-NEXT: store i16 %{{[0-9]+}}, ptr %s2, align 2
auto s2 = reinterpret_cast<__private short>(s1);
- // CHECK: %{{[0-9]+}} = load i16, i16* %s1, align 2
- // CHECK-NEXT: store i16 %{{[0-9]+}}, i16* %s3, align 2
+ // CHECK: %{{[0-9]+}} = load i16, ptr %s1, align 2
+ // CHECK-NEXT: store i16 %{{[0-9]+}}, ptr %s3, align 2
auto s3 = reinterpret_cast<decltype(s1)>(s1);
- // CHECK: %{{[0-9]+}} = load i16, i16* %s1, align 2
- // CHECK-NEXT: store i16 %{{[0-9]+}}, i16* %s4, align 2
+ // CHECK: %{{[0-9]+}} = load i16, ptr %s1, align 2
+ // CHECK-NEXT: store i16 %{{[0-9]+}}, ptr %s4, align 2
auto s4 = reinterpret_cast<__global short>(s1);
int4 i4;
diff --git a/clang/test/CodeGenOpenCLCXX/template-address-spaces.clcpp b/clang/test/CodeGenOpenCLCXX/template-address-spaces.clcpp
index 052fce45c3ae4..b6a177a176a8f 100644
--- a/clang/test/CodeGenOpenCLCXX/template-address-spaces.clcpp
+++ b/clang/test/CodeGenOpenCLCXX/template-address-spaces.clcpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -no-opaque-pointers %s -emit-llvm -o - -O0 -triple spir-unknown-unknown | FileCheck %s
+// RUN: %clang_cc1 %s -emit-llvm -o - -O0 -triple spir-unknown-unknown | FileCheck %s
template <typename T>
struct S{
@@ -10,15 +10,15 @@ template<typename T>
T S<T>::foo() { return a;}
// CHECK: %struct.S = type { i32 }
-// CHECK: %struct.S.0 = type { i32 addrspace(4)* }
-// CHECK: %struct.S.1 = type { i32 addrspace(1)* }
+// CHECK: %struct.S.0 = type { ptr addrspace(4) }
+// CHECK: %struct.S.1 = type { ptr addrspace(1) }
-// CHECK: [[A1:%[.a-z0-9]+]] = addrspacecast %struct.S* %{{[a-z0-9]+}} to %struct.S addrspace(4)*
-// CHECK: %call = call spir_func noundef i32 @_ZNU3AS41SIiE3fooEv(%struct.S addrspace(4)* {{[^,]*}} [[A1]]) #1
-// CHECK: [[A2:%[.a-z0-9]+]] = addrspacecast %struct.S.0* %{{[a-z0-9]+}} to %struct.S.0 addrspace(4)*
-// CHECK: %call1 = call spir_func noundef i32 addrspace(4)* @_ZNU3AS41SIPU3AS4iE3fooEv(%struct.S.0 addrspace(4)* {{[^,]*}} [[A2]]) #1
-// CHECK: [[A3:%[.a-z0-9]+]] = addrspacecast %struct.S.1* %{{[a-z0-9]+}} to %struct.S.1 addrspace(4)*
-// CHECK: %call2 = call spir_func noundef i32 addrspace(1)* @_ZNU3AS41SIPU3AS1iE3fooEv(%struct.S.1 addrspace(4)* {{[^,]*}} [[A3]]) #1
+// CHECK: [[A1:%[.a-z0-9]+]] = addrspacecast ptr %{{[a-z0-9]+}} to ptr addrspace(4)
+// CHECK: %call = call spir_func noundef i32 @_ZNU3AS41SIiE3fooEv(ptr addrspace(4) {{[^,]*}} [[A1]]) #1
+// CHECK: [[A2:%[.a-z0-9]+]] = addrspacecast ptr %{{[a-z0-9]+}} to ptr addrspace(4)
+// CHECK: %call1 = call spir_func noundef ptr addrspace(4) @_ZNU3AS41SIPU3AS4iE3fooEv(ptr addrspace(4) {{[^,]*}} [[A2]]) #1
+// CHECK: [[A3:%[.a-z0-9]+]] = addrspacecast ptr %{{[a-z0-9]+}} to ptr addrspace(4)
+// CHECK: %call2 = call spir_func noundef ptr addrspace(1) @_ZNU3AS41SIPU3AS1iE3fooEv(ptr addrspace(4) {{[^,]*}} [[A3]]) #1
void bar(){
S<int> sint;
More information about the cfe-commits
mailing list