Skip to content

Commit 628487a

Browse files
aelovikov-intelaelovikov-intel
committed
[SYCL] Add a test showing miscompilation
1 parent 7eb082c commit 628487a

File tree

1 file changed

+31
-13
lines changed

1 file changed

+31
-13
lines changed

‎sycl/test/check_device_code/extensions/address_cast.cpp‎

Lines changed: 31 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
2-
// RUN: %clangxx -O3 -fsycl -fsycl-device-only -fno-discard-value-names -S -emit-llvm -fno-sycl-instrument-device-code -o - %s | FileCheck %s
2+
// RUN: %clangxx -D__ENABLE_USM_ADDR_SPACE__ -O3 -fsycl -fsycl-device-only -fno-discard-value-names -S -emit-llvm -fno-sycl-instrument-device-code -o - %s | FileCheck %s
33

44
// Linux/Windows have minor differences in the generated IR (e.g. TBAA
55
// metadata). Having linux-only checks eases the maintenance without sacrifising
@@ -23,7 +23,7 @@ namespace static_as_cast {
2323
// CHECK-NEXT: [[ENTRY:.*:]]
2424
// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[P]], align 8, !tbaa [[TBAA8:![0-9]+]]
2525
// CHECK-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr addrspace(4)
26-
// CHECK-NEXT: [[CALL_I_I_I_I:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z33__spirv_GenericCastToPtr_ToGlobalPvi(ptr addrspace(4) noundef [[TMP1]], i32 noundef 5) #[[ATTR5:[0-9]+]]
26+
// CHECK-NEXT: [[CALL_I_I_I_I:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z33__spirv_GenericCastToPtr_ToGlobalPvi(ptr addrspace(4) noundef [[TMP1]], i32 noundef 5) #[[ATTR6:[0-9]+]]
2727
// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA12:![0-9]+]], !alias.scope [[META14:![0-9]+]]
2828
// CHECK-NEXT: ret void
2929
//
@@ -33,7 +33,7 @@ SYCL_EXTERNAL auto to_global_decorated(decorated_generic_ptr<int> p) {
3333
// CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast23to_global_not_decoratedEPi(
3434
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.1") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR2:[0-9]+]] !srcloc [[META19:![0-9]+]] !sycl_fixed_targets [[META7]] {
3535
// CHECK-NEXT: [[ENTRY:.*:]]
36-
// CHECK-NEXT: [[CALL_I_I_I:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z33__spirv_GenericCastToPtr_ToGlobalPvi(ptr addrspace(4) noundef [[P]], i32 noundef 5) #[[ATTR5]]
36+
// CHECK-NEXT: [[CALL_I_I_I:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z33__spirv_GenericCastToPtr_ToGlobalPvi(ptr addrspace(4) noundef [[P]], i32 noundef 5) #[[ATTR6]]
3737
// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA20:![0-9]+]], !alias.scope [[META22:![0-9]+]]
3838
// CHECK-NEXT: ret void
3939
//
@@ -59,45 +59,63 @@ SYCL_EXTERNAL auto to_generic_decorated(decorated_generic_ptr<int> p) {
5959
SYCL_EXTERNAL auto to_generic_not_decorated(int *p) {
6060
return static_address_cast<access::address_space::generic_space>(p);
6161
}
62+
63+
// CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast16to_global_deviceEPi(
64+
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture readnone sret(%"class.sycl::_V1::multi_ptr.3") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) nocapture noundef readnone [[P:%.*]]) local_unnamed_addr #[[ATTR5:[0-9]+]] !srcloc [[META35:![0-9]+]] !sycl_fixed_targets [[META7]] {
65+
// CHECK-NEXT: [[ENTRY:.*:]]
66+
// CHECK-NEXT: unreachable
67+
//
68+
SYCL_EXTERNAL auto to_global_device(int *p) {
69+
return static_address_cast<access::address_space::ext_intel_global_device_space>(p);
70+
}
71+
72+
// CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast14to_global_hostEPi(
73+
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture readnone sret(%"class.sycl::_V1::multi_ptr.4") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) nocapture noundef readnone [[P:%.*]]) local_unnamed_addr #[[ATTR5]] !srcloc [[META36:![0-9]+]] !sycl_fixed_targets [[META7]] {
74+
// CHECK-NEXT: [[ENTRY:.*:]]
75+
// CHECK-NEXT: unreachable
76+
//
77+
SYCL_EXTERNAL auto to_global_host(int *p) {
78+
return static_address_cast<access::address_space::ext_intel_global_host_space>(p);
79+
}
6280
} // namespace static_as_cast
6381

6482
namespace dynamic_as_cast {
6583
// CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast19to_global_decoratedEN4sycl3_V19multi_ptrIiLNS1_6access13address_spaceE6ELNS3_9decoratedE1EEE(
66-
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr") align 8 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 [[P:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META35:![0-9]+]] !sycl_fixed_targets [[META7]] {
84+
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr") align 8 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 [[P:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META37:![0-9]+]] !sycl_fixed_targets [[META7]] {
6785
// CHECK-NEXT: [[ENTRY:.*:]]
6886
// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[P]], align 8, !tbaa [[TBAA8]]
6987
// CHECK-NEXT: [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr addrspace(4)
70-
// CHECK-NEXT: [[CALL_I_I_I_I:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z41__spirv_GenericCastToPtrExplicit_ToGlobalPvi(ptr addrspace(4) noundef [[TMP1]], i32 noundef 5) #[[ATTR5]]
71-
// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA12]], !alias.scope [[META36:![0-9]+]]
88+
// CHECK-NEXT: [[CALL_I_I_I_I:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z41__spirv_GenericCastToPtrExplicit_ToGlobalPvi(ptr addrspace(4) noundef [[TMP1]], i32 noundef 5) #[[ATTR6]]
89+
// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA12]], !alias.scope [[META38:![0-9]+]]
7290
// CHECK-NEXT: ret void
7391
//
7492
SYCL_EXTERNAL auto to_global_decorated(decorated_generic_ptr<int> p) {
7593
return dynamic_address_cast<access::address_space::global_space>(p);
7694
}
7795
// CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast23to_global_not_decoratedEPi(
78-
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.1") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR2]] !srcloc [[META41:![0-9]+]] !sycl_fixed_targets [[META7]] {
96+
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.1") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR2]] !srcloc [[META43:![0-9]+]] !sycl_fixed_targets [[META7]] {
7997
// CHECK-NEXT: [[ENTRY:.*:]]
80-
// CHECK-NEXT: [[CALL_I_I_I:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z41__spirv_GenericCastToPtrExplicit_ToGlobalPvi(ptr addrspace(4) noundef [[P]], i32 noundef 5) #[[ATTR5]]
81-
// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA20]], !alias.scope [[META42:![0-9]+]]
98+
// CHECK-NEXT: [[CALL_I_I_I:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z41__spirv_GenericCastToPtrExplicit_ToGlobalPvi(ptr addrspace(4) noundef [[P]], i32 noundef 5) #[[ATTR6]]
99+
// CHECK-NEXT: store ptr addrspace(1) [[CALL_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA20]], !alias.scope [[META44:![0-9]+]]
82100
// CHECK-NEXT: ret void
83101
//
84102
SYCL_EXTERNAL auto to_global_not_decorated(int *p) {
85103
return dynamic_address_cast<access::address_space::global_space>(p);
86104
}
87105
// CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast20to_generic_decoratedEN4sycl3_V19multi_ptrIiLNS1_6access13address_spaceE6ELNS3_9decoratedE1EEE(
88-
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.0") align 8 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 [[P:%.*]]) local_unnamed_addr #[[ATTR3]] !srcloc [[META45:![0-9]+]] !sycl_fixed_targets [[META7]] {
106+
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.0") align 8 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 [[P:%.*]]) local_unnamed_addr #[[ATTR3]] !srcloc [[META47:![0-9]+]] !sycl_fixed_targets [[META7]] {
89107
// CHECK-NEXT: [[ENTRY:.*:]]
90108
// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[P]], align 8, !tbaa [[TBAA8]]
91-
// CHECK-NEXT: store i64 [[TMP0]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA8]], !alias.scope [[META46:![0-9]+]]
109+
// CHECK-NEXT: store i64 [[TMP0]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA8]], !alias.scope [[META48:![0-9]+]]
92110
// CHECK-NEXT: ret void
93111
//
94112
SYCL_EXTERNAL auto to_generic_decorated(decorated_generic_ptr<int> p) {
95113
return dynamic_address_cast<access::address_space::generic_space>(p);
96114
}
97115
// CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast24to_generic_not_decoratedEPi(
98-
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.2") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR4]] !srcloc [[META49:![0-9]+]] !sycl_fixed_targets [[META7]] {
116+
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.2") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR4]] !srcloc [[META51:![0-9]+]] !sycl_fixed_targets [[META7]] {
99117
// CHECK-NEXT: [[ENTRY:.*:]]
100-
// CHECK-NEXT: store ptr addrspace(4) [[P]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA30]], !alias.scope [[META50:![0-9]+]]
118+
// CHECK-NEXT: store ptr addrspace(4) [[P]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA30]], !alias.scope [[META52:![0-9]+]]
101119
// CHECK-NEXT: ret void
102120
//
103121
SYCL_EXTERNAL auto to_generic_not_decorated(int *p) {

0 commit comments

Comments
 (0)