Skip to content

Commit

Permalink
Merge pull request #604 from hvdijk/gep-nuw-flag
Browse files Browse the repository at this point in the history 
[LLVM 20] Allow getelementptr nuw flag.
  • Loading branch information
@hvdijk
hvdijk authored Dec 6, 2024
2 parents 6bd190e + 6896e75 commit e44c154
Show file tree
Hide file tree
Showing 4 changed files with 33 additions and 33 deletions.
2 changes: 1 addition & 1 deletion modules/compiler/vecz/test/lit/llvm/constant_address.ll
View file
Original file line number Diff line number Diff line change
Expand Up @@ -54,5 +54,5 @@ attributes #1 = { nounwind readnone }
; CHECK-NEXT: entry:
; CHECK-NEXT: %gid = call i64 @__mux_get_global_id(i32 0)
; CHECK-NEXT: %conv = trunc i64 %gid to i32
; CHECK-NEXT: %arrayidx = getelementptr inbounds {{i32|i8}}, ptr addrspace(1) %out, i64 {{3|12}}
; CHECK-NEXT: %arrayidx = getelementptr inbounds {{(nuw )?}}{{i32|i8}}, ptr addrspace(1) %out, i64 {{3|12}}
; CHECK-NEXT: store i32 %conv, ptr addrspace(1) %arrayidx, align 4
2 changes: 1 addition & 1 deletion modules/compiler/vecz/test/lit/llvm/constant_address_with_uniform.ll
View file
Original file line number Diff line number Diff line change
Expand Up @@ -36,6 +36,6 @@ entry:
; CHECK: define spir_kernel void @__vecz_v4_test
; CHECK-NEXT: entry:
; CHECK-NEXT: %gid = call i32 @__mux_get_global_id(i32 0)
; CHECK-NEXT: %arrayidx = getelementptr inbounds {{i32|i8}}, ptr addrspace(1) %out, i32 {{3|12}}
; CHECK-NEXT: %arrayidx = getelementptr inbounds {{(nuw )?}}{{i32|i8}}, ptr addrspace(1) %out, i32 {{3|12}}
; CHECK: store i32 %gid, ptr addrspace(1) %arrayidx, align 4
; CHECK: store <4 x ptr addrspace(1)> %{{.+}}, ptr addrspace(1) %{{.+}}
60 changes: 30 additions & 30 deletions modules/compiler/vecz/test/lit/llvm/emit_memintrinsics.ll
View file
Original file line number Diff line number Diff line change
Expand Up @@ -98,45 +98,45 @@ declare i64 @__mux_get_local_id(i32)
; Check if the generated loads and stores are in place
; Check the stores for the first memset
; CHECK: store i64 %ms64val, ptr %sa
; CHECK: %[[V14:[0-9]+]] = getelementptr inbounds i8, ptr %sa, i64 8
; CHECK: %[[V14:[0-9]+]] = getelementptr inbounds {{(nuw )?}}i8, ptr %sa, i64 8
; CHECK: store i64 %ms64val, ptr %[[V14]]
; CHECK: %[[V15:[0-9]+]] = getelementptr inbounds i8, ptr %sa, i64 16
; CHECK: %[[V15:[0-9]+]] = getelementptr inbounds {{(nuw )?}}i8, ptr %sa, i64 16
; CHECK: store i64 %ms64val, ptr %[[V15]]
; CHECK: %[[V16:[0-9]+]] = getelementptr inbounds i8, ptr %sa, i64 24
; CHECK: %[[V16:[0-9]+]] = getelementptr inbounds {{(nuw )?}}i8, ptr %sa, i64 24
; CHECK: store i64 %ms64val, ptr %[[V16]]
; CHECK: %[[V17:[0-9]+]] = getelementptr inbounds i8, ptr %sa, i64 32
; CHECK: %[[V17:[0-9]+]] = getelementptr inbounds {{(nuw )?}}i8, ptr %sa, i64 32
; CHECK: store i64 %ms64val, ptr %[[V17]]
; CHECK: %[[V18:[0-9]+]] = getelementptr inbounds i8, ptr %sa, i64 40
; CHECK: %[[V18:[0-9]+]] = getelementptr inbounds {{(nuw )?}}i8, ptr %sa, i64 40
; CHECK: store i64 %ms64val, ptr %[[V18]]
; CHECK: %[[V19:[0-9]+]] = getelementptr inbounds i8, ptr %sa, i64 48
; CHECK: %[[V19:[0-9]+]] = getelementptr inbounds {{(nuw )?}}i8, ptr %sa, i64 48
; CHECK: store i64 %ms64val, ptr %[[V19]]
; CHECK: %[[V20:[0-9]+]] = getelementptr inbounds i8, ptr %sa, i64 56
; CHECK-EQ14: %[[V20:[0-9]+]] = getelementptr inbounds %struct.S2, %struct.S2* %sa, i64 0, i32 3, i64 8
; CHECK: %[[V21:[0-9]+]] = getelementptr inbounds i8, ptr %sa, i64 64
; CHECK: %[[V22:[0-9]+]] = getelementptr inbounds i8, ptr %sa, i64 72
; CHECK: %[[V20:[0-9]+]] = getelementptr inbounds {{(nuw )?}}i8, ptr %sa, i64 56
; CHECK-EQ14: %[[V20:[0-9]+]] = getelementptr inbounds {{(nuw )?}}%struct.S2, %struct.S2* %sa, i64 0, i32 3, i64 8
; CHECK: %[[V21:[0-9]+]] = getelementptr inbounds {{(nuw )?}}i8, ptr %sa, i64 64
; CHECK: %[[V22:[0-9]+]] = getelementptr inbounds {{(nuw )?}}i8, ptr %sa, i64 72

; Check the stores for the second memset
; CHECK: store i64 0, ptr addrspace(1) %[[SB_I8AS]]
; CHECK: %[[V24:[0-9]+]] = getelementptr inbounds i8, ptr addrspace(1) %[[SB_I8AS]], i64 8
; CHECK: %[[V24:[0-9]+]] = getelementptr inbounds {{(nuw )?}}i8, ptr addrspace(1) %[[SB_I8AS]], i64 8
; CHECK: store i64 0, ptr addrspace(1) %[[V24]]
; CHECK: %[[V26:[0-9]+]] = getelementptr inbounds i8, ptr addrspace(1) %[[SB_I8AS]], i64 16
; CHECK: %[[V26:[0-9]+]] = getelementptr inbounds {{(nuw )?}}i8, ptr addrspace(1) %[[SB_I8AS]], i64 16
; CHECK: store i64 0, ptr addrspace(1) %[[V26]]
; CHECK: %[[V28:[0-9]+]] = getelementptr inbounds i8, ptr addrspace(1) %[[SB_I8AS]], i64 24
; CHECK: %[[V28:[0-9]+]] = getelementptr inbounds {{(nuw )?}}i8, ptr addrspace(1) %[[SB_I8AS]], i64 24
; CHECK: store i64 0, ptr addrspace(1) %[[V28]]
; CHECK: %[[V30:[0-9]+]] = getelementptr inbounds i8, ptr addrspace(1) %[[SB_I8AS]], i64 32
; CHECK: %[[V30:[0-9]+]] = getelementptr inbounds {{(nuw )?}}i8, ptr addrspace(1) %[[SB_I8AS]], i64 32
; CHECK: store i64 0, ptr addrspace(1) %[[V30]]
; CHECK: %[[V32:[0-9]+]] = getelementptr inbounds i8, ptr addrspace(1) %[[SB_I8AS]], i64 40
; CHECK: %[[V32:[0-9]+]] = getelementptr inbounds {{(nuw )?}}i8, ptr addrspace(1) %[[SB_I8AS]], i64 40
; CHECK: store i64 0, ptr addrspace(1) %[[V32]]
; CHECK: %[[V33:[0-9]+]] = getelementptr inbounds i8, ptr addrspace(1) %[[SB_I8AS]], i64 48
; CHECK: %[[V33:[0-9]+]] = getelementptr inbounds {{(nuw )?}}i8, ptr addrspace(1) %[[SB_I8AS]], i64 48
; CHECK: store i64 0, ptr addrspace(1) %[[V33]]
; CHECK: %[[V35T:[0-9]+]] = getelementptr inbounds i8, ptr addrspace(1) %[[SB_I8AS]], i64 56
; CHECK-EQ14: %[[V35T:[0-9]+]] = getelementptr inbounds %struct.S2, %struct.S2* %sb, i64 0, i32 3, i64 8
; CHECK: %[[V35T:[0-9]+]] = getelementptr inbounds {{(nuw )?}}i8, ptr addrspace(1) %[[SB_I8AS]], i64 56
; CHECK-EQ14: %[[V35T:[0-9]+]] = getelementptr inbounds {{(nuw )?}}%struct.S2, %struct.S2* %sb, i64 0, i32 3, i64 8
; CHECK-EQ14: %[[V35:[0-9]+]] = bitcast i8* %[[V35T]] to i64*
; CHECK-EQ14: %[[SB_I8AS18:.+]] = addrspacecast i64* %[[V35]] to i64 addrspace(1)*
; CHECK: store i64 0, ptr addrspace(1) %[[V35T]]
; CHECK: %[[V36:[0-9]+]] = getelementptr inbounds i8, ptr addrspace(1) %[[SB_I8AS]], i64 64
; CHECK: %[[V36:[0-9]+]] = getelementptr inbounds {{(nuw )?}}i8, ptr addrspace(1) %[[SB_I8AS]], i64 64
; CHECK: store i64 0, ptr addrspace(1) %[[V36]]
; CHECK: %[[V38:[0-9]+]] = getelementptr inbounds i8, ptr addrspace(1) %[[SB_I8AS]], i64 72
; CHECK: %[[V38:[0-9]+]] = getelementptr inbounds {{(nuw )?}}i8, ptr addrspace(1) %[[SB_I8AS]], i64 72
; CHECK: store i64 0, ptr addrspace(1) %[[V38]]


Expand Down Expand Up @@ -167,32 +167,32 @@ declare i64 @__mux_get_local_id(i32)
; CHECK:end: ; preds = %middle, %entry
; CHECK: %[[SB_I8AS42:.+]] = load i64, ptr addrspace(1) %[[SB_I8AS]]
; CHECK: store i64 %[[SB_I8AS42]], ptr %result2
; CHECK: %[[V42:[0-9]+]] = getelementptr inbounds i8, ptr %result2, i64 8
; CHECK: %[[V42:[0-9]+]] = getelementptr inbounds {{(nuw )?}}i8, ptr %result2, i64 8
; CHECK: %[[SB_I8AS44:.+]] = load i64, ptr addrspace(1) %[[V24]]
; CHECK: store i64 %[[SB_I8AS44]], ptr %[[V42]]
; CHECK: %[[V43:[0-9]+]] = getelementptr inbounds i8, ptr %result2, i64 16
; CHECK: %[[V43:[0-9]+]] = getelementptr inbounds {{(nuw )?}}i8, ptr %result2, i64 16
; CHECK: %[[SB_I8AS46:.+]] = load i64, ptr addrspace(1) %[[V26]]
; CHECK: store i64 %[[SB_I8AS46]], ptr %[[V43]]
; CHECK: %[[V44:[0-9]+]] = getelementptr inbounds i8, ptr %result2, i64 24
; CHECK: %[[V44:[0-9]+]] = getelementptr inbounds {{(nuw )?}}i8, ptr %result2, i64 24
; CHECK: %[[SB_I8AS48:.+]] = load i64, ptr addrspace(1) %[[V28]]
; CHECK: store i64 %[[SB_I8AS48]], ptr %[[V44]]
; CHECK: %[[V45:[0-9]+]] = getelementptr inbounds i8, ptr %result2, i64 32
; CHECK: %[[V45:[0-9]+]] = getelementptr inbounds {{(nuw )?}}i8, ptr %result2, i64 32
; CHECK: %[[SB_I8AS50:.+]] = load i64, ptr addrspace(1) %[[V30]]
; CHECK: store i64 %[[SB_I8AS50]], ptr %[[V45]]
; CHECK: %[[V46:[0-9]+]] = getelementptr inbounds i8, ptr %result2, i64 40
; CHECK: %[[V46:[0-9]+]] = getelementptr inbounds {{(nuw )?}}i8, ptr %result2, i64 40
; CHECK: %[[SB_I8AS52:.+]] = load i64, ptr addrspace(1) %[[V32]]
; CHECK: store i64 %[[SB_I8AS52]], ptr %[[V46]]
; CHECK: %[[V47:[0-9]+]] = getelementptr inbounds i8, ptr %result2, i64 48
; CHECK: %[[V47:[0-9]+]] = getelementptr inbounds {{(nuw )?}}i8, ptr %result2, i64 48
; CHECK: %[[SB_I8AS54:.+]] = load i64, ptr addrspace(1) %[[V33]]
; CHECK: store i64 %[[SB_I8AS54]], ptr %[[V47]]
; CHECK: %[[V48:[0-9]+]] = getelementptr inbounds i8, ptr %result2, i64 56
; CHECK-EQ14: %[[V48:[0-9]+]] = getelementptr inbounds %struct.S2, %struct.S2* %result2, i64 0, i32 3, i64 8
; CHECK: %[[V48:[0-9]+]] = getelementptr inbounds {{(nuw )?}}i8, ptr %result2, i64 56
; CHECK-EQ14: %[[V48:[0-9]+]] = getelementptr inbounds {{(nuw )?}}%struct.S2, %struct.S2* %result2, i64 0, i32 3, i64 8
; CHECK: %[[SB_I8AS56:.+]] = load i64, ptr addrspace(1) %[[V35T]]
; CHECK: store i64 %[[SB_I8AS56]], ptr %[[V48]]
; CHECK: %[[V49:[0-9]+]] = getelementptr inbounds i8, ptr %result2, i64 64
; CHECK: %[[V49:[0-9]+]] = getelementptr inbounds {{(nuw )?}}i8, ptr %result2, i64 64
; CHECK: %[[SB_I8AS58:.+]] = load i64, ptr addrspace(1) %[[V36]]
; CHECK: store i64 %[[SB_I8AS58]], ptr %[[V49]]
; CHECK: %[[V50:[0-9]+]] = getelementptr inbounds i8, ptr %result2, i64 72
; CHECK: %[[V50:[0-9]+]] = getelementptr inbounds {{(nuw )?}}i8, ptr %result2, i64 72
; CHECK: %[[SB_I8AS60:.+]] = load i64, ptr addrspace(1) %[[V38]]
; CHECK: store i64 %[[SB_I8AS60]], ptr %[[V50]]

Expand Down
2 changes: 1 addition & 1 deletion modules/compiler/vecz/test/lit/llvm/gep_duplication.ll
View file
Original file line number Diff line number Diff line change
Expand Up @@ -26,7 +26,7 @@ target datalayout = "e-p:64:64:64-m:e-i64:64-f80:128-n8:16:32:64-S128"
; combination of instcombine and GVN).
; CHECK: spir_kernel void @__vecz_v{{[0-9]+}}_gep_duplication
; CHECK: entry:
; CHECK: getelementptr inbounds {{\[2 x i32]|i8}}, ptr %myStruct, {{i64 0, i64 1|i64 4}}
; CHECK: getelementptr inbounds {{(nuw )?}}{{\[2 x i32]|i8}}, ptr %myStruct, {{i64 0, i64 1|i64 4}}
; CHECK-NOT: getelementptr {{.*}}%myStruct
define spir_kernel void @gep_duplication(ptr addrspace(1) align 4 %out) {
entry:
Expand Down

0 comments on commit e44c154

Please sign in to comment.