Skip to content

Commit ccc1c90

Browse files
author
iclsrc
committed
Merge from 'sycl' to 'sycl-web'
2 parents 56b58d0 + 6ba7b52 commit ccc1c90

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

49 files changed

+1490
-212
lines changed

‎clang/lib/Driver/ToolChains/SYCL.h

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -194,7 +194,8 @@ class LLVM_LIBRARY_VISIBILITY SYCLToolChain : public ToolChain {
194194
return false;
195195
}
196196
llvm::codegenoptions::DebugInfoFormat getDefaultDebugFormat() const override {
197-
if (this->IsSYCLNativeCPU)
197+
if (this->IsSYCLNativeCPU ||
198+
this->HostTC.getTriple().isWindowsMSVCEnvironment())
198199
return this->HostTC.getDefaultDebugFormat();
199200
return ToolChain::getDefaultDebugFormat();
200201
}

‎clang/test/Driver/sycl.c

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -154,3 +154,12 @@
154154
// RUN: %clang -fsycl -fsycl-targets=nvptx64-nvidia-cuda,amdgcn-amd-amdhsa -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx908 -Xsycl-target-backend=nvptx64-nvidia-cuda --offload-arch=sm_86 -c -ccc-print-phases %s 2>&1 | FileCheck %s --check-prefix=MULTIPLE_TARGETS
155155
// MULTIPLE_TARGETS: offload, "device-sycl (nvptx64-nvidia-cuda:sm_86)"
156156
// MULTIPLE_TARGETS: offload, "device-sycl (amdgcn-amd-amdhsa:gfx908)"
157+
158+
/// Verify debug format for spir target.
159+
// RUN: %clang -### -target x86_64-windows-msvc -fsycl -g -c %s 2>&1 \
160+
// RUN: | FileCheck -check-prefix=DEBUG-WIN %s
161+
// RUN: %clang_cl -### -fsycl -Zi -c %s 2>&1 \
162+
// RUN: | FileCheck -check-prefix=DEBUG-WIN %s
163+
// DEBUG-WIN: {{.*}}"-fsycl-is-device"{{.*}}"-gcodeview"
164+
// DEBUG-WIN: {{.*}}"-fsycl-is-host"{{.*}}"-gcodeview"
165+
// DEBUG-WIN-NOT: dwarf-version

‎llvm/lib/SYCLLowerIR/CompileTimeProperties.def

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -54,3 +54,9 @@ SYCL_COMPILE_TIME_PROPERTY("sycl-stride-size", 5883, DecorValueTy::uint32)
5454
SYCL_COMPILE_TIME_PROPERTY("sycl-word-size", 5884, DecorValueTy::uint32)
5555
SYCL_COMPILE_TIME_PROPERTY("sycl-bi-directional-ports-true", 5885,
5656
DecorValueTy::none)
57+
58+
// The corresponding SPIR-V OpCodes for sycl_ext_oneapi_prefetch
59+
// SYCL Spec: https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_prefetch.asciidoc
60+
// SPIR-V Spec: https://github.com/KhronosGroup/SPIRV-Registry/blob/main/extensions/INTEL/SPV_INTEL_cache_controls.asciidoc
61+
SYCL_COMPILE_TIME_PROPERTY("sycl-prefetch-hint", 6442, DecorValueTy::uint32)
62+
SYCL_COMPILE_TIME_PROPERTY("sycl-prefetch-hint-nt", 6442, DecorValueTy::uint32)

‎llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp

Lines changed: 9 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -645,7 +645,15 @@ bool CompileTimePropertiesPass::transformSYCLPropertiesAnnotation(
645645
// always valid, even if the decoration parameters are not strings.
646646
NewAnnotString += "{" + std::to_string(DecorCode);
647647
if (PropVal)
648-
NewAnnotString += ":\"" + PropVal->str() + "\"";
648+
NewAnnotString += ":\"" + PropVal->str();
649+
650+
if (PropName == "sycl-prefetch-hint")
651+
NewAnnotString += ",1"; // CachedINTEL
652+
if (PropName == "sycl-prefetch-hint-nt")
653+
NewAnnotString += ",3"; // InvalidateAfterReadINTEL
654+
655+
if (PropVal)
656+
NewAnnotString += "\"";
649657
NewAnnotString += "}";
650658
}
651659

‎llvm/test/SYCLLowerIR/CompileTimePropertiesPass/prefetch.ll

Lines changed: 66 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,66 @@
1+
; RUN: opt -passes=compile-time-properties %s -S | FileCheck %s
2+
3+
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
4+
target triple = "spir64-unknown-unknown"
5+
6+
$_ZTSZ4mainEUlvE_ = comdat any
7+
8+
@.str = private unnamed_addr addrspace(1) constant [16 x i8] c"sycl-properties\00", section "llvm.metadata"
9+
@.str.1 = private unnamed_addr addrspace(1) constant [16 x i8] c"../prefetch.hpp\00", section "llvm.metadata"
10+
@.str.2 = private unnamed_addr addrspace(1) constant [19 x i8] c"sycl-prefetch-hint\00", section "llvm.metadata"
11+
@.str.3 = private unnamed_addr addrspace(1) constant [2 x i8] c"0\00", section "llvm.metadata"
12+
@.args = private unnamed_addr addrspace(1) constant { ptr addrspace(1), ptr addrspace(1) } { ptr addrspace(1) @.str.2, ptr addrspace(1) @.str.3 }, section "llvm.metadata"
13+
@.str.4 = private unnamed_addr addrspace(1) constant [2 x i8] c"1\00", section "llvm.metadata"
14+
@.args.5 = private unnamed_addr addrspace(1) constant { ptr addrspace(1), ptr addrspace(1) } { ptr addrspace(1) @.str.2, ptr addrspace(1) @.str.4 }, section "llvm.metadata"
15+
@.str.6 = private unnamed_addr addrspace(1) constant [22 x i8] c"sycl-prefetch-hint-nt\00", section "llvm.metadata"
16+
@.str.7 = private unnamed_addr addrspace(1) constant [2 x i8] c"2\00", section "llvm.metadata"
17+
@.args.8 = private unnamed_addr addrspace(1) constant { ptr addrspace(1), ptr addrspace(1) } { ptr addrspace(1) @.str.6, ptr addrspace(1) @.str.7 }, section "llvm.metadata"
18+
19+
; CHECK: @[[NewAnnotStr1:.*]] = private unnamed_addr addrspace(1) constant [13 x i8] c"{6442:\220,1\22}\00"
20+
; CHECK: @[[NewAnnotStr2:.*]] = private unnamed_addr addrspace(1) constant [13 x i8] c"{6442:\221,1\22}\00"
21+
; CHECK: @[[NewAnnotStr3:.*]] = private unnamed_addr addrspace(1) constant [13 x i8] c"{6442:\222,3\22}\00"
22+
23+
; Function Attrs: convergent norecurse nounwind
24+
define weak_odr dso_local spir_kernel void @_ZTSZ4mainEUlvE_(ptr addrspace(1) noundef align 1 %_arg_dataPtr) local_unnamed_addr comdat !srcloc !5 !kernel_arg_buffer_location !6 !sycl_fixed_targets !7 !sycl_kernel_omit_args !8 {
25+
entry:
26+
%0 = addrspacecast ptr addrspace(1) %_arg_dataPtr to ptr addrspace(4)
27+
%call.i.i.i.i = tail call spir_func noundef ptr addrspace(1) @_Z41__spirv_GenericCastToPtrExplicit_ToGlobalPvi(ptr addrspace(4) noundef %0, i32 noundef 5)
28+
%1 = tail call ptr addrspace(1) @llvm.ptr.annotation.p1.p1(ptr addrspace(1) %call.i.i.i.i, ptr addrspace(1) @.str, ptr addrspace(1) @.str.1, i32 76, ptr addrspace(1) @.args)
29+
; CHECK %{{.*}} = call ptr addrspace(1) @llvm.ptr.annotation.p1.p1(ptr {{.*}}, ptr addrspace(1) @[[NewAnnotStr1]]{{.*}}
30+
tail call spir_func void @_Z20__spirv_ocl_prefetchPU3AS1Kcm(ptr addrspace(1) noundef %1, i64 noundef 1)
31+
%arrayidx3.i = getelementptr inbounds i8, ptr addrspace(4) %0, i64 1
32+
%call.i.i.i13.i = tail call spir_func noundef ptr addrspace(1) @_Z41__spirv_GenericCastToPtrExplicit_ToGlobalPvi(ptr addrspace(4) noundef %arrayidx3.i, i32 noundef 5)
33+
%2 = tail call ptr addrspace(1) @llvm.ptr.annotation.p1.p1(ptr addrspace(1) %call.i.i.i13.i, ptr addrspace(1) @.str, ptr addrspace(1) @.str.1, i32 80, ptr addrspace(1) @.args.5)
34+
; CHECK %{{.*}} = call ptr addrspace(1) @llvm.ptr.annotation.p1.p1(ptr {{.*}}, ptr addrspace(1) @[[NewAnnotStr2]]{{.*}}
35+
tail call spir_func void @_Z20__spirv_ocl_prefetchPU3AS1Kcm(ptr addrspace(1) noundef %2, i64 noundef 1)
36+
%arrayidx7.i = getelementptr inbounds i8, ptr addrspace(4) %0, i64 2
37+
%call.i.i.i16.i = tail call spir_func noundef ptr addrspace(1) @_Z41__spirv_GenericCastToPtrExplicit_ToGlobalPvi(ptr addrspace(4) noundef %arrayidx7.i, i32 noundef 5)
38+
%3 = tail call ptr addrspace(1) @llvm.ptr.annotation.p1.p1(ptr addrspace(1) %call.i.i.i16.i, ptr addrspace(1) @.str, ptr addrspace(1) @.str.1, i32 80, ptr addrspace(1) @.args.8)
39+
; CHECK %{{.*}} = call ptr addrspace(1) @llvm.ptr.annotation.p1.p1(ptr {{.*}}, ptr addrspace(1) @[[NewAnnotStr3]]{{.*}}
40+
tail call spir_func void @_Z20__spirv_ocl_prefetchPU3AS1Kcm(ptr addrspace(1) noundef %3, i64 noundef 2)
41+
ret void
42+
}
43+
44+
; Function Attrs: mustprogress nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: readwrite)
45+
declare ptr addrspace(1) @llvm.ptr.annotation.p1.p1(ptr addrspace(1), ptr addrspace(1), ptr addrspace(1), i32, ptr addrspace(1))
46+
47+
; Function Attrs: convergent nounwind
48+
declare dso_local spir_func void @_Z20__spirv_ocl_prefetchPU3AS1Kcm(ptr addrspace(1) noundef, i64 noundef) local_unnamed_addr
49+
50+
; Function Attrs: convergent mustprogress nofree nounwind willreturn memory(none)
51+
declare dso_local spir_func noundef ptr addrspace(1) @_Z41__spirv_GenericCastToPtrExplicit_ToGlobalPvi(ptr addrspace(4) noundef, i32 noundef) local_unnamed_addr
52+
53+
!llvm.module.flags = !{!0, !1}
54+
!opencl.spir.version = !{!2}
55+
!spirv.Source = !{!3}
56+
!llvm.ident = !{!4}
57+
58+
!0 = !{i32 1, !"wchar_size", i32 4}
59+
!1 = !{i32 7, !"frame-pointer", i32 2}
60+
!2 = !{i32 1, i32 2}
61+
!3 = !{i32 4, i32 100000}
62+
!4 = !{!"clang version 18.0.0"}
63+
!5 = !{i32 1522}
64+
!6 = !{i32 -1}
65+
!7 = !{}
66+
!8 = !{i1 false}

‎llvm/test/tools/sycl-post-link/spec-constants/nested-struct.ll

Lines changed: 149 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,149 @@
1+
; This test demonstrates that multiple padding elements can be
2+
; inserted in the spec constant metadata
3+
4+
; RUN: sycl-post-link --spec-const=native -S %s -o %t.table
5+
; RUN: FileCheck %s -input-file=%t_0.ll
6+
7+
; CHECK: %[[#SCV1:]] = call i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID1:]], i8 120)
8+
; CHECK: %[[#SCV2:]] = call i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID2:]], i8 121)
9+
; CHECK: %[[#SCV3:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#SCID3:]], i32 122)
10+
; CHECK: %[[#SCV4:]] = call i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID4:]], i8 97)
11+
; CHECK: %[[#SCV5:]] = call %struct.anon @_Z29__spirv_SpecConstantCompositeaia_Rstruct.anon(i8 %[[#SCV2:]], i32 %[[#SCV3:]], i8 %[[#SCV4:]])
12+
; CHECK: %[[#SCV6:]] = call i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID5:]], i8 98)
13+
; CHECK: call %struct.user_defined_type3 @_Z29__spirv_SpecConstantCompositeastruct.anona_Rstruct.user_defined_type3(i8 %[[#SCV1:]], %struct.anon %[[#SCV5:]], i8 %[[#SCV6:]])
14+
15+
; CHECK: !sycl.specialization-constants = !{![[#SC:]]}
16+
; CHECK: ![[#SC]] = !{!"uid0a28d8a0a23067ab____ZL8spec_id3",
17+
; CHECK-SAME: i32 [[#SCID1:]], i32 0, i32 1,
18+
; CHECK-SAME: i32 [[#SCID2:]], i32 4, i32 1,
19+
; CHECK-SAME: i32 [[#SCID3:]], i32 8, i32 4,
20+
; CHECK-SAME: i32 [[#SCID4:]], i32 12, i32 1,
21+
; CHECK-SAME: i32 -1, i32 13, i32 3,
22+
; CHECK-SAME: i32 [[#SCID5:]], i32 16, i32 1,
23+
; CHECK-SAME: i32 -1, i32 17, i32 3}
24+
25+
; ModuleID = '..\sycl\test-e2e\SpecConstants\2020\nested-non-packed-struct.cpp'
26+
source_filename = "..\\sycl\\test-e2e\\SpecConstants\\2020\\nested-non-packed-struct.cpp"
27+
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
28+
target triple = "spir64-unknown-unknown"
29+
30+
%"class.sycl::_V1::specialization_id.24" = type { %struct.user_defined_type3 }
31+
%struct.user_defined_type3 = type { i8, %struct.anon, i8 }
32+
%struct.anon = type { i8, i32, i8 }
33+
%"class.sycl::_V1::id" = type { %"class.sycl::_V1::detail::array" }
34+
%"class.sycl::_V1::detail::array" = type { [1 x i64] }
35+
36+
@__usid_str.2 = private unnamed_addr constant [35 x i8] c"uid0a28d8a0a23067ab____ZL8spec_id3\00", align 1
37+
@_ZL8spec_id3 = internal addrspace(1) constant %"class.sycl::_V1::specialization_id.24" { %struct.user_defined_type3 { i8 120, %struct.anon { i8 121, i32 122, i8 97 }, i8 98 } }, align 4
38+
39+
; Function Attrs: convergent norecurse nounwind
40+
define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE2_clES2_EUlNS0_14kernel_handlerEE_(ptr addrspace(1) noundef align 4 %_arg_acc, ptr noundef byval(%"class.sycl::_V1::id") align 8 %_arg_acc3) local_unnamed_addr #0 !srcloc !78 !kernel_arg_buffer_location !63 !kernel_arg_runtime_aligned !64 !kernel_arg_exclusive_ptr !64 !sycl_fixed_targets !65 !sycl_kernel_omit_args !66 {
41+
entry:
42+
%ref.tmp.i = alloca %struct.user_defined_type3, align 4
43+
%0 = load i64, ptr %_arg_acc3, align 8
44+
%add.ptr.i = getelementptr inbounds %struct.user_defined_type3, ptr addrspace(1) %_arg_acc, i64 %0
45+
%ref.tmp.ascast.i = addrspacecast ptr %ref.tmp.i to ptr addrspace(4)
46+
call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI18user_defined_type3ET_PKcPKvS5_(ptr addrspace(4) sret(%struct.user_defined_type3) align 4 %ref.tmp.ascast.i, ptr addrspace(4) noundef addrspacecast (ptr @__usid_str.2 to ptr addrspace(4)), ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @_ZL8spec_id3 to ptr addrspace(4)), ptr addrspace(4) noundef null) #5
47+
call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 4 %add.ptr.i, ptr align 4 %ref.tmp.i, i64 20, i1 false), !tbaa.struct !79
48+
ret void
49+
}
50+
51+
; Function Attrs: convergent nounwind
52+
declare dso_local spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI18user_defined_type3ET_PKcPKvS5_(ptr addrspace(4) sret(%struct.user_defined_type3) align 4, ptr addrspace(4) noundef, ptr addrspace(4) noundef, ptr addrspace(4) noundef) local_unnamed_addr #2
53+
54+
; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: readwrite)
55+
declare void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) noalias nocapture writeonly, ptr noalias nocapture readonly, i64, i1 immarg) #3
56+
57+
attributes #0 = { convergent norecurse nounwind "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="..\\sycl\\test-e2e\\SpecConstants\\2020\\nested-non-packed-struct.cpp" "sycl-optlevel"="2" "sycl-single-task" "uniform-work-group-size"="true" }
58+
attributes #2 = { convergent nounwind "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
59+
attributes #3 = { nocallback nofree nounwind willreturn memory(argmem: readwrite) }
60+
attributes #4 = { nounwind }
61+
attributes #5 = { convergent nounwind }
62+
63+
!llvm.dependent-libraries = !{!0}
64+
!llvm.module.flags = !{!1, !2}
65+
!opencl.spir.version = !{!3}
66+
!spirv.Source = !{!4}
67+
!sycl_aspects = !{!5, !6, !7, !8, !9, !10, !11, !12, !13, !14, !15, !16, !17, !18, !19, !20, !21, !22, !23, !24, !25, !26, !27, !28, !29, !30, !31, !32, !33, !34, !35, !36, !37, !38, !39, !40, !41, !42, !43, !44, !45, !46, !47, !48, !49, !50, !51, !52, !53, !54, !55, !56, !57, !58, !59, !60}
68+
!llvm.ident = !{!61}
69+
70+
!0 = !{!"libcpmt"}
71+
!1 = !{i32 1, !"wchar_size", i32 2}
72+
!2 = !{i32 7, !"frame-pointer", i32 2}
73+
!3 = !{i32 1, i32 2}
74+
!4 = !{i32 4, i32 100000}
75+
!5 = !{!"cpu", i32 1}
76+
!6 = !{!"gpu", i32 2}
77+
!7 = !{!"accelerator", i32 3}
78+
!8 = !{!"custom", i32 4}
79+
!9 = !{!"fp16", i32 5}
80+
!10 = !{!"fp64", i32 6}
81+
!11 = !{!"image", i32 9}
82+
!12 = !{!"online_compiler", i32 10}
83+
!13 = !{!"online_linker", i32 11}
84+
!14 = !{!"queue_profiling", i32 12}
85+
!15 = !{!"usm_device_allocations", i32 13}
86+
!16 = !{!"usm_host_allocations", i32 14}
87+
!17 = !{!"usm_shared_allocations", i32 15}
88+
!18 = !{!"usm_system_allocations", i32 17}
89+
!19 = !{!"ext_intel_pci_address", i32 18}
90+
!20 = !{!"ext_intel_gpu_eu_count", i32 19}
91+
!21 = !{!"ext_intel_gpu_eu_simd_width", i32 20}
92+
!22 = !{!"ext_intel_gpu_slices", i32 21}
93+
!23 = !{!"ext_intel_gpu_subslices_per_slice", i32 22}
94+
!24 = !{!"ext_intel_gpu_eu_count_per_subslice", i32 23}
95+
!25 = !{!"ext_intel_max_mem_bandwidth", i32 24}
96+
!26 = !{!"ext_intel_mem_channel", i32 25}
97+
!27 = !{!"usm_atomic_host_allocations", i32 26}
98+
!28 = !{!"usm_atomic_shared_allocations", i32 27}
99+
!29 = !{!"atomic64", i32 28}
100+
!30 = !{!"ext_intel_device_info_uuid", i32 29}
101+
!31 = !{!"ext_oneapi_srgb", i32 30}
102+
!32 = !{!"ext_oneapi_native_assert", i32 31}
103+
!33 = !{!"host_debuggable", i32 32}
104+
!34 = !{!"ext_intel_gpu_hw_threads_per_eu", i32 33}
105+
!35 = !{!"ext_oneapi_cuda_async_barrier", i32 34}
106+
!36 = !{!"ext_oneapi_bfloat16_math_functions", i32 35}
107+
!37 = !{!"ext_intel_free_memory", i32 36}
108+
!38 = !{!"ext_intel_device_id", i32 37}
109+
!39 = !{!"ext_intel_memory_clock_rate", i32 38}
110+
!40 = !{!"ext_intel_memory_bus_width", i32 39}
111+
!41 = !{!"emulated", i32 40}
112+
!42 = !{!"ext_intel_legacy_image", i32 41}
113+
!43 = !{!"ext_oneapi_bindless_images", i32 42}
114+
!44 = !{!"ext_oneapi_bindless_images_shared_usm", i32 43}
115+
!45 = !{!"ext_oneapi_bindless_images_1d_usm", i32 44}
116+
!46 = !{!"ext_oneapi_bindless_images_2d_usm", i32 45}
117+
!47 = !{!"ext_oneapi_interop_memory_import", i32 46}
118+
!48 = !{!"ext_oneapi_interop_memory_export", i32 47}
119+
!49 = !{!"ext_oneapi_interop_semaphore_import", i32 48}
120+
!50 = !{!"ext_oneapi_interop_semaphore_export", i32 49}
121+
!51 = !{!"ext_oneapi_mipmap", i32 50}
122+
!52 = !{!"ext_oneapi_mipmap_anisotropy", i32 51}
123+
!53 = !{!"ext_oneapi_mipmap_level_reference", i32 52}
124+
!54 = !{!"ext_intel_esimd", i32 53}
125+
!55 = !{!"ext_oneapi_non_uniform_groups", i32 54}
126+
!56 = !{!"int64_base_atomics", i32 7}
127+
!57 = !{!"int64_extended_atomics", i32 8}
128+
!58 = !{!"usm_system_allocator", i32 17}
129+
!59 = !{!"usm_restricted_shared_allocations", i32 16}
130+
!60 = !{!"host", i32 0}
131+
!61 = !{!"clang version 18.0.0 (https://github.com/intel/llvm.git c92b6b0c266b6a0d5bca1d61a63f06e2bce37904)"}
132+
!62 = !{i32 2091}
133+
!63 = !{i32 -1, i32 -1}
134+
!64 = !{i1 true, i1 false}
135+
!65 = !{}
136+
!66 = !{i1 false, i1 true, i1 true, i1 false, i1 true}
137+
!67 = !{i64 0, i64 4, !68, i64 4, i64 1, !72, i64 8, i64 4, !73}
138+
!68 = !{!69, !69, i64 0}
139+
!69 = !{!"float", !70, i64 0}
140+
!70 = !{!"omnipotent char", !71, i64 0}
141+
!71 = !{!"Simple C++ TBAA"}
142+
!72 = !{!70, !70, i64 0}
143+
!73 = !{!74, !74, i64 0}
144+
!74 = !{!"int", !70, i64 0}
145+
!75 = !{i32 2451}
146+
!76 = !{i64 0, i64 4, !68, i64 32, i64 4, !68, i64 36, i64 1, !72, i64 40, i64 4, !73, i64 64, i64 4, !73, i64 68, i64 1, !72}
147+
!77 = !{i32 2882}
148+
!78 = !{i32 3230}
149+
!79 = !{i64 0, i64 1, !72, i64 4, i64 1, !72, i64 8, i64 4, !73, i64 12, i64 1, !72, i64 16, i64 1, !72}

0 commit comments

Comments
 (0)