[SYCL] Update SYCL clang tests

Replace SPIR-V type names with built-in OpenCL type names. Use mangled class names in IR checking. Signed-off-by: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>
MrSidims · May 29, 2019 · 39c82eb · 39c82eb
1 parent a5f2fda
commit 39c82eb
Show file tree

Hide file tree

Showing 9 changed files with 39 additions and 38 deletions.
diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
@@ -8,18 +8,18 @@
 
 // Dummy runtime classes to model SYCL API.
 namespace cl {
-namespace __spirv {
-class OpTypeSampler;
-}
-
 namespace sycl {
 struct sampler_impl {
-  __spirv::OpTypeSampler* m_Sampler;
+#ifdef __SYCL_DEVICE_ONLY__
+  __ocl_sampler_t m_Sampler;
+#endif
 };
 
 class sampler {
   struct sampler_impl impl;
-  void __init(__spirv::OpTypeSampler* Sampler) { impl.m_Sampler = Sampler; }
+#ifdef __SYCL_DEVICE_ONLY__
+  void __init(__ocl_sampler_t Sampler) { impl.m_Sampler = Sampler; }
+#endif
 
 public:
   void use(void) const {}

diff --git a/clang/test/CodeGenSYCL/device-functions.cpp b/clang/test/CodeGenSYCL/device-functions.cpp
@@ -22,6 +22,6 @@ int main() {
   return 0;
 }
 // CHECK: define spir_kernel void @_ZTSZ4mainE11fake_kernel()
-// CHECK: define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%class.anon* %this)
+// CHECK: define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}.anon"* %this)
 // CHECK: define spir_func void @_Z3foov()
 // CHECK: define linkonce_odr spir_func i32 @_Z3barIiET_S0_(i32 %arg)
diff --git a/clang/test/CodeGenSYCL/intel-fpga-local.cpp b/clang/test/CodeGenSYCL/intel-fpga-local.cpp
@@ -43,31 +43,31 @@ struct foo_two {
 
 void bar() {
   struct foo_two s1;
-  //CHECK: %[[FIELD1:.*]] = getelementptr inbounds %struct.foo_two{{.*}}
+  //CHECK: %[[FIELD1:.*]] = getelementptr inbounds %struct.{{.*}}.foo_two{{.*}}
   //CHECK: %[[CAST:.*]] = bitcast{{.*}}%[[FIELD1]]
   //CHECK: call i8* @llvm.ptr.annotation.p0i8{{.*}}%[[CAST]]{{.*}}[[ANN1]]
   s1.f1 = 0;
-  //CHECK: %[[FIELD2:.*]] = getelementptr inbounds %struct.foo_two{{.*}}
+  //CHECK: %[[FIELD2:.*]] = getelementptr inbounds %struct.{{.*}}.foo_two{{.*}}
   //CHECK: %[[CAST:.*]] = bitcast{{.*}}%[[FIELD2]]
   //CHECK: call i8* @llvm.ptr.annotation.p0i8{{.*}}%[[CAST]]{{.*}}[[ANN2]]
   s1.f2 = 0;
-  //CHECK: %[[FIELD3:.*]] = getelementptr inbounds %struct.foo_two{{.*}}
+  //CHECK: %[[FIELD3:.*]] = getelementptr inbounds %struct.{{.*}}.foo_two{{.*}}
   //CHECK: %[[CAST:.*]] = bitcast{{.*}}%[[FIELD3]]
   //CHECK: call i8* @llvm.ptr.annotation.p0i8{{.*}}%[[CAST]]{{.*}}[[ANN3]]
   s1.f3 = 0;
-  //CHECK: %[[FIELD4:.*]] = getelementptr inbounds %struct.foo_two{{.*}}
+  //CHECK: %[[FIELD4:.*]] = getelementptr inbounds %struct.{{.*}}.foo_two{{.*}}
   //CHECK: %[[CAST:.*]] = bitcast{{.*}}%[[FIELD4]]
   //CHECK: call i8* @llvm.ptr.annotation.p0i8{{.*}}%[[CAST]]{{.*}}[[ANN4]]
   s1.f4 = 0;
-  //CHECK: %[[FIELD5:.*]] = getelementptr inbounds %struct.foo_two{{.*}}
+  //CHECK: %[[FIELD5:.*]] = getelementptr inbounds %struct.{{.*}}.foo_two{{.*}}
   //CHECK: %[[CAST:.*]] = bitcast{{.*}}%[[FIELD5]]
   //CHECK: call i8* @llvm.ptr.annotation.p0i8{{.*}}%[[CAST]]{{.*}}[[ANN5]]
   s1.f5 = 0;
-  //CHECK: %[[FIELD6:.*]] = getelementptr inbounds %struct.foo_two{{.*}}
+  //CHECK: %[[FIELD6:.*]] = getelementptr inbounds %struct.{{.*}}.foo_two{{.*}}
   //CHECK: %[[CAST:.*]] = bitcast{{.*}}%[[FIELD6]]
   //CHECK: call i8* @llvm.ptr.annotation.p0i8{{.*}}%[[CAST]]{{.*}}[[ANN10]]
   s1.f6 = 0;
-  //CHECK: %[[FIELD7:.*]] = getelementptr inbounds %struct.foo_two{{.*}}
+  //CHECK: %[[FIELD7:.*]] = getelementptr inbounds %struct.{{.*}}.foo_two{{.*}}
   //CHECK: %[[CAST:.*]] = bitcast{{.*}}%[[FIELD7]]
   //CHECK: call i8* @llvm.ptr.annotation.p0i8{{.*}}%[[CAST]]{{.*}}[[ANN11]]
   s1.f7 = 0;

diff --git a/clang/test/CodeGenSYCL/kernel-with-id.cpp b/clang/test/CodeGenSYCL/kernel-with-id.cpp
@@ -9,7 +9,7 @@ __attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
 
 int main() {
   cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write> accessorA;
-// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.cl::sycl::accessor"* %{{.*}}, i32 addrspace(1)* %{{.*}}, %"struct.cl::sycl::range"* byval align 1 %{{.*}}, %"struct.cl::sycl::id"* byval align 1 %{{.*}})
+// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor"* %{{.*}}, i32 addrspace(1)* %{{.*}}, %"struct.{{.*}}.cl::sycl::range"* byval align 1 %{{.*}}, %"struct.{{.*}}.cl::sycl::id"* byval align 1 %{{.*}})
     kernel<class kernel_function>(
       [=]() {
         accessorA.use();

diff --git a/clang/test/CodeGenSYCL/sampler.cpp b/clang/test/CodeGenSYCL/sampler.cpp
@@ -1,14 +1,14 @@
 // RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -I %S/Inputs  -fsycl-is-device -disable-llvm-passes -emit-llvm -x c++ %s -o - | FileCheck --enable-var-scope %s
-// CHECK: define spir_kernel void @{{[a-zA-Z0-9_]+}}(%spirv.Sampler* [[SAMPLER_ARG:%[a-zA-Z0-9_]+]])
+// CHECK: define spir_kernel void @{{[a-zA-Z0-9_]+}}(%opencl.sampler_t addrspace(2)* [[SAMPLER_ARG:%[a-zA-Z0-9_]+]])
 // CHECK-NEXT: entry:
-// CHECK-NEXT: [[SAMPLER_ARG]].addr = alloca %spirv.Sampler*, align 8
-// CHECK-NEXT: [[ANON:%[0-9]+]] = alloca %class.anon, align 8
-// CHECK-NEXT: store %spirv.Sampler* [[SAMPLER_ARG]], %spirv.Sampler** [[SAMPLER_ARG]].addr, align 8, !tbaa !9
-// CHECK-NEXT: [[BITCAST:%[0-9]+]] = bitcast %class.anon* [[ANON]] to i8*
+// CHECK-NEXT: [[SAMPLER_ARG]].addr = alloca %opencl.sampler_t addrspace(2)*, align 8
+// CHECK-NEXT: [[ANON:%[0-9]+]] = alloca %"class.{{.*}}.anon", align 8
+// CHECK-NEXT: store %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG]], %opencl.sampler_t addrspace(2)** [[SAMPLER_ARG]].addr, align 8, !tbaa !9
+// CHECK-NEXT: [[BITCAST:%[0-9]+]] = bitcast %"class.{{.*}}.anon"* [[ANON]] to i8*
 // CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* [[BITCAST]]) #4
-// CHECK-NEXT: [[GEP:%[0-9]+]]  = getelementptr inbounds %class.anon, %class.anon* [[ANON]], i32 0, i32 0
-// CHECK-NEXT: [[LOAD_SAMPLER_ARG:%[0-9]+]] = load %spirv.Sampler*, %spirv.Sampler** [[SAMPLER_ARG]].addr, align 8, !tbaa !9
-// CHECK-NEXT: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.cl::sycl::sampler"* [[GEP]], %spirv.Sampler* [[LOAD_SAMPLER_ARG]])
+// CHECK-NEXT: [[GEP:%[0-9]+]]  = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[ANON]], i32 0, i32 0
+// CHECK-NEXT: [[LOAD_SAMPLER_ARG:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[SAMPLER_ARG]].addr, align 8, !tbaa !9
+// CHECK-NEXT: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.{{.*}}.cl::sycl::sampler"* [[GEP]], %opencl.sampler_t addrspace(2)* [[LOAD_SAMPLER_ARG]])
 //
 #include "sycl.hpp"
 

diff --git a/clang/test/CodeGenSYCL/spir-calling-conv.cpp b/clang/test/CodeGenSYCL/spir-calling-conv.cpp
@@ -9,9 +9,9 @@ int main() {
 
   // CHECK: define spir_kernel void @_ZTSZ4mainE15kernel_function()
 
-  // CHECK: call spir_func void @"_ZZ4mainENK3$_0clEv"(%class.anon* %0)
+  // CHECK: call spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}.anon"* %0)
 
-  // CHECK: define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%class.anon* %this)
+  // CHECK: define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}anon"* %this)
 
   kernel_single_task<class kernel_function>([]() {});
   return 0;

diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp
@@ -8,9 +8,6 @@
 #endif
 
 namespace cl {
-namespace __spirv {
-class OpTypeSampler;
-}
 namespace sycl {
 namespace access {
 
@@ -75,12 +72,16 @@ class accessor {
 };
 
 struct sampler_impl {
-  __spirv::OpTypeSampler *m_Sampler;
+#ifdef __SYCL_DEVICE_ONLY__
+  __ocl_sampler_t m_Sampler;
+#endif
 };
 
 class sampler {
   struct sampler_impl impl;
-  void __init(__spirv::OpTypeSampler *Sampler) { impl.m_Sampler = Sampler; }
+#ifdef __SYCL_DEVICE_ONLY__
+  void __init(__ocl_sampler_t Sampler) { impl.m_Sampler = Sampler; }
+#endif
 
 public:
   void use(void) const {}

diff --git a/clang/test/SemaSYCL/sampler.cpp b/clang/test/SemaSYCL/sampler.cpp
@@ -16,18 +16,18 @@ int main() {
 }
 
 // Check declaration of the test kernel
-// CHECK: FunctionDecl {{.*}}use_kernel_for_test 'void (__spirv::OpTypeSampler *)'
+// CHECK: FunctionDecl {{.*}}use_kernel_for_test 'void (__ocl_sampler_t)'
 //
 // Check parameters of the test kernel
-// CHECK: ParmVarDecl {{.*}} used [[_arg_sampler:[0-9a-zA-Z_]+]] '__spirv::OpTypeSampler *'
+// CHECK: ParmVarDecl {{.*}} used [[_arg_sampler:[0-9a-zA-Z_]+]] '__ocl_sampler_t'
 //
 // Check that sampler field of the test kernel object is initialized using __init method
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (__spirv::OpTypeSampler *)' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}} 'void (__ocl_sampler_t)' lvalue .__init
 // CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::sampler':'cl::sycl::sampler' lvalue
 // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}sampler.cpp{{.*}})' lvalue Var {{.*}} '(lambda at {{.*}}sampler.cpp{{.*}})'
 //
 // Check the parameters of __init method
-// CHECK-NEXT: ImplicitCastExpr {{.*}} '__spirv::OpTypeSampler *' <LValueToRValue>
-// CHECK-NEXT: ImplicitCastExpr {{.*}} '__spirv::OpTypeSampler *' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} '__spirv::OpTypeSampler *' lvalue ParmVar {{.*}} '[[_arg_sampler]]' '__spirv::OpTypeSampler
+// CHECK-NEXT: ImplicitCastExpr {{.*}} '__ocl_sampler_t':'sampler_t' <LValueToRValue>
+// CHECK-NEXT: ImplicitCastExpr {{.*}} '__ocl_sampler_t':'sampler_t' lvalue <NoOp>
+// CHECK-NEXT: DeclRefExpr {{.*}} '__ocl_sampler_t':'sampler_t' lvalue ParmVar {{.*}} '[[_arg_sampler]]' '__ocl_sampler_t':'sampler_t'
diff --git a/clang/test/SemaSYCL/spir-enum.cpp b/clang/test/SemaSYCL/spir-enum.cpp
@@ -22,8 +22,8 @@ int main() {
 
   // CHECK: define spir_kernel void @_ZTSZ4test9enum_typeE15kernel_function(i32 %_arg_)
 
-  // CHECK: getelementptr inbounds %class.anon, %class.anon*
-  // CHECK: call spir_func void @"_ZZ4test9enum_typeENK3$_0clEv"(%class.anon* %0)
+  // CHECK: getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"*
+  // CHECK: call spir_func void @"_ZZ4test9enum_typeENK3$_0clEv"(%"class.{{.*}}.anon"* %0)
 
 
   test( enum_type::B );