Skip to content

[NFC][OpaquePtrs] Update SYCLLowerIR tests to use opaque pointers #10687

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 4 commits into from
Aug 17, 2023
Merged

[NFC][OpaquePtrs] Update SYCLLowerIR tests to use opaque pointers #10687

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
509fe11
[NFC][OpaquePtrs] Update SYCLLowerIR tests to use opaque pointers
Aug 3, 2023
887ab89
Revert opaquify changes for tests that run ESIMDLowerVecArg pass. Th…
Aug 8, 2023
e9841bc
Opaquify removed GEP insn. Remove associated llvm.dbg.value call.
Aug 8, 2023
e4d1d24
Disable SYCL alignment properties testcase temporarily. Alignment no…
Aug 8, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
24 changes: 12 additions & 12 deletions llvm/test/SYCLLowerIR/CompileTimePropertiesPass/device-globals/test_global_variable.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,7 +9,7 @@ source_filename = "test_global_variable.cpp"
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"
target triple = "spir64-unknown-unknown"

%"class.cl::sycl::ext::oneapi::device_global.0" = type { i32 addrspace(4)* }
%"class.cl::sycl::ext::oneapi::device_global.0" = type { ptr addrspace(4) }
%"class.cl::sycl::ext::oneapi::device_global.1" = type { i8 }
%class.anon.0 = type { i8 }

Expand All @@ -24,24 +24,24 @@ target triple = "spir64-unknown-unknown"
@_ZL7no_dg_int1 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8 #6
; CHECK-IR: @_ZL7no_dg_int1 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8, !spirv.Decorations ![[#MN12:]]

define internal spir_func void @_ZZ4mainENKUlvE_clEv(%class.anon.0 addrspace(4)* align 1 dereferenceable_or_null(1) %this) #4 align 2 {
define internal spir_func void @_ZZ4mainENKUlvE_clEv(ptr addrspace(4) align 1 dereferenceable_or_null(1) %this) #4 align 2 {
entry:
%this.addr = alloca %class.anon.0 addrspace(4)*, align 8
%this.addr.ascast = addrspacecast %class.anon.0 addrspace(4)** %this.addr to %class.anon.0 addrspace(4)* addrspace(4)*
store %class.anon.0 addrspace(4)* %this, %class.anon.0 addrspace(4)* addrspace(4)* %this.addr.ascast, align 8
%this1 = load %class.anon.0 addrspace(4)*, %class.anon.0 addrspace(4)* addrspace(4)* %this.addr.ascast, align 8
%call1 = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(4)* align 8 dereferenceable_or_null(8) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(1)* @_ZL7dg_int1 to %"class.cl::sycl::ext::oneapi::device_global.0" addrspace(4)*)) #5
%call2 = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(4)* align 8 dereferenceable_or_null(8) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(1)* @_ZL7dg_int2 to %"class.cl::sycl::ext::oneapi::device_global.0" addrspace(4)*)) #5
%call3 = call spir_func align 1 dereferenceable(1) i8 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIbJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global.1" addrspace(4)* align 1 dereferenceable_or_null(1) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.1" addrspace(1)* @_ZL8dg_bool3 to %"class.cl::sycl::ext::oneapi::device_global.1" addrspace(4)*)) #5
%call4 = call spir_func align 1 dereferenceable(1) i8 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIbJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global.1" addrspace(4)* align 1 dereferenceable_or_null(1) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.1" addrspace(1)* @_ZL8dg_bool4 to %"class.cl::sycl::ext::oneapi::device_global.1" addrspace(4)*)) #5
%this.addr = alloca ptr addrspace(4), align 8
%this.addr.ascast = addrspacecast ptr %this.addr to ptr addrspace(4)
store ptr addrspace(4) %this, ptr addrspace(4) %this.addr.ascast, align 8
%this1 = load ptr addrspace(4), ptr addrspace(4) %this.addr.ascast, align 8
%call1 = call spir_func align 4 dereferenceable(4) ptr addrspace(4) @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(ptr addrspace(4) align 8 dereferenceable_or_null(8) addrspacecast (ptr addrspace(1) @_ZL7dg_int1 to ptr addrspace(4))) #5
%call2 = call spir_func align 4 dereferenceable(4) ptr addrspace(4) @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(ptr addrspace(4) align 8 dereferenceable_or_null(8) addrspacecast (ptr addrspace(1) @_ZL7dg_int2 to ptr addrspace(4))) #5
%call3 = call spir_func align 1 dereferenceable(1) ptr addrspace(4) @_ZNK2cl4sycl3ext6oneapi13device_globalIbJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(ptr addrspace(4) align 1 dereferenceable_or_null(1) addrspacecast (ptr addrspace(1) @_ZL8dg_bool3 to ptr addrspace(4))) #5
%call4 = call spir_func align 1 dereferenceable(1) ptr addrspace(4) @_ZNK2cl4sycl3ext6oneapi13device_globalIbJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(ptr addrspace(4) align 1 dereferenceable_or_null(1) addrspacecast (ptr addrspace(1) @_ZL8dg_bool4 to ptr addrspace(4))) #5
ret void
}

; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
declare spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(4)* align 8 dereferenceable_or_null(8) %this) #4 align 2
declare spir_func align 4 dereferenceable(4) ptr addrspace(4) @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(ptr addrspace(4) align 8 dereferenceable_or_null(8) %this) #4 align 2

; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
declare spir_func align 1 dereferenceable(1) i8 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIbJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global.1" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #4 align 2
declare spir_func align 1 dereferenceable(1) ptr addrspace(4) @_ZNK2cl4sycl3ext6oneapi13device_globalIbJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(ptr addrspace(4) align 1 dereferenceable_or_null(1) %this) #4 align 2

attributes #0 = { "sycl-unique-id"="6da74a122db9f35d____ZL7dg_int1" "sycl-device-image-scope"="false" "sycl-host-access"="1" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-device-global-size"="4" }
attributes #1 = { "sycl-unique-id"="7da74a1187b9f35d____ZL7dg_int2" "sycl-implement-in-csr"="false" "sycl-init-mode"="1" "sycl-device-global-size"="4" }
Expand Down
17 changes: 6 additions & 11 deletions llvm/test/SYCLLowerIR/CompileTimePropertiesPass/sycl-latency-control.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -26,19 +26,14 @@ entry:
%1 = alloca ptr , align 8
store ptr %0, ptr %1, align 8
%2 = load ptr , ptr %1, align 8
%3 = getelementptr inbounds %struct.__spirv_Something, ptr %2, i32 0, i32 0
%4 = bitcast ptr %3 to ptr
%5 = call ptr @llvm.ptr.annotation.p0.p0(ptr %4, ptr @.str, ptr @.str.1, i32 5, ptr @.args)
%3 = call ptr @llvm.ptr.annotation.p0.p0(ptr %2, ptr @.str, ptr @.str.1, i32 5, ptr @.args)
; CHECK: %{{.*}} = call ptr @llvm.ptr.annotation.p0.p0(ptr %[[#]], ptr @[[NewAnnotStr1]], ptr @.str.1, i32 5, ptr null)
%6 = bitcast ptr %5 to ptr
%7 = load i32, ptr %6, align 8
%8 = load ptr , ptr %1, align 8
%9 = getelementptr inbounds %struct.__spirv_Something, ptr %8, i32 0, i32 1
%10 = bitcast ptr %9 to ptr
%11 = call ptr @llvm.ptr.annotation.p0.p0(ptr %10, ptr @.str, ptr @.str.1, i32 5, ptr @.args.9)
%4 = load i32, ptr %3, align 8
%5 = load ptr , ptr %1, align 8
%6 = getelementptr inbounds %struct.__spirv_Something, ptr %5, i32 0, i32 1
%7 = call ptr @llvm.ptr.annotation.p0.p0(ptr %6, ptr @.str, ptr @.str.1, i32 5, ptr @.args.9)
; CHECK: %{{.*}} = call ptr @llvm.ptr.annotation.p0.p0(ptr %[[#]], ptr @[[NewAnnotStr2]], ptr @.str.1, i32 5, ptr null)
%12 = bitcast ptr %11 to ptr
%13 = load i32, ptr %12, align 8
%8 = load i32, ptr %7, align 8
ret void
}

Expand Down
89 changes: 42 additions & 47 deletions llvm/test/SYCLLowerIR/CompileTimePropertiesPass/sycl-properties-alignment-loadstore.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,10 +3,13 @@
;
; Tests the translation of "sycl-alignment" to alignment attributes on load/store

; FIXME: Alignment properties not preserved after testcase was opaquified
; REQUIRES: TEMPORARY_DISABLED

target triple = "spir64_fpga-unknown-unknown"

%struct.MyIP = type { %class.ann_ptr }
%class.ann_ptr = type { i32 addrspace(4)* }
%class.ann_ptr = type { ptr addrspace(4) }

$_ZN7ann_refIiEC2EPi = comdat any
$_ZN7ann_refIiEcvRiEv = comdat any
Expand All @@ -16,73 +19,65 @@ $_ZN7ann_refIiEC2EPi1= comdat any
@.str.1 = private unnamed_addr addrspace(1) constant [9 x i8] c"main.cpp\00", section "llvm.metadata"
@.str.2 = private unnamed_addr addrspace(1) constant [15 x i8] c"sycl-alignment\00", section "llvm.metadata"
@.str.3 = private unnamed_addr addrspace(1) constant [3 x i8] c"64\00", section "llvm.metadata"
@.args = private unnamed_addr addrspace(1) constant { [15 x i8] addrspace(1)*, [3 x i8] addrspace(1)* } { [15 x i8] addrspace(1)* @.str.2, [3 x i8] addrspace(1)* @.str.3 }, section "llvm.met
@.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.met
adata"

; Function Attrs: nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: readwrite)
declare i8 addrspace(4)* @llvm.ptr.annotation.p4i8.p1i8(i8 addrspace(4)*, i8 addrspace(1)*, i8 addrspace(1)*, i32, i8 addrspace(1)*) #5
declare ptr addrspace(4) @llvm.ptr.annotation.p4.p1(ptr addrspace(4), ptr addrspace(1), ptr addrspace(1), i32, ptr addrspace(1)) #5

define weak_odr dso_local spir_kernel void @_MyIP(i32 addrspace(1)* noundef "sycl-alignment"="64" %_arg_a) {
define weak_odr dso_local spir_kernel void @_MyIP(ptr addrspace(1) noundef "sycl-alignment"="64" %_arg_a) {
; CHECK: define{{.*}}@_MyIP{{.*}}align 64{{.*}} {
ret void
}

; Function Attrs: convergent mustprogress norecurse nounwind
define linkonce_odr dso_local spir_func noundef align 4 dereferenceable(4) i32 addrspace(4)* @_ZN7ann_refIiEcvRiEv(%class.ann_ptr addrspace(4)* noundef align 8 dereferenceable_or_null(8) %this) #3 comdat align 2 {
define linkonce_odr dso_local spir_func noundef align 4 dereferenceable(4) ptr addrspace(4) @_ZN7ann_refIiEcvRiEv(ptr addrspace(4) noundef align 8 dereferenceable_or_null(8) %this) #3 comdat align 2 {
entry:
%retval = alloca i32 addrspace(4)*, align 8
%this.addr = alloca %class.ann_ptr addrspace(4)*, align 8
%retval.ascast = addrspacecast i32 addrspace(4)** %retval to i32 addrspace(4)* addrspace(4)*
%this.addr.ascast = addrspacecast %class.ann_ptr addrspace(4)** %this.addr to %class.ann_ptr addrspace(4)* addrspace(4)*
store %class.ann_ptr addrspace(4)* %this, %class.ann_ptr addrspace(4)* addrspace(4)* %this.addr.ascast, align 8
%this1 = load %class.ann_ptr addrspace(4)*, %class.ann_ptr addrspace(4)* addrspace(4)* %this.addr.ascast, align 8
%p = getelementptr inbounds %class.ann_ptr, %class.ann_ptr addrspace(4)* %this1, i32 0, i32 0
%0 = load i32 addrspace(4)*, i32 addrspace(4)* addrspace(4)* %p, align 8
%1 = bitcast i32 addrspace(4)* %0 to i8 addrspace(4)*
%2 = call i8 addrspace(4)* @llvm.ptr.annotation.p4i8.p1i8(i8 addrspace(4)* %1, i8 addrspace(1)* getelementptr inbounds ([16 x i8], [16 x i8] addrspace(1)* @.str, i32 0, i32 0), i8 addrspace(1)* getelementptr inbounds ([9 x i8], [9 x i8] addrspace(1)* @.str.1, i32 0, i32 0), i32 22, i8 addrspace(1)* bitcast ({ [15 x i8] addrspace(1)*, [3 x i8] addrspace(1)* } addrspace(1)* @.args to i8 addrspace(1)*))
%3 = bitcast i8 addrspace(4)* %2 to i32 addrspace(4)*
%4 = load i32, i32 addrspace(4)* %3, align 8
%retval = alloca ptr addrspace(4), align 8
%this.addr = alloca ptr addrspace(4), align 8
%retval.ascast = addrspacecast ptr %retval to ptr addrspace(4)
%this.addr.ascast = addrspacecast ptr %this.addr to ptr addrspace(4)
store ptr addrspace(4) %this, ptr addrspace(4) %this.addr.ascast, align 8
%this1 = load ptr addrspace(4), ptr addrspace(4) %this.addr.ascast, align 8
%0 = load ptr addrspace(4), ptr addrspace(4) %this1, align 8
%1 = call ptr addrspace(4) @llvm.ptr.annotation.p4.p1(ptr addrspace(4) %0, ptr addrspace(1) @.str, ptr addrspace(1) @.str.1, i32 22, ptr addrspace(1) @.args)
%2 = load i32, ptr addrspace(4) %1, align 8
; CHECK: load {{.*}}, align 64
ret i32 addrspace(4)* %3
ret ptr addrspace(4) %1
}

; Function Attrs: convergent norecurse nounwind
define linkonce_odr dso_local spir_func void @_ZN7ann_refIiEC2EPi(%class.ann_ptr addrspace(4)* noundef align 8 dereferenceable_or_null(8) %this, i32 addrspace(4)* noundef %ptr) unnamed_addr #2 comdat align 2 {
define linkonce_odr dso_local spir_func void @_ZN7ann_refIiEC2EPi(ptr addrspace(4) noundef align 8 dereferenceable_or_null(8) %this, ptr addrspace(4) noundef %ptr) unnamed_addr #2 comdat align 2 {
entry:
%this.addr = alloca %class.ann_ptr addrspace(4)*, align 8
%ptr.addr = alloca i32 addrspace(4)*, align 8
%this.addr.ascast = addrspacecast %class.ann_ptr addrspace(4)** %this.addr to %class.ann_ptr addrspace(4)* addrspace(4)*
%ptr.addr.ascast = addrspacecast i32 addrspace(4)** %ptr.addr to i32 addrspace(4)* addrspace(4)*
store %class.ann_ptr addrspace(4)* %this, %class.ann_ptr addrspace(4)* addrspace(4)* %this.addr.ascast, align 8
store i32 addrspace(4)* %ptr, i32 addrspace(4)* addrspace(4)* %ptr.addr.ascast, align 8
%this1 = load %class.ann_ptr addrspace(4)*, %class.ann_ptr addrspace(4)* addrspace(4)* %this.addr.ascast, align 8
%p = getelementptr inbounds %class.ann_ptr, %class.ann_ptr addrspace(4)* %this1, i32 0, i32 0
%0 = load i32 addrspace(4)*, i32 addrspace(4)* addrspace(4)* %p, align 8
%1 = bitcast i32 addrspace(4)* %0 to i8 addrspace(4)*
%2 = call i8 addrspace(4)* @llvm.ptr.annotation.p4i8.p1i8(i8 addrspace(4)* %1, i8 addrspace(1)* getelementptr inbounds ([16 x i8], [16 x i8] addrspace(1)* @.str, i32 0, i32 0), i8 addrspace(1)* getelementptr inbounds ([9 x i8], [9 x i8] addrspace(1)* @.str.1, i32 0, i32 0), i32 22, i8 addrspace(1)* bitcast ({ [15 x i8] addrspace(1)*, [3 x i8] addrspace(1)* } addrspace(1)* @.args to i8 addrspace(1)*))
%3 = bitcast i8 addrspace(4)* %2 to i32 addrspace(4)*
store i32 5, i32 addrspace(4)* %3, align 8
%this.addr = alloca ptr addrspace(4), align 8
%ptr.addr = alloca ptr addrspace(4), align 8
%this.addr.ascast = addrspacecast ptr %this.addr to ptr addrspace(4)
%ptr.addr.ascast = addrspacecast ptr %ptr.addr to ptr addrspace(4)
store ptr addrspace(4) %this, ptr addrspace(4) %this.addr.ascast, align 8
store ptr addrspace(4) %ptr, ptr addrspace(4) %ptr.addr.ascast, align 8
%this1 = load ptr addrspace(4), ptr addrspace(4) %this.addr.ascast, align 8
%0 = load ptr addrspace(4), ptr addrspace(4) %this1, align 8
%1 = call ptr addrspace(4) @llvm.ptr.annotation.p4.p1(ptr addrspace(4) %0, ptr addrspace(1) @.str, ptr addrspace(1) @.str.1, i32 22, ptr addrspace(1) @.args)
store i32 5, ptr addrspace(4) %1, align 8
; CHECK: store {{.*}}, align 64
ret void
}

; Function Attrs: convergent norecurse nounwind
define linkonce_odr dso_local spir_func void @_ZN7ann_refIiEC2EPi1(%class.ann_ptr addrspace(4)* noundef align 8 dereferenceable_or_null(8) %this, i32 addrspace(4)* noundef %ptr, i8 addrspace(4)* %h) comdat align 2 {
define linkonce_odr dso_local spir_func void @_ZN7ann_refIiEC2EPi1(ptr addrspace(4) noundef align 8 dereferenceable_or_null(8) %this, ptr addrspace(4) noundef %ptr, ptr addrspace(4) %h) comdat align 2 {
entry:
%this.addr = alloca %class.ann_ptr addrspace(4)*, align 8
%ptr.addr = alloca i32 addrspace(4)*, align 8
%this.addr.ascast = addrspacecast %class.ann_ptr addrspace(4)** %this.addr to %class.ann_ptr addrspace(4)* addrspace(4)*
%ptr.addr.ascast = addrspacecast i32 addrspace(4)** %ptr.addr to i32 addrspace(4)* addrspace(4)*
store %class.ann_ptr addrspace(4)* %this, %class.ann_ptr addrspace(4)* addrspace(4)* %this.addr.ascast, align 8
store i32 addrspace(4)* %ptr, i32 addrspace(4)* addrspace(4)* %ptr.addr.ascast, align 8
%this1 = load %class.ann_ptr addrspace(4)*, %class.ann_ptr addrspace(4)* addrspace(4)* %this.addr.ascast, align 8
%p = getelementptr inbounds %class.ann_ptr, %class.ann_ptr addrspace(4)* %this1, i32 0, i32 0
%0 = load i32 addrspace(4)*, i32 addrspace(4)* addrspace(4)* %p, align 8
%1 = bitcast i32 addrspace(4)* %0 to i8 addrspace(4)*
%2 = call i8 addrspace(4)* @llvm.ptr.annotation.p4i8.p1i8(i8 addrspace(4)* %1, i8 addrspace(1)* getelementptr inbounds ([16 x i8], [16 x i8] addrspace(1)* @.str, i32 0, i32 0), i8 addrspace(1)* getelementptr inbounds ([9 x i8], [9 x i8] addrspace(1)* @.str.1, i32 0, i32 0), i32 22, i8 addrspace(1)* bitcast ({ [15 x i8] addrspace(1)*, [3 x i8] addrspace(1)* } addrspace(1)* @.args to i8 addrspace(1)*))
call void @llvm.memcpy.p4i8.p4i8.i32(i8 addrspace(4)* %2, i8 addrspace(4)* %h, i32 1, i1 false)
; CHECK: call void @llvm.memcpy.p4i8.p4i8.i32(i8 addrspace(4)* align 64 %1, i8 addrspace(4)* %h, i32 1, i1 false)
%this.addr = alloca ptr addrspace(4), align 8
%ptr.addr = alloca ptr addrspace(4), align 8
%this.addr.ascast = addrspacecast ptr %this.addr to ptr addrspace(4)
%ptr.addr.ascast = addrspacecast ptr %ptr.addr to ptr addrspace(4)
store ptr addrspace(4) %this, ptr addrspace(4) %this.addr.ascast, align 8
store ptr addrspace(4) %ptr, ptr addrspace(4) %ptr.addr.ascast, align 8
%this1 = load ptr addrspace(4), ptr addrspace(4) %this.addr.ascast, align 8
%0 = load ptr addrspace(4), ptr addrspace(4) %this1, align 8
%1 = call ptr addrspace(4) @llvm.ptr.annotation.p4.p1(ptr addrspace(4) %0, ptr addrspace(1) @.str, ptr addrspace(1) @.str.1, i32 22, ptr addrspace(1) @.args)
call void @llvm.memcpy.p4.p4.i32(ptr addrspace(4) %1, ptr addrspace(4) %h, i32 1, i1 false)
; CHECK: call void @llvm.memcpy.p4.p4.i32(ptr addrspace(4) align 64 %0, ptr addrspace(4) %h, i32 1, i1 false)
ret void
}

declare void @llvm.memcpy.p4i8.p4i8.i32(i8 addrspace(4)*, i8 addrspace(4)*, i32, i1)
declare void @llvm.memcpy.p4.p4.i32(ptr addrspace(4), ptr addrspace(4), i32, i1)
Loading