Skip to content

Commit 6a6ec9f

Browse files
vmaksimovmaksimo
committed
Fix SYCL lit-tests after 1b1c8d8 by adding missing noundef attribute
1 parent bd98ada commit 6a6ec9f

File tree

6 files changed

+21
-21
lines changed

6 files changed

+21
-21
lines changed

‎sycl/test/check_device_code/atomic_fence.cpp

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -6,19 +6,19 @@ int main() {
66
sycl::queue Q;
77

88
Q.single_task([] {
9-
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 2, i32 896) #{{.*}}
9+
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 noundef 2, i32 noundef 896) #{{.*}}
1010
sycl::atomic_fence(sycl::memory_order::relaxed,
1111
sycl::memory_scope::work_group);
12-
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 2, i32 898) #{{.*}}
12+
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 noundef 2, i32 noundef 898) #{{.*}}
1313
sycl::atomic_fence(sycl::memory_order::acquire,
1414
sycl::memory_scope::work_group);
15-
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 2, i32 900) #{{.*}}
15+
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 noundef 2, i32 noundef 900) #{{.*}}
1616
sycl::atomic_fence(sycl::memory_order::release,
1717
sycl::memory_scope::work_group);
18-
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 2, i32 904) #{{.*}}
18+
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 noundef 2, i32 noundef 904) #{{.*}}
1919
sycl::atomic_fence(sycl::memory_order::acq_rel,
2020
sycl::memory_scope::work_group);
21-
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 2, i32 912) #{{.*}}
21+
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 noundef 2, i32 noundef 912) #{{.*}}
2222
sycl::atomic_fence(sycl::memory_order::seq_cst,
2323
sycl::memory_scope::work_group);
2424
});

‎sycl/test/check_device_code/islessgreater.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -6,4 +6,4 @@ SYCL_EXTERNAL void test_islessgreater(float a, float b) {
66
sycl::islessgreater(a, b);
77
}
88
// CHECK-NOT: __spirv_LessOrGreater
9-
// CHECK: {{.*}} = call spir_func zeroext i1 @_Z20__spirv_FOrdNotEqualff(float {{.*}}, float {{.*}})
9+
// CHECK: {{.*}} = call spir_func noundef zeroext i1 @_Z20__spirv_FOrdNotEqualff(float {{.*}}, float {{.*}})

‎sycl/test/check_device_code/no_offset.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -15,7 +15,7 @@ int main() {
1515
sycl::ext::oneapi::accessor_property_list PL{sycl::ext::oneapi::no_offset, sycl::no_init};
1616
sycl::accessor acc_a(a, cgh, sycl::write_only, PL);
1717
sycl::accessor acc_b{b, cgh, sycl::read_only};
18-
// CHECK: define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_EUlT_E_(i32 addrspace(1)* {{.*}}, i32 addrspace(1)* readonly {{.*}}, %"class.cl::sycl::id"* byval(%"class.cl::sycl::id") align 8 {{.*}})
18+
// CHECK: define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_EUlT_E_(i32 addrspace(1)* {{.*}}, i32 addrspace(1)* noundef readonly {{.*}}, %"class.cl::sycl::id"* noundef byval(%"class.cl::sycl::id") align 8 {{.*}})
1919
cgh.parallel_for(size, [=](auto i) {
2020
acc_a[i] = acc_b[i];
2121
});
@@ -33,7 +33,7 @@ int main() {
3333
q.submit([&](sycl::handler &cgh) {
3434
sycl::accessor acc_a(a, cgh, sycl::write_only);
3535
sycl::accessor acc_b{b, cgh, sycl::read_only};
36-
// CHECK: define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE0_clES2_EUlT_E_(i32 addrspace(1)* {{.*}}, %"class.cl::sycl::id"* byval(%"class.cl::sycl::id") align 8 {{.*}}, i32 addrspace(1)* readonly {{.*}}, %"class.cl::sycl::id"* byval(%"class.cl::sycl::id") align 8 {{.*}})
36+
// CHECK: define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE0_clES2_EUlT_E_(i32 addrspace(1)* {{.*}}, %"class.cl::sycl::id"* noundef byval(%"class.cl::sycl::id") align 8 {{.*}}, i32 addrspace(1)* noundef readonly {{.*}}, %"class.cl::sycl::id"* noundef byval(%"class.cl::sycl::id") align 8 {{.*}})
3737
cgh.parallel_for(size, [=](auto i) {
3838
acc_a[i] = acc_b[i];
3939
});

‎sycl/test/check_device_code/usm_pointers.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -10,15 +10,15 @@
1010
// CHECK-ENABLE: %[[DEVPTR_T:.*]] = type { i8 addrspace(5)* }
1111
// CHECK-ENABLE: %[[HOSTPTR_T:.*]] = type { i8 addrspace(6)* }
1212
//
13-
// CHECK-LABEL: define {{.*}} spir_func i8 addrspace(4)* @{{.*}}multi_ptr{{.*}}
13+
// CHECK-LABEL: define {{.*}} spir_func noundef i8 addrspace(4)* @{{.*}}multi_ptr{{.*}}
1414
// CHECK: %[[M_PTR:.*]] = getelementptr inbounds %[[DEVPTR_T]]
1515
// CHECK-DISABLE-NEXT: %[[DEVLOAD:[0-9]+]] = load i8 addrspace(1)*, i8 addrspace(1)* addrspace(4)* %[[M_PTR]]
1616
// CHECK-DISABLE-NEXT: %[[DEVCAST:[0-9]+]] = addrspacecast i8 addrspace(1)* %[[DEVLOAD]] to i8 addrspace(4)*
1717
// CHECK-ENABLE-NEXT: %[[DEVLOAD:[0-9]+]] = load i8 addrspace(5)*, i8 addrspace(5)* addrspace(4)* %[[M_PTR]]
1818
// CHECK-ENABLE-NEXT: %[[DEVCAST:[0-9]+]] = addrspacecast i8 addrspace(5)* %[[DEVLOAD]] to i8 addrspace(4)*
1919
// ret i8 addrspace(4)* %[[DEVCAST]]
2020
//
21-
// CHECK-LABEL: define {{.*}} spir_func i8 addrspace(4)* @{{.*}}multi_ptr{{.*}}
21+
// CHECK-LABEL: define {{.*}} spir_func noundef i8 addrspace(4)* @{{.*}}multi_ptr{{.*}}
2222
// CHECK: %[[M_PTR]] = getelementptr inbounds %[[HOSTPTR_T]]
2323
// CHECK-DISABLE-NEXT: %[[HOSTLOAD:[0-9]+]] = load i8 addrspace(1)*, i8 addrspace(1)* addrspace(4)* %[[M_PTR]]
2424
// CHECK-DISABLE-NEXT: %[[HOSTCAST:[0-9]+]] = addrspacecast i8 addrspace(1)* %[[HOSTLOAD]] to i8 addrspace(4)*

‎sycl/test/esimd/lane_id.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -25,7 +25,7 @@ SYCL_ESIMD_FUNCTION SYCL_EXTERNAL simd<int, 16> foo(int x) {
2525
simd<int, 16> v = 0;
2626
SIMT_BEGIN(16, lane)
2727
//CHECK: define internal spir_func void @_ZZ3fooiENKUlvE_clEv({{.*}}) {{.*}} #[[ATTR:[0-9]+]]
28-
//CHECK: %{{[0-9a-zA-Z_.]+}} = tail call spir_func i32 @_Z15__esimd_lane_idv()
28+
//CHECK: %{{[0-9a-zA-Z_.]+}} = tail call spir_func noundef i32 @_Z15__esimd_lane_idv()
2929
v.select<1, 1>(lane) = x++;
3030
SIMT_END
3131
return v;

‎sycl/test/esimd/math_impl.cpp

Lines changed: 10 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -14,13 +14,13 @@ using namespace sycl::ext::intel::experimental::esimd;
1414
// Math sin,cos,log,exp functions are translated into scalar __spirv_ocl_ calls
1515
SYCL_ESIMD_FUNCTION SYCL_EXTERNAL simd<float, 16> sycl_math(simd<float, 16> x) {
1616
simd<float, 16> v = 0;
17-
//CHECK: call spir_func float @_Z15__spirv_ocl_cosf
17+
//CHECK: call spir_func noundef float @_Z15__spirv_ocl_cosf
1818
v = sycl::cos(x);
19-
//CHECK: call spir_func float @_Z15__spirv_ocl_sinf
19+
//CHECK: call spir_func noundef float @_Z15__spirv_ocl_sinf
2020
v = sycl::sin(v);
21-
//CHECK: call spir_func float @_Z15__spirv_ocl_logf
21+
//CHECK: call spir_func noundef float @_Z15__spirv_ocl_logf
2222
v = sycl::log(v);
23-
//CHECK: call spir_func float @_Z15__spirv_ocl_expf
23+
//CHECK: call spir_func noundef float @_Z15__spirv_ocl_expf
2424
v = sycl::exp(v);
2525
return v;
2626
}
@@ -30,13 +30,13 @@ SYCL_ESIMD_FUNCTION SYCL_EXTERNAL simd<float, 16> sycl_math(simd<float, 16> x) {
3030
SYCL_ESIMD_FUNCTION SYCL_EXTERNAL simd<float, 16>
3131
esimd_math(simd<float, 16> x) {
3232
simd<float, 16> v = 0;
33-
//CHECK: call spir_func <16 x float> @_Z11__esimd_cos
33+
//CHECK: call spir_func noundef <16 x float> @_Z11__esimd_cos
3434
v = esimd::cos(x);
35-
//CHECK: call spir_func <16 x float> @_Z11__esimd_sin
35+
//CHECK: call spir_func noundef <16 x float> @_Z11__esimd_sin
3636
v = esimd::sin(v);
37-
//CHECK: call spir_func <16 x float> @_Z11__esimd_log
37+
//CHECK: call spir_func noundef <16 x float> @_Z11__esimd_log
3838
v = esimd::log2(v);
39-
//CHECK: call spir_func <16 x float> @_Z11__esimd_exp
39+
//CHECK: call spir_func noundef <16 x float> @_Z11__esimd_exp
4040
v = esimd::exp2(v);
4141
return v;
4242
}
@@ -46,9 +46,9 @@ esimd_math(simd<float, 16> x) {
4646
SYCL_ESIMD_FUNCTION SYCL_EXTERNAL simd<float, 16>
4747
esimd_math_emu(simd<float, 16> x) {
4848
simd<float, 16> v = 0;
49-
//CHECK: call spir_func <16 x float> @_Z11__esimd_log
49+
//CHECK: call spir_func noundef <16 x float> @_Z11__esimd_log
5050
v = esimd::log(x);
51-
//CHECK: call spir_func <16 x float> @_Z11__esimd_exp
51+
//CHECK: call spir_func noundef <16 x float> @_Z11__esimd_exp
5252
v = esimd::exp(v);
5353
return v;
5454
}

0 commit comments

Comments
 (0)