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; _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
