Skip to content

Commit

Permalink
[NFC][OpaquePtrs] Update SYCLLowerIR tests to use opaque pointers (#1…
Browse files Browse the repository at this point in the history 
…0687)

Update SYCLLowerIR tests to use opaque pointers.

---------

Signed-off-by: Lu, John <john.lu@intel.com>
  • Loading branch information
@LU-JOHN
LU-JOHN committed Aug 17, 2023
1 parent d82c2a1 commit cdd4dd3
Show file tree
Hide file tree
Showing 33 changed files with 528 additions and 575 deletions.
24 changes: 12 additions & 12 deletions llvm/test/SYCLLowerIR/CompileTimePropertiesPass/device-globals/test_global_variable.ll
View file
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
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
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

0 comments on commit cdd4dd3

Please sign in to comment.