diff --git a/sycl/test/check_device_code/extensions/properties/properties_cache_control.cpp b/sycl/test/check_device_code/extensions/properties/properties_cache_control.cpp index 76faf52dfb18b..6497dd0e70730 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_cache_control.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_cache_control.cpp @@ -1,5 +1,5 @@ -// RUN: %clangxx -fsycl-device-only -S -Xclang -emit-llvm %s -o - | \ -// RUN: FileCheck %s --check-prefix CHECK-IR +// RUN: %clangxx -fsycl-device-only -fsycl-targets=spir64 -S -emit-llvm %s -o - | \ +// RUN: FileCheck %s #include @@ -127,7 +127,7 @@ void cache_control_load_store_func() { auto d_h = d_buf_h; auto kernel = - [=](nd_item<2> item) [[intel::reqd_sub_group_size(SG_SIZE)]] { + [=](nd_item<2> item) [[sycl::reqd_sub_group_size(SG_SIZE)]] { const int global_tid = item.get_global_id(0); const int row_st = global_tid * SG_SIZE; @@ -170,58 +170,58 @@ SYCL_EXTERNAL void annotated_ptr_func_param_test(float *p) { *(store_hint{p}) = 42.0f; } -// CHECK-IR: spir_func{{.*}}annotated_ptr_func_param_test -// CHECK-IR: {{.*}}call ptr addrspace(4) @llvm.ptr.annotation.p4.p1{{.*}}!spirv.Decorations [[WHINT:.*]] -// CHECK-IR: ret void - -// CHECK-IR: spir_kernel{{.*}}cache_control_read_hint_func -// CHECK-IR: {{.*}}addrspacecast ptr addrspace(1){{.*}}!spirv.Decorations [[RHINT:.*]] -// CHECK-IR: ret void - -// CHECK-IR: spir_kernel{{.*}}cache_control_read_assertion_func -// CHECK-IR: {{.*}}addrspacecast ptr addrspace(1){{.*}}!spirv.Decorations [[RASSERT:.*]] -// CHECK-IR: ret void - -// CHECK-IR: spir_kernel{{.*}}cache_control_write_hint_func -// CHECK-IR: {{.*}}addrspacecast ptr addrspace(1){{.*}}!spirv.Decorations [[WHINT]] -// CHECK-IR: ret void - -// CHECK-IR: spir_kernel{{.*}}cache_control_read_write_func -// CHECK-IR: {{.*}}addrspacecast ptr addrspace(1){{.*}}!spirv.Decorations [[RWHINT:.*]] -// CHECK-IR: ret void - -// CHECK-IR: spir_kernel{{.*}}cache_control_load_store_func -// CHECK-IR: {{.*}}getelementptr{{.*}}addrspace(4){{.*}}!spirv.Decorations [[LDSTHINT_A:.*]] -// CHECK-IR: {{.*}}getelementptr{{.*}}addrspace(4){{.*}}!spirv.Decorations [[LDSTHINT_B:.*]] -// CHECK-IR: ret void - -// CHECK-IR: [[WHINT]] = !{[[WHINT1:.*]], [[WHINT2:.*]], [[WHINT3:.*]], [[WHINT4:.*]]} -// CHECK-IR: [[WHINT1]] = !{i32 6443, i32 3, i32 3} -// CHECK-IR: [[WHINT2]] = !{i32 6443, i32 0, i32 1} -// CHECK-IR: [[WHINT3]] = !{i32 6443, i32 1, i32 2} -// CHECK-IR: [[WHINT4]] = !{i32 6443, i32 2, i32 2} - -// CHECK-IR: [[RHINT]] = !{[[RHINT1:.*]], [[RHINT2:.*]], [[RHINT3:.*]]} -// CHECK-IR: [[RHINT1]] = !{i32 6442, i32 1, i32 0} -// CHECK-IR: [[RHINT2]] = !{i32 6442, i32 2, i32 0} -// CHECK-IR: [[RHINT3]] = !{i32 6442, i32 0, i32 1} - -// CHECK-IR: [[RASSERT]] = !{[[RASSERT1:.*]], [[RASSERT2:.*]], [[RASSERT3:.*]]} -// CHECK-IR: [[RASSERT1]] = !{i32 6442, i32 1, i32 3} -// CHECK-IR: [[RASSERT2]] = !{i32 6442, i32 2, i32 3} -// CHECK-IR: [[RASSERT3]] = !{i32 6442, i32 0, i32 4} - -// CHECK-IR: [[RWHINT]] = !{[[RWHINT1:.*]], [[RWHINT2:.*]], [[RWHINT3:.*]]} -// CHECK-IR: [[RWHINT1]] = !{i32 6442, i32 2, i32 1} -// CHECK-IR: [[RWHINT2]] = !{i32 6442, i32 3, i32 4} -// CHECK-IR: [[RWHINT3]] = !{i32 6443, i32 3, i32 1} - -// CHECK-IR: [[LDSTHINT_A]] = !{[[RHINT1]], [[RHINT2]], [[RHINT3]], [[LDSTHINT_A1:.*]], [[LDSTHINT_A2:.*]], [[LDSTHINT_A3:.*]]} -// CHECK-IR: [[LDSTHINT_A1]] = !{i32 6443, i32 0, i32 0} -// CHECK-IR: [[LDSTHINT_A2]] = !{i32 6443, i32 1, i32 0} -// CHECK-IR: [[LDSTHINT_A3]] = !{i32 6443, i32 2, i32 0} - -// CHECK-IR: [[LDSTHINT_B]] = !{[[LDSTHINT_B1:.*]], [[RWHINT1]], [[LDSTHINT_B2:.*]], [[LDSTHINT_A2]], [[LDSTHINT_A3]], [[LDSTHINT_B3:.*]]} -// CHECK-IR: [[LDSTHINT_B1]] = !{i32 6442, i32 1, i32 1} -// CHECK-IR: [[LDSTHINT_B2]] = !{i32 6442, i32 0, i32 2} -// CHECK-IR: [[LDSTHINT_B3]] = !{i32 6443, i32 0, i32 2} +// CHECK: spir_func{{.*}}annotated_ptr_func_param_test +// CHECK: {{.*}}call ptr addrspace(4) @llvm.ptr.annotation.p4.p1{{.*}}!spirv.Decorations [[WHINT:.*]] +// CHECK: ret void + +// CHECK: spir_kernel{{.*}}cache_control_read_hint_func +// CHECK: {{.*}}addrspacecast ptr addrspace(1){{.*}}!spirv.Decorations [[RHINT:.*]] +// CHECK: ret void + +// CHECK: spir_kernel{{.*}}cache_control_read_assertion_func +// CHECK: {{.*}}addrspacecast ptr addrspace(1){{.*}}!spirv.Decorations [[RASSERT:.*]] +// CHECK: ret void + +// CHECK: spir_kernel{{.*}}cache_control_write_hint_func +// CHECK: {{.*}}addrspacecast ptr addrspace(1){{.*}}!spirv.Decorations [[WHINT]] +// CHECK: ret void + +// CHECK: spir_kernel{{.*}}cache_control_read_write_func +// CHECK: {{.*}}addrspacecast ptr addrspace(1){{.*}}!spirv.Decorations [[RWHINT:.*]] +// CHECK: ret void + +// CHECK: spir_kernel{{.*}}cache_control_load_store_func +// CHECK: {{.*}}getelementptr{{.*}}addrspace(4){{.*}}!spirv.Decorations [[LDSTHINT_A:.*]] +// CHECK: {{.*}}getelementptr{{.*}}addrspace(4){{.*}}!spirv.Decorations [[LDSTHINT_B:.*]] +// CHECK: ret void + +// CHECK: [[WHINT]] = !{[[WHINT1:.*]], [[WHINT2:.*]], [[WHINT3:.*]], [[WHINT4:.*]]} +// CHECK: [[WHINT1]] = !{i32 6443, i32 3, i32 3} +// CHECK: [[WHINT2]] = !{i32 6443, i32 0, i32 1} +// CHECK: [[WHINT3]] = !{i32 6443, i32 1, i32 2} +// CHECK: [[WHINT4]] = !{i32 6443, i32 2, i32 2} + +// CHECK: [[RHINT]] = !{[[RHINT1:.*]], [[RHINT2:.*]], [[RHINT3:.*]]} +// CHECK: [[RHINT1]] = !{i32 6442, i32 1, i32 0} +// CHECK: [[RHINT2]] = !{i32 6442, i32 2, i32 0} +// CHECK: [[RHINT3]] = !{i32 6442, i32 0, i32 1} + +// CHECK: [[RASSERT]] = !{[[RASSERT1:.*]], [[RASSERT2:.*]], [[RASSERT3:.*]]} +// CHECK: [[RASSERT1]] = !{i32 6442, i32 1, i32 3} +// CHECK: [[RASSERT2]] = !{i32 6442, i32 2, i32 3} +// CHECK: [[RASSERT3]] = !{i32 6442, i32 0, i32 4} + +// CHECK: [[RWHINT]] = !{[[RWHINT1:.*]], [[RWHINT2:.*]], [[RWHINT3:.*]]} +// CHECK: [[RWHINT1]] = !{i32 6442, i32 2, i32 1} +// CHECK: [[RWHINT2]] = !{i32 6442, i32 3, i32 4} +// CHECK: [[RWHINT3]] = !{i32 6443, i32 3, i32 1} + +// CHECK: [[LDSTHINT_A]] = !{[[RHINT1]], [[RHINT2]], [[RHINT3]], [[LDSTHINT_A1:.*]], [[LDSTHINT_A2:.*]], [[LDSTHINT_A3:.*]]} +// CHECK: [[LDSTHINT_A1]] = !{i32 6443, i32 0, i32 0} +// CHECK: [[LDSTHINT_A2]] = !{i32 6443, i32 1, i32 0} +// CHECK: [[LDSTHINT_A3]] = !{i32 6443, i32 2, i32 0} + +// CHECK: [[LDSTHINT_B]] = !{[[LDSTHINT_B1:.*]], [[RWHINT1]], [[LDSTHINT_B2:.*]], [[LDSTHINT_A2]], [[LDSTHINT_A3]], [[LDSTHINT_B3:.*]]} +// CHECK: [[LDSTHINT_B1]] = !{i32 6442, i32 1, i32 1} +// CHECK: [[LDSTHINT_B2]] = !{i32 6442, i32 0, i32 2} +// CHECK: [[LDSTHINT_B3]] = !{i32 6443, i32 0, i32 2}