Skip to content

Commit 39c82eb

Browse files
Fznamznonbader
Fznamznon
authored and
bader
committed
[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>
1 parent a5f2fda commit 39c82eb

File tree

9 files changed

+39
-38
lines changed

9 files changed

+39
-38
lines changed

‎clang/test/CodeGenSYCL/Inputs/sycl.hpp

+6-6
Original file line numberDiff line numberDiff line change
@@ -8,18 +8,18 @@
88

99
// Dummy runtime classes to model SYCL API.
1010
namespace cl {
11-
namespace __spirv {
12-
class OpTypeSampler;
13-
}
14-
1511
namespace sycl {
1612
struct sampler_impl {
17-
__spirv::OpTypeSampler* m_Sampler;
13+
#ifdef __SYCL_DEVICE_ONLY__
14+
__ocl_sampler_t m_Sampler;
15+
#endif
1816
};
1917

2018
class sampler {
2119
struct sampler_impl impl;
22-
void __init(__spirv::OpTypeSampler* Sampler) { impl.m_Sampler = Sampler; }
20+
#ifdef __SYCL_DEVICE_ONLY__
21+
void __init(__ocl_sampler_t Sampler) { impl.m_Sampler = Sampler; }
22+
#endif
2323

2424
public:
2525
void use(void) const {}

‎clang/test/CodeGenSYCL/device-functions.cpp

+1-1
Original file line numberDiff line numberDiff line change
@@ -22,6 +22,6 @@ int main() {
2222
return 0;
2323
}
2424
// CHECK: define spir_kernel void @_ZTSZ4mainE11fake_kernel()
25-
// CHECK: define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%class.anon* %this)
25+
// CHECK: define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}.anon"* %this)
2626
// CHECK: define spir_func void @_Z3foov()
2727
// CHECK: define linkonce_odr spir_func i32 @_Z3barIiET_S0_(i32 %arg)

‎clang/test/CodeGenSYCL/intel-fpga-local.cpp

+7-7
Original file line numberDiff line numberDiff line change
@@ -43,31 +43,31 @@ struct foo_two {
4343

4444
void bar() {
4545
struct foo_two s1;
46-
//CHECK: %[[FIELD1:.*]] = getelementptr inbounds %struct.foo_two{{.*}}
46+
//CHECK: %[[FIELD1:.*]] = getelementptr inbounds %struct.{{.*}}.foo_two{{.*}}
4747
//CHECK: %[[CAST:.*]] = bitcast{{.*}}%[[FIELD1]]
4848
//CHECK: call i8* @llvm.ptr.annotation.p0i8{{.*}}%[[CAST]]{{.*}}[[ANN1]]
4949
s1.f1 = 0;
50-
//CHECK: %[[FIELD2:.*]] = getelementptr inbounds %struct.foo_two{{.*}}
50+
//CHECK: %[[FIELD2:.*]] = getelementptr inbounds %struct.{{.*}}.foo_two{{.*}}
5151
//CHECK: %[[CAST:.*]] = bitcast{{.*}}%[[FIELD2]]
5252
//CHECK: call i8* @llvm.ptr.annotation.p0i8{{.*}}%[[CAST]]{{.*}}[[ANN2]]
5353
s1.f2 = 0;
54-
//CHECK: %[[FIELD3:.*]] = getelementptr inbounds %struct.foo_two{{.*}}
54+
//CHECK: %[[FIELD3:.*]] = getelementptr inbounds %struct.{{.*}}.foo_two{{.*}}
5555
//CHECK: %[[CAST:.*]] = bitcast{{.*}}%[[FIELD3]]
5656
//CHECK: call i8* @llvm.ptr.annotation.p0i8{{.*}}%[[CAST]]{{.*}}[[ANN3]]
5757
s1.f3 = 0;
58-
//CHECK: %[[FIELD4:.*]] = getelementptr inbounds %struct.foo_two{{.*}}
58+
//CHECK: %[[FIELD4:.*]] = getelementptr inbounds %struct.{{.*}}.foo_two{{.*}}
5959
//CHECK: %[[CAST:.*]] = bitcast{{.*}}%[[FIELD4]]
6060
//CHECK: call i8* @llvm.ptr.annotation.p0i8{{.*}}%[[CAST]]{{.*}}[[ANN4]]
6161
s1.f4 = 0;
62-
//CHECK: %[[FIELD5:.*]] = getelementptr inbounds %struct.foo_two{{.*}}
62+
//CHECK: %[[FIELD5:.*]] = getelementptr inbounds %struct.{{.*}}.foo_two{{.*}}
6363
//CHECK: %[[CAST:.*]] = bitcast{{.*}}%[[FIELD5]]
6464
//CHECK: call i8* @llvm.ptr.annotation.p0i8{{.*}}%[[CAST]]{{.*}}[[ANN5]]
6565
s1.f5 = 0;
66-
//CHECK: %[[FIELD6:.*]] = getelementptr inbounds %struct.foo_two{{.*}}
66+
//CHECK: %[[FIELD6:.*]] = getelementptr inbounds %struct.{{.*}}.foo_two{{.*}}
6767
//CHECK: %[[CAST:.*]] = bitcast{{.*}}%[[FIELD6]]
6868
//CHECK: call i8* @llvm.ptr.annotation.p0i8{{.*}}%[[CAST]]{{.*}}[[ANN10]]
6969
s1.f6 = 0;
70-
//CHECK: %[[FIELD7:.*]] = getelementptr inbounds %struct.foo_two{{.*}}
70+
//CHECK: %[[FIELD7:.*]] = getelementptr inbounds %struct.{{.*}}.foo_two{{.*}}
7171
//CHECK: %[[CAST:.*]] = bitcast{{.*}}%[[FIELD7]]
7272
//CHECK: call i8* @llvm.ptr.annotation.p0i8{{.*}}%[[CAST]]{{.*}}[[ANN11]]
7373
s1.f7 = 0;

‎clang/test/CodeGenSYCL/kernel-with-id.cpp

+1-1
Original file line numberDiff line numberDiff line change
@@ -9,7 +9,7 @@ __attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
99

1010
int main() {
1111
cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write> accessorA;
12-
// 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 %{{.*}})
12+
// 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 %{{.*}})
1313
kernel<class kernel_function>(
1414
[=]() {
1515
accessorA.use();

‎clang/test/CodeGenSYCL/sampler.cpp

+8-8
Original file line numberDiff line numberDiff line change
@@ -1,14 +1,14 @@
11
// 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
2-
// CHECK: define spir_kernel void @{{[a-zA-Z0-9_]+}}(%spirv.Sampler* [[SAMPLER_ARG:%[a-zA-Z0-9_]+]])
2+
// CHECK: define spir_kernel void @{{[a-zA-Z0-9_]+}}(%opencl.sampler_t addrspace(2)* [[SAMPLER_ARG:%[a-zA-Z0-9_]+]])
33
// CHECK-NEXT: entry:
4-
// CHECK-NEXT: [[SAMPLER_ARG]].addr = alloca %spirv.Sampler*, align 8
5-
// CHECK-NEXT: [[ANON:%[0-9]+]] = alloca %class.anon, align 8
6-
// CHECK-NEXT: store %spirv.Sampler* [[SAMPLER_ARG]], %spirv.Sampler** [[SAMPLER_ARG]].addr, align 8, !tbaa !9
7-
// CHECK-NEXT: [[BITCAST:%[0-9]+]] = bitcast %class.anon* [[ANON]] to i8*
4+
// CHECK-NEXT: [[SAMPLER_ARG]].addr = alloca %opencl.sampler_t addrspace(2)*, align 8
5+
// CHECK-NEXT: [[ANON:%[0-9]+]] = alloca %"class.{{.*}}.anon", align 8
6+
// CHECK-NEXT: store %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG]], %opencl.sampler_t addrspace(2)** [[SAMPLER_ARG]].addr, align 8, !tbaa !9
7+
// CHECK-NEXT: [[BITCAST:%[0-9]+]] = bitcast %"class.{{.*}}.anon"* [[ANON]] to i8*
88
// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* [[BITCAST]]) #4
9-
// CHECK-NEXT: [[GEP:%[0-9]+]] = getelementptr inbounds %class.anon, %class.anon* [[ANON]], i32 0, i32 0
10-
// CHECK-NEXT: [[LOAD_SAMPLER_ARG:%[0-9]+]] = load %spirv.Sampler*, %spirv.Sampler** [[SAMPLER_ARG]].addr, align 8, !tbaa !9
11-
// CHECK-NEXT: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.cl::sycl::sampler"* [[GEP]], %spirv.Sampler* [[LOAD_SAMPLER_ARG]])
9+
// CHECK-NEXT: [[GEP:%[0-9]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[ANON]], i32 0, i32 0
10+
// 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
11+
// CHECK-NEXT: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.{{.*}}.cl::sycl::sampler"* [[GEP]], %opencl.sampler_t addrspace(2)* [[LOAD_SAMPLER_ARG]])
1212
//
1313
#include "sycl.hpp"
1414

‎clang/test/CodeGenSYCL/spir-calling-conv.cpp

+2-2
Original file line numberDiff line numberDiff line change
@@ -9,9 +9,9 @@ int main() {
99

1010
// CHECK: define spir_kernel void @_ZTSZ4mainE15kernel_function()
1111

12-
// CHECK: call spir_func void @"_ZZ4mainENK3$_0clEv"(%class.anon* %0)
12+
// CHECK: call spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}.anon"* %0)
1313

14-
// CHECK: define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%class.anon* %this)
14+
// CHECK: define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}anon"* %this)
1515

1616
kernel_single_task<class kernel_function>([]() {});
1717
return 0;

‎clang/test/SemaSYCL/Inputs/sycl.hpp

+6-5
Original file line numberDiff line numberDiff line change
@@ -8,9 +8,6 @@
88
#endif
99

1010
namespace cl {
11-
namespace __spirv {
12-
class OpTypeSampler;
13-
}
1411
namespace sycl {
1512
namespace access {
1613

@@ -75,12 +72,16 @@ class accessor {
7572
};
7673

7774
struct sampler_impl {
78-
__spirv::OpTypeSampler *m_Sampler;
75+
#ifdef __SYCL_DEVICE_ONLY__
76+
__ocl_sampler_t m_Sampler;
77+
#endif
7978
};
8079

8180
class sampler {
8281
struct sampler_impl impl;
83-
void __init(__spirv::OpTypeSampler *Sampler) { impl.m_Sampler = Sampler; }
82+
#ifdef __SYCL_DEVICE_ONLY__
83+
void __init(__ocl_sampler_t Sampler) { impl.m_Sampler = Sampler; }
84+
#endif
8485

8586
public:
8687
void use(void) const {}

‎clang/test/SemaSYCL/sampler.cpp

+6-6
Original file line numberDiff line numberDiff line change
@@ -16,18 +16,18 @@ int main() {
1616
}
1717

1818
// Check declaration of the test kernel
19-
// CHECK: FunctionDecl {{.*}}use_kernel_for_test 'void (__spirv::OpTypeSampler *)'
19+
// CHECK: FunctionDecl {{.*}}use_kernel_for_test 'void (__ocl_sampler_t)'
2020
//
2121
// Check parameters of the test kernel
22-
// CHECK: ParmVarDecl {{.*}} used [[_arg_sampler:[0-9a-zA-Z_]+]] '__spirv::OpTypeSampler *'
22+
// CHECK: ParmVarDecl {{.*}} used [[_arg_sampler:[0-9a-zA-Z_]+]] '__ocl_sampler_t'
2323
//
2424
// Check that sampler field of the test kernel object is initialized using __init method
2525
// CHECK: CXXMemberCallExpr {{.*}} 'void'
26-
// CHECK-NEXT: MemberExpr {{.*}} 'void (__spirv::OpTypeSampler *)' lvalue .__init
26+
// CHECK-NEXT: MemberExpr {{.*}} 'void (__ocl_sampler_t)' lvalue .__init
2727
// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::sampler':'cl::sycl::sampler' lvalue
2828
// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}sampler.cpp{{.*}})' lvalue Var {{.*}} '(lambda at {{.*}}sampler.cpp{{.*}})'
2929
//
3030
// Check the parameters of __init method
31-
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__spirv::OpTypeSampler *' <LValueToRValue>
32-
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__spirv::OpTypeSampler *' lvalue <NoOp>
33-
// CHECK-NEXT: DeclRefExpr {{.*}} '__spirv::OpTypeSampler *' lvalue ParmVar {{.*}} '[[_arg_sampler]]' '__spirv::OpTypeSampler
31+
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__ocl_sampler_t':'sampler_t' <LValueToRValue>
32+
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__ocl_sampler_t':'sampler_t' lvalue <NoOp>
33+
// CHECK-NEXT: DeclRefExpr {{.*}} '__ocl_sampler_t':'sampler_t' lvalue ParmVar {{.*}} '[[_arg_sampler]]' '__ocl_sampler_t':'sampler_t'

‎clang/test/SemaSYCL/spir-enum.cpp

+2-2
Original file line numberDiff line numberDiff line change
@@ -22,8 +22,8 @@ int main() {
2222

2323
// CHECK: define spir_kernel void @_ZTSZ4test9enum_typeE15kernel_function(i32 %_arg_)
2424

25-
// CHECK: getelementptr inbounds %class.anon, %class.anon*
26-
// CHECK: call spir_func void @"_ZZ4test9enum_typeENK3$_0clEv"(%class.anon* %0)
25+
// CHECK: getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"*
26+
// CHECK: call spir_func void @"_ZZ4test9enum_typeENK3$_0clEv"(%"class.{{.*}}.anon"* %0)
2727

2828

2929
test( enum_type::B );

0 commit comments

Comments
 (0)