Skip to content

Commit ab05021

Browse files
baderbader
committed
Update SYCL lit tests.
1 parent 3c43979 commit ab05021

File tree

4 files changed

+38
-38
lines changed

4 files changed

+38
-38
lines changed

‎sycl/test/check_device_code/id_queries_fit_int.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -4,7 +4,7 @@
44

55
using namespace sycl;
66

7-
// CHECK: define {{.*}}dso_local spir_func void @{{.*}}testItem{{.*}}(%"class.{{.*}}item"*{{.*}}%{{.*}})
7+
// CHECK: define {{.*}}dso_local spir_func void @{{.*}}testItem{{.*}}(%"class.{{.*}}item"{{.*}}%{{.*}})
88
SYCL_EXTERNAL void testItem(item<1> TestItem) {
99
// CHECK: call void @llvm.assume(i1 {{.*}})
1010
int Id = TestItem.get_id(0);
@@ -14,7 +14,7 @@ SYCL_EXTERNAL void testItem(item<1> TestItem) {
1414
int LinearId = TestItem.get_linear_id();
1515
}
1616

17-
// CHECK: define {{.*}}dso_local spir_func void @{{.*}}testNDItem{{.*}}(%"class.{{.*}}nd_item"*{{.*}}%{{.*}})
17+
// CHECK: define {{.*}}dso_local spir_func void @{{.*}}testNDItem{{.*}}(%"class.{{.*}}nd_item"{{.*}}%{{.*}})
1818
SYCL_EXTERNAL void testNDItem(nd_item<1> TestNDItem) {
1919
// CHECK: call void @llvm.assume(i1 {{.*}})
2020
int GlobalId = TestNDItem.get_global_id(0);

‎sycl/test/check_device_code/kernel_arguments_as.cpp

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1,16 +1,16 @@
1-
// RUN: %clangxx -fsycl-device-only -Xclang -fsycl-is-device -emit-llvm %s -S -o %t.ll -I %sycl_include -Wno-sycl-strict -Xclang -verify-ignore-unexpected=note,warning -Xclang -disable-llvm-passes
1+
// RUN: %clangxx -fsycl-device-only -Xclang -opaque-pointers -emit-llvm %s -S -o %t.ll -I %sycl_include -Wno-sycl-strict -Xclang -verify-ignore-unexpected=note,warning -Xclang -disable-llvm-passes
22
// RUN: FileCheck %s --input-file %t.ll --check-prefixes=CHECK,CHECK-DISABLE
3-
// RUN: %clangxx -fsycl-device-only -Xclang -fsycl-is-device -emit-llvm %s -S -o %t.ll -I %sycl_include -Wno-sycl-strict -Xclang -verify-ignore-unexpected=note,warning -Xclang -disable-llvm-passes -D__ENABLE_USM_ADDR_SPACE__
3+
// RUN: %clangxx -fsycl-device-only -Xclang -opaque-pointers -emit-llvm %s -S -o %t.ll -I %sycl_include -Wno-sycl-strict -Xclang -verify-ignore-unexpected=note,warning -Xclang -disable-llvm-passes -D__ENABLE_USM_ADDR_SPACE__
44
// RUN: FileCheck %s --input-file %t.ll --check-prefixes=CHECK,CHECK-ENABLE
55
//
66
// Check the address space of the pointer in accessor class.
77
//
88
// CHECK: %struct.AccWrapper = type { %"class.cl::sycl::accessor[[NUMBER_SUFFIX:\.?[0-9]*]]" }
99
// CHECK: %"class.cl::sycl::accessor[[NUMBER_SUFFIX]]" = type { %"class{{.*}}AccessorImplDevice", %[[UNION:.*]] }
10-
// CHECK-DISABLE: %[[UNION]] = type { i32 addrspace(1)* }
11-
// CHECK-ENABLE: %[[UNION]] = type { i32 addrspace(5)* }
10+
// CHECK-DISABLE: %[[UNION]] = type { ptr addrspace(1) }
11+
// CHECK-ENABLE: %[[UNION]] = type { ptr addrspace(5) }
1212
// CHECK: %struct.AccWrapper.{{[0-9]+}} = type { %"class.cl::sycl::accessor.[[NUM:[0-9]+]]" }
13-
// CHECK-NEXT: %"class.cl::sycl::accessor.[[NUM]]" = type { %"class{{.*}}LocalAccessorBaseDevice", i32 addrspace(3)* }
13+
// CHECK-NEXT: %"class.cl::sycl::accessor.[[NUM]]" = type { %"class{{.*}}LocalAccessorBaseDevice", ptr addrspace(3) }
1414
//
1515
// Check that kernel arguments doesn't have generic address space.
1616
//

‎sycl/test/check_device_code/no_offset.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clangxx -fsycl-device-only -fsycl-early-optimizations -fsycl-dead-args-optimization -D__SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ -S -emit-llvm -o - %s | FileCheck %s
1+
// RUN: %clangxx -fsycl-device-only -Xclang -opaque-pointers -fsycl-early-optimizations -fsycl-dead-args-optimization -D__SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ -S -emit-llvm -o - %s | FileCheck %s
22

33
#include <sycl/sycl.hpp>
44

@@ -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)* noundef readonly {{.*}}, %"class.cl::sycl::id"* noundef byval(%"class.cl::sycl::id") align 8 {{.*}})
18+
// CHECK: define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_EUlT_E_(ptr addrspace(1) {{.*}}, ptr addrspace(1) noundef readonly {{.*}}, ptr 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"* 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 {{.*}})
36+
// CHECK: define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE0_clES2_EUlT_E_(ptr addrspace(1) {{.*}}, ptr noundef byval(%"class.cl::sycl::id") align 8 {{.*}}, ptr addrspace(1) noundef readonly {{.*}}, ptr 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/extensions/sub_group_as.cpp

Lines changed: 28 additions & 28 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
1-
// RUN: %clangxx -fsycl -fsycl-device-only -O3 -S -emit-llvm -x c++ -Xclang -no-enable-noundef-analysis %s -o - | FileCheck %s --check-prefix CHECK-O3
2-
// RUN: %clangxx -fsycl -fsycl-device-only -O0 -S -emit-llvm -x c++ -Xclang -no-enable-noundef-analysis %s -o - | FileCheck %s --check-prefix CHECK-O0
1+
// RUN: %clangxx -fsycl-device-only -O3 -S -emit-llvm -Xclang -opaque-pointers -Xclang -no-enable-noundef-analysis %s -o - | FileCheck %s --check-prefix CHECK-O3
2+
// RUN: %clangxx -fsycl-device-only -O0 -S -emit-llvm -Xclang -opaque-pointers -Xclang -no-enable-noundef-analysis %s -o - | FileCheck %s --check-prefix CHECK-O0
33
// Test compilation with -O3 when all methods are inlined in kernel function
44
// and -O0 when helper methods are preserved.
55
#include <cassert>
@@ -62,54 +62,54 @@ int main(int argc, char *argv[]) {
6262
// CHECK-O3: call spir_func void {{.*}}spirv_ControlBarrierjjj
6363

6464
// load() for global address space
65-
// CHECK-O3: call spir_func i8 addrspace(3)* {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(i8 addrspace(4)*
65+
// CHECK-O3: call spir_func ptr addrspace(3) {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(ptr addrspace(4)
6666
// CHECK-O3: {{.*}}SubgroupLocalInvocationId
67-
// CHECK-O3: call spir_func i8 addrspace(1)* {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(i8 addrspace(4)*
68-
// CHECK-O3: call spir_func i32 {{.*}}spirv_SubgroupBlockRead{{.*}}(i32 addrspace(1)*
67+
// CHECK-O3: call spir_func ptr addrspace(1) {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4)
68+
// CHECK-O3: call spir_func i32 {{.*}}spirv_SubgroupBlockRead{{.*}}(ptr addrspace(1)
6969
// CHECK-O3: call spir_func void {{.*}}assert
7070

7171

7272
// load() for local address space
73-
// CHECK-O3: call spir_func i8 addrspace(3)* {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(i8 addrspace(4)*
73+
// CHECK-O3: call spir_func ptr addrspace(3) {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(ptr addrspace(4)
7474
// CHECK-O3: {{.*}}SubgroupLocalInvocationId
75-
// CHECK-O3: call spir_func i8 addrspace(1)* {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(i8 addrspace(4)*
76-
// CHECK-O3: call spir_func i32 {{.*}}spirv_SubgroupBlockRead{{.*}}(i32 addrspace(1)*
75+
// CHECK-O3: call spir_func ptr addrspace(1) {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4)
76+
// CHECK-O3: call spir_func i32 {{.*}}spirv_SubgroupBlockRead{{.*}}(ptr addrspace(1)
7777
// CHECK-O3: call spir_func void {{.*}}assert
7878

7979
// load() for private address space
80-
// CHECK-O3: call spir_func i8 addrspace(3)* {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(i8 addrspace(4)*
80+
// CHECK-O3: call spir_func ptr addrspace(3) {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(ptr addrspace(4)
8181
// CHECK-O3: {{.*}}SubgroupLocalInvocationId
82-
// CHECK-O3: call spir_func i8 addrspace(1)* {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(i8 addrspace(4)*
83-
// CHECK-O3: call spir_func i32 {{.*}}spirv_SubgroupBlockRead{{.*}}(i32 addrspace(1)*
82+
// CHECK-O3: call spir_func ptr addrspace(1) {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4)
83+
// CHECK-O3: call spir_func i32 {{.*}}spirv_SubgroupBlockRead{{.*}}(ptr addrspace(1)
8484
// CHECK-O3: call spir_func void {{.*}}assert
8585

8686
// store() for global address space
87-
// CHECK-O3: call spir_func i8 addrspace(3)* {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(i8 addrspace(4)*
87+
// CHECK-O3: call spir_func ptr addrspace(3) {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(ptr addrspace(4)
8888
// CHECK-O3: {{.*}}SubgroupLocalInvocationId
89-
// CHECK-O3: call spir_func i8 addrspace(1)* {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(i8 addrspace(4)*
90-
// CHECK-O3: call spir_func void {{.*}}spirv_SubgroupBlockWriteINTEL{{.*}}(i32 addrspace(1)*
89+
// CHECK-O3: call spir_func ptr addrspace(1) {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4)
90+
// CHECK-O3: call spir_func void {{.*}}spirv_SubgroupBlockWriteINTEL{{.*}}(ptr addrspace(1)
9191
// CHECK-O3: call spir_func void {{.*}}assert
9292

9393
// load() accepting raw pointers method
94-
// CHECK-O0: define{{.*}}spir_func i32 {{.*}}cl4sycl3ext6oneapi9sub_group4load{{.*}}addrspace(4)* %
95-
// CHECK-O0: call spir_func i32 addrspace(3)* {{.*}}SYCL_GenericCastToPtrExplicit_ToLocal{{.*}}(i8 addrspace(4)*
96-
// CHECK-O0: call spir_func i32 {{.*}}sycl3ext6oneapi9sub_group4load{{.*}}i32 addrspace(3)* %
97-
// CHECK-O0: call spir_func i32 addrspace(1)* {{.*}}SYCL_GenericCastToPtrExplicit_ToGlobal{{.*}}(i8 addrspace(4)*
98-
// CHECK-O0: call spir_func i32 {{.*}}sycl3ext6oneapi9sub_group4load{{.*}}i32 addrspace(1)* %
94+
// CHECK-O0: define{{.*}}spir_func i32 {{.*}}cl4sycl3ext6oneapi9sub_group4load{{.*}}addrspace(4) %
95+
// CHECK-O0: call spir_func ptr addrspace(3) {{.*}}SYCL_GenericCastToPtrExplicit_ToLocal{{.*}}(ptr addrspace(4)
96+
// CHECK-O0: call spir_func i32 {{.*}}sycl3ext6oneapi9sub_group4load{{.*}}ptr addrspace(3) %
97+
// CHECK-O0: call spir_func ptr addrspace(1) {{.*}}SYCL_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4)
98+
// CHECK-O0: call spir_func i32 {{.*}}sycl3ext6oneapi9sub_group4load{{.*}}ptr addrspace(1) %
9999
// CHECK-O0: call spir_func void {{.*}}assert
100100

101101
// store() accepting raw pointers method
102-
// CHECK-O0: define{{.*}}spir_func void {{.*}}cl4sycl3ext6oneapi9sub_group5store{{.*}}i32 addrspace(4)* %
103-
// CHECK-O0: call spir_func i32 addrspace(3)* {{.*}}SYCL_GenericCastToPtrExplicit_ToLocal{{.*}}(i8 addrspace(4)*
104-
// CHECK-O0: call spir_func void {{.*}}cl4sycl3ext6oneapi9sub_group5store{{.*}}, i32 addrspace(3)* %
105-
// CHECK-O0: call spir_func i32 addrspace(1)* {{.*}}SYCL_GenericCastToPtrExplicit_ToGlobal{{.*}}(i8 addrspace(4)*
106-
// CHECK-O0: call spir_func void {{.*}}cl4sycl3ext6oneapi9sub_group5store{{.*}}, i32 addrspace(1)* %
102+
// CHECK-O0: define{{.*}}spir_func void {{.*}}cl4sycl3ext6oneapi9sub_group5store{{.*}}ptr addrspace(4) %
103+
// CHECK-O0: call spir_func ptr addrspace(3) {{.*}}SYCL_GenericCastToPtrExplicit_ToLocal{{.*}}(ptr addrspace(4)
104+
// CHECK-O0: call spir_func void {{.*}}cl4sycl3ext6oneapi9sub_group5store{{.*}}, ptr addrspace(3) %
105+
// CHECK-O0: call spir_func ptr addrspace(1) {{.*}}SYCL_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4)
106+
// CHECK-O0: call spir_func void {{.*}}cl4sycl3ext6oneapi9sub_group5store{{.*}}, ptr addrspace(1) %
107107
// CHECK-O0: call spir_func void {{.*}}assert
108108

109-
// CHECK-O0: define {{.*}}spir_func i32 addrspace(3)* {{.*}}SYCL_GenericCastToPtrExplicit_ToLocal{{.*}}(i8 addrspace(4)* %
110-
// CHECK-O0: call spir_func i8 addrspace(3)* {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(i8 addrspace(4)*
111-
// CHECK-O0: define {{.*}}spir_func i32 addrspace(1)* {{.*}}SYCL_GenericCastToPtrExplicit_ToGlobal{{.*}}(i8 addrspace(4)* %
112-
// CHECK-O0: call spir_func i8 addrspace(1)* {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(i8 addrspace(4)*
109+
// CHECK-O0: define {{.*}}spir_func ptr addrspace(3) {{.*}}SYCL_GenericCastToPtrExplicit_ToLocal{{.*}}(ptr addrspace(4) %
110+
// CHECK-O0: call spir_func ptr addrspace(3) {{.*}}spirv_GenericCastToPtrExplicit_ToLocal{{.*}}(ptr addrspace(4)
111+
// CHECK-O0: define {{.*}}spir_func ptr addrspace(1) {{.*}}SYCL_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4) %
112+
// CHECK-O0: call spir_func ptr addrspace(1) {{.*}}spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}(ptr addrspace(4)
113113
// clang-format off
114114

115115
return 0;

0 commit comments

Comments
 (0)