Author: Nikita Popov Date: 2026-02-19T16:02:54+01:00 New Revision: 87eee80dad79417e079c369b9ff5578873019b78
URL: https://github.com/llvm/llvm-project/commit/87eee80dad79417e079c369b9ff5578873019b78 DIFF: https://github.com/llvm/llvm-project/commit/87eee80dad79417e079c369b9ff5578873019b78.diff LOG: [AMDGPUEmitPrintf] Use CreatePtrDiff() (#182283) Use CreatePtrDiff() to emit the pointer subtraction, which will use ptrtoaddr instead of ptrtoint. Add a conservative cast to i64 as the return value of CreatePtrDiff is no longer guaranteed to be a i64. Added: Modified: clang/test/CodeGenHIP/printf.cpp clang/test/CodeGenHIP/printf_nonhostcall.cpp llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp Removed: ################################################################################ diff --git a/clang/test/CodeGenHIP/printf.cpp b/clang/test/CodeGenHIP/printf.cpp index 2dc08aa1e5dd9..1f8e8ef9c3fb7 100644 --- a/clang/test/CodeGenHIP/printf.cpp +++ b/clang/test/CodeGenHIP/printf.cpp @@ -27,8 +27,8 @@ extern "C" __device__ int printf(const char *format, ...); // AMDGCN-NEXT: [[TMP7:%.*]] = icmp eq i8 [[TMP6]], 0 // AMDGCN-NEXT: br i1 [[TMP7]], label %[[STRLEN_WHILE_DONE:.*]], label %[[STRLEN_WHILE]] // AMDGCN: [[STRLEN_WHILE_DONE]]: -// AMDGCN-NEXT: [[TMP8:%.*]] = ptrtoint ptr [[TMP4]] to i64 -// AMDGCN-NEXT: [[TMP9:%.*]] = sub i64 [[TMP8]], ptrtoint (ptr addrspacecast (ptr addrspace(4) @.str.1 to ptr) to i64) +// AMDGCN-NEXT: [[TMP8:%.*]] = ptrtoaddr ptr [[TMP4]] to i64 +// AMDGCN-NEXT: [[TMP9:%.*]] = sub i64 ptrtoaddr (ptr addrspacecast (ptr addrspace(4) @.str.1 to ptr) to i64), [[TMP8]] // AMDGCN-NEXT: [[TMP10:%.*]] = add i64 [[TMP9]], 1 // AMDGCN-NEXT: br label %[[STRLEN_JOIN]] // AMDGCN: [[STRLEN_JOIN]]: @@ -47,9 +47,9 @@ extern "C" __device__ int printf(const char *format, ...); // AMDGCN-NEXT: [[TMP21:%.*]] = icmp eq i8 [[TMP20]], 0 // AMDGCN-NEXT: br i1 [[TMP21]], label %[[STRLEN_WHILE_DONE3:.*]], label %[[STRLEN_WHILE2]] // AMDGCN: [[STRLEN_WHILE_DONE3]]: -// AMDGCN-NEXT: [[TMP22:%.*]] = ptrtoint ptr [[TMP0]] to i64 -// AMDGCN-NEXT: [[TMP23:%.*]] = ptrtoint ptr [[TMP18]] to i64 -// AMDGCN-NEXT: [[TMP24:%.*]] = sub i64 [[TMP23]], [[TMP22]] +// AMDGCN-NEXT: [[TMP22:%.*]] = ptrtoaddr ptr [[TMP0]] to i64 +// AMDGCN-NEXT: [[TMP23:%.*]] = ptrtoaddr ptr [[TMP18]] to i64 +// AMDGCN-NEXT: [[TMP24:%.*]] = sub i64 [[TMP22]], [[TMP23]] // AMDGCN-NEXT: [[TMP25:%.*]] = add i64 [[TMP24]], 1 // AMDGCN-NEXT: br label %[[STRLEN_JOIN1]] // AMDGCN: [[STRLEN_JOIN1]]: @@ -78,8 +78,8 @@ extern "C" __device__ int printf(const char *format, ...); // AMDGCNSPIRV-NEXT: [[TMP7:%.*]] = icmp eq i8 [[TMP6]], 0 // AMDGCNSPIRV-NEXT: br i1 [[TMP7]], label %[[STRLEN_WHILE_DONE:.*]], label %[[STRLEN_WHILE]] // AMDGCNSPIRV: [[STRLEN_WHILE_DONE]]: -// AMDGCNSPIRV-NEXT: [[TMP8:%.*]] = ptrtoint ptr addrspace(4) [[TMP4]] to i64 -// AMDGCNSPIRV-NEXT: [[TMP9:%.*]] = sub i64 [[TMP8]], ptrtoint (ptr addrspace(4) addrspacecast (ptr addrspace(1) @.str.1 to ptr addrspace(4)) to i64) +// AMDGCNSPIRV-NEXT: [[TMP8:%.*]] = ptrtoaddr ptr addrspace(4) [[TMP4]] to i64 +// AMDGCNSPIRV-NEXT: [[TMP9:%.*]] = sub i64 ptrtoaddr (ptr addrspace(4) addrspacecast (ptr addrspace(1) @.str.1 to ptr addrspace(4)) to i64), [[TMP8]] // AMDGCNSPIRV-NEXT: [[TMP10:%.*]] = add i64 [[TMP9]], 1 // AMDGCNSPIRV-NEXT: br label %[[STRLEN_JOIN]] // AMDGCNSPIRV: [[STRLEN_JOIN]]: @@ -98,9 +98,9 @@ extern "C" __device__ int printf(const char *format, ...); // AMDGCNSPIRV-NEXT: [[TMP21:%.*]] = icmp eq i8 [[TMP20]], 0 // AMDGCNSPIRV-NEXT: br i1 [[TMP21]], label %[[STRLEN_WHILE_DONE3:.*]], label %[[STRLEN_WHILE2]] // AMDGCNSPIRV: [[STRLEN_WHILE_DONE3]]: -// AMDGCNSPIRV-NEXT: [[TMP22:%.*]] = ptrtoint ptr addrspace(4) [[TMP0]] to i64 -// AMDGCNSPIRV-NEXT: [[TMP23:%.*]] = ptrtoint ptr addrspace(4) [[TMP18]] to i64 -// AMDGCNSPIRV-NEXT: [[TMP24:%.*]] = sub i64 [[TMP23]], [[TMP22]] +// AMDGCNSPIRV-NEXT: [[TMP22:%.*]] = ptrtoaddr ptr addrspace(4) [[TMP0]] to i64 +// AMDGCNSPIRV-NEXT: [[TMP23:%.*]] = ptrtoaddr ptr addrspace(4) [[TMP18]] to i64 +// AMDGCNSPIRV-NEXT: [[TMP24:%.*]] = sub i64 [[TMP22]], [[TMP23]] // AMDGCNSPIRV-NEXT: [[TMP25:%.*]] = add i64 [[TMP24]], 1 // AMDGCNSPIRV-NEXT: br label %[[STRLEN_JOIN1]] // AMDGCNSPIRV: [[STRLEN_JOIN1]]: @@ -133,8 +133,8 @@ __device__ char *dstr; // AMDGCN-NEXT: [[TMP7:%.*]] = icmp eq i8 [[TMP6]], 0 // AMDGCN-NEXT: br i1 [[TMP7]], label %[[STRLEN_WHILE_DONE:.*]], label %[[STRLEN_WHILE]] // AMDGCN: [[STRLEN_WHILE_DONE]]: -// AMDGCN-NEXT: [[TMP8:%.*]] = ptrtoint ptr [[TMP4]] to i64 -// AMDGCN-NEXT: [[TMP9:%.*]] = sub i64 [[TMP8]], ptrtoint (ptr addrspacecast (ptr addrspace(4) @.str.2 to ptr) to i64) +// AMDGCN-NEXT: [[TMP8:%.*]] = ptrtoaddr ptr [[TMP4]] to i64 +// AMDGCN-NEXT: [[TMP9:%.*]] = sub i64 ptrtoaddr (ptr addrspacecast (ptr addrspace(4) @.str.2 to ptr) to i64), [[TMP8]] // AMDGCN-NEXT: [[TMP10:%.*]] = add i64 [[TMP9]], 1 // AMDGCN-NEXT: br label %[[STRLEN_JOIN]] // AMDGCN: [[STRLEN_JOIN]]: @@ -149,9 +149,9 @@ __device__ char *dstr; // AMDGCN-NEXT: [[TMP17:%.*]] = icmp eq i8 [[TMP16]], 0 // AMDGCN-NEXT: br i1 [[TMP17]], label %[[STRLEN_WHILE_DONE3:.*]], label %[[STRLEN_WHILE2]] // AMDGCN: [[STRLEN_WHILE_DONE3]]: -// AMDGCN-NEXT: [[TMP18:%.*]] = ptrtoint ptr [[TMP0]] to i64 -// AMDGCN-NEXT: [[TMP19:%.*]] = ptrtoint ptr [[TMP14]] to i64 -// AMDGCN-NEXT: [[TMP20:%.*]] = sub i64 [[TMP19]], [[TMP18]] +// AMDGCN-NEXT: [[TMP18:%.*]] = ptrtoaddr ptr [[TMP0]] to i64 +// AMDGCN-NEXT: [[TMP19:%.*]] = ptrtoaddr ptr [[TMP14]] to i64 +// AMDGCN-NEXT: [[TMP20:%.*]] = sub i64 [[TMP18]], [[TMP19]] // AMDGCN-NEXT: [[TMP21:%.*]] = add i64 [[TMP20]], 1 // AMDGCN-NEXT: br label %[[STRLEN_JOIN1]] // AMDGCN: [[STRLEN_JOIN1]]: @@ -177,8 +177,8 @@ __device__ char *dstr; // AMDGCNSPIRV-NEXT: [[TMP7:%.*]] = icmp eq i8 [[TMP6]], 0 // AMDGCNSPIRV-NEXT: br i1 [[TMP7]], label %[[STRLEN_WHILE_DONE:.*]], label %[[STRLEN_WHILE]] // AMDGCNSPIRV: [[STRLEN_WHILE_DONE]]: -// AMDGCNSPIRV-NEXT: [[TMP8:%.*]] = ptrtoint ptr addrspace(4) [[TMP4]] to i64 -// AMDGCNSPIRV-NEXT: [[TMP9:%.*]] = sub i64 [[TMP8]], ptrtoint (ptr addrspace(4) addrspacecast (ptr addrspace(1) @.str.2 to ptr addrspace(4)) to i64) +// AMDGCNSPIRV-NEXT: [[TMP8:%.*]] = ptrtoaddr ptr addrspace(4) [[TMP4]] to i64 +// AMDGCNSPIRV-NEXT: [[TMP9:%.*]] = sub i64 ptrtoaddr (ptr addrspace(4) addrspacecast (ptr addrspace(1) @.str.2 to ptr addrspace(4)) to i64), [[TMP8]] // AMDGCNSPIRV-NEXT: [[TMP10:%.*]] = add i64 [[TMP9]], 1 // AMDGCNSPIRV-NEXT: br label %[[STRLEN_JOIN]] // AMDGCNSPIRV: [[STRLEN_JOIN]]: @@ -193,9 +193,9 @@ __device__ char *dstr; // AMDGCNSPIRV-NEXT: [[TMP17:%.*]] = icmp eq i8 [[TMP16]], 0 // AMDGCNSPIRV-NEXT: br i1 [[TMP17]], label %[[STRLEN_WHILE_DONE3:.*]], label %[[STRLEN_WHILE2]] // AMDGCNSPIRV: [[STRLEN_WHILE_DONE3]]: -// AMDGCNSPIRV-NEXT: [[TMP18:%.*]] = ptrtoint ptr addrspace(4) [[TMP0]] to i64 -// AMDGCNSPIRV-NEXT: [[TMP19:%.*]] = ptrtoint ptr addrspace(4) [[TMP14]] to i64 -// AMDGCNSPIRV-NEXT: [[TMP20:%.*]] = sub i64 [[TMP19]], [[TMP18]] +// AMDGCNSPIRV-NEXT: [[TMP18:%.*]] = ptrtoaddr ptr addrspace(4) [[TMP0]] to i64 +// AMDGCNSPIRV-NEXT: [[TMP19:%.*]] = ptrtoaddr ptr addrspace(4) [[TMP14]] to i64 +// AMDGCNSPIRV-NEXT: [[TMP20:%.*]] = sub i64 [[TMP18]], [[TMP19]] // AMDGCNSPIRV-NEXT: [[TMP21:%.*]] = add i64 [[TMP20]], 1 // AMDGCNSPIRV-NEXT: br label %[[STRLEN_JOIN1]] // AMDGCNSPIRV: [[STRLEN_JOIN1]]: diff --git a/clang/test/CodeGenHIP/printf_nonhostcall.cpp b/clang/test/CodeGenHIP/printf_nonhostcall.cpp index a05b8166eda8a..e252bc4019c02 100644 --- a/clang/test/CodeGenHIP/printf_nonhostcall.cpp +++ b/clang/test/CodeGenHIP/printf_nonhostcall.cpp @@ -27,9 +27,9 @@ extern "C" __device__ int printf(const char *format, ...); // CHECK-NEXT: [[TMP6:%.*]] = icmp eq i8 [[TMP5]], 0 // CHECK-NEXT: br i1 [[TMP6]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]] // CHECK: strlen.while.done: -// CHECK-NEXT: [[TMP7:%.*]] = ptrtoint ptr [[TMP0]] to i64 -// CHECK-NEXT: [[TMP8:%.*]] = ptrtoint ptr [[TMP3]] to i64 -// CHECK-NEXT: [[TMP9:%.*]] = sub i64 [[TMP8]], [[TMP7]] +// CHECK-NEXT: [[TMP7:%.*]] = ptrtoaddr ptr [[TMP0]] to i64 +// CHECK-NEXT: [[TMP8:%.*]] = ptrtoaddr ptr [[TMP3]] to i64 +// CHECK-NEXT: [[TMP9:%.*]] = sub i64 [[TMP7]], [[TMP8]] // CHECK-NEXT: [[TMP10:%.*]] = add i64 [[TMP9]], 1 // CHECK-NEXT: br label [[STRLEN_JOIN]] // CHECK: strlen.join: @@ -82,9 +82,9 @@ extern "C" __device__ int printf(const char *format, ...); // CHECK_CONSTRAINED-NEXT: [[TMP6:%.*]] = icmp eq i8 [[TMP5]], 0 // CHECK_CONSTRAINED-NEXT: br i1 [[TMP6]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]] // CHECK_CONSTRAINED: strlen.while.done: -// CHECK_CONSTRAINED-NEXT: [[TMP7:%.*]] = ptrtoint ptr [[TMP0]] to i64 -// CHECK_CONSTRAINED-NEXT: [[TMP8:%.*]] = ptrtoint ptr [[TMP3]] to i64 -// CHECK_CONSTRAINED-NEXT: [[TMP9:%.*]] = sub i64 [[TMP8]], [[TMP7]] +// CHECK_CONSTRAINED-NEXT: [[TMP7:%.*]] = ptrtoaddr ptr [[TMP0]] to i64 +// CHECK_CONSTRAINED-NEXT: [[TMP8:%.*]] = ptrtoaddr ptr [[TMP3]] to i64 +// CHECK_CONSTRAINED-NEXT: [[TMP9:%.*]] = sub i64 [[TMP7]], [[TMP8]] // CHECK_CONSTRAINED-NEXT: [[TMP10:%.*]] = add i64 [[TMP9]], 1 // CHECK_CONSTRAINED-NEXT: br label [[STRLEN_JOIN]] // CHECK_CONSTRAINED: strlen.join: @@ -143,9 +143,9 @@ __device__ const // CHECK-NEXT: [[TMP6:%.*]] = icmp eq i8 [[TMP5]], 0 // CHECK-NEXT: br i1 [[TMP6]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]] // CHECK: strlen.while.done: -// CHECK-NEXT: [[TMP7:%.*]] = ptrtoint ptr [[TMP0]] to i64 -// CHECK-NEXT: [[TMP8:%.*]] = ptrtoint ptr [[TMP3]] to i64 -// CHECK-NEXT: [[TMP9:%.*]] = sub i64 [[TMP8]], [[TMP7]] +// CHECK-NEXT: [[TMP7:%.*]] = ptrtoaddr ptr [[TMP0]] to i64 +// CHECK-NEXT: [[TMP8:%.*]] = ptrtoaddr ptr [[TMP3]] to i64 +// CHECK-NEXT: [[TMP9:%.*]] = sub i64 [[TMP7]], [[TMP8]] // CHECK-NEXT: [[TMP10:%.*]] = add i64 [[TMP9]], 1 // CHECK-NEXT: br label [[STRLEN_JOIN]] // CHECK: strlen.join: @@ -193,9 +193,9 @@ __device__ const // CHECK_CONSTRAINED-NEXT: [[TMP6:%.*]] = icmp eq i8 [[TMP5]], 0 // CHECK_CONSTRAINED-NEXT: br i1 [[TMP6]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]] // CHECK_CONSTRAINED: strlen.while.done: -// CHECK_CONSTRAINED-NEXT: [[TMP7:%.*]] = ptrtoint ptr [[TMP0]] to i64 -// CHECK_CONSTRAINED-NEXT: [[TMP8:%.*]] = ptrtoint ptr [[TMP3]] to i64 -// CHECK_CONSTRAINED-NEXT: [[TMP9:%.*]] = sub i64 [[TMP8]], [[TMP7]] +// CHECK_CONSTRAINED-NEXT: [[TMP7:%.*]] = ptrtoaddr ptr [[TMP0]] to i64 +// CHECK_CONSTRAINED-NEXT: [[TMP8:%.*]] = ptrtoaddr ptr [[TMP3]] to i64 +// CHECK_CONSTRAINED-NEXT: [[TMP9:%.*]] = sub i64 [[TMP7]], [[TMP8]] // CHECK_CONSTRAINED-NEXT: [[TMP10:%.*]] = add i64 [[TMP9]], 1 // CHECK_CONSTRAINED-NEXT: br label [[STRLEN_JOIN]] // CHECK_CONSTRAINED: strlen.join: @@ -385,9 +385,9 @@ __device__ int foo3() { // CHECK-NEXT: [[TMP5:%.*]] = icmp eq i8 [[TMP4]], 0 // CHECK-NEXT: br i1 [[TMP5]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]] // CHECK: strlen.while.done: -// CHECK-NEXT: [[TMP6:%.*]] = ptrtoint ptr [[TMP0]] to i64 -// CHECK-NEXT: [[TMP7:%.*]] = ptrtoint ptr [[TMP2]] to i64 -// CHECK-NEXT: [[TMP8:%.*]] = sub i64 [[TMP7]], [[TMP6]] +// CHECK-NEXT: [[TMP6:%.*]] = ptrtoaddr ptr [[TMP0]] to i64 +// CHECK-NEXT: [[TMP7:%.*]] = ptrtoaddr ptr [[TMP2]] to i64 +// CHECK-NEXT: [[TMP8:%.*]] = sub i64 [[TMP6]], [[TMP7]] // CHECK-NEXT: [[TMP9:%.*]] = add i64 [[TMP8]], 1 // CHECK-NEXT: br label [[STRLEN_JOIN]] // CHECK: strlen.join: @@ -428,9 +428,9 @@ __device__ int foo3() { // CHECK_CONSTRAINED-NEXT: [[TMP5:%.*]] = icmp eq i8 [[TMP4]], 0 // CHECK_CONSTRAINED-NEXT: br i1 [[TMP5]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]] // CHECK_CONSTRAINED: strlen.while.done: -// CHECK_CONSTRAINED-NEXT: [[TMP6:%.*]] = ptrtoint ptr [[TMP0]] to i64 -// CHECK_CONSTRAINED-NEXT: [[TMP7:%.*]] = ptrtoint ptr [[TMP2]] to i64 -// CHECK_CONSTRAINED-NEXT: [[TMP8:%.*]] = sub i64 [[TMP7]], [[TMP6]] +// CHECK_CONSTRAINED-NEXT: [[TMP6:%.*]] = ptrtoaddr ptr [[TMP0]] to i64 +// CHECK_CONSTRAINED-NEXT: [[TMP7:%.*]] = ptrtoaddr ptr [[TMP2]] to i64 +// CHECK_CONSTRAINED-NEXT: [[TMP8:%.*]] = sub i64 [[TMP6]], [[TMP7]] // CHECK_CONSTRAINED-NEXT: [[TMP9:%.*]] = add i64 [[TMP8]], 1 // CHECK_CONSTRAINED-NEXT: br label [[STRLEN_JOIN]] // CHECK_CONSTRAINED: strlen.join: diff --git a/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp b/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp index a25632acbfcc3..466f5b6878e92 100644 --- a/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp +++ b/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp @@ -136,9 +136,8 @@ static Value *getStrlenWithNull(IRBuilder<> &Builder, Value *Str) { // Add one to the computed length. Builder.SetInsertPoint(WhileDone, WhileDone->begin()); - auto Begin = Builder.CreatePtrToInt(Str, Int64Ty); - auto End = Builder.CreatePtrToInt(PtrPhi, Int64Ty); - auto Len = Builder.CreateSub(End, Begin); + auto Len = Builder.CreatePtrDiff(Str, PtrPhi); + Len = Builder.CreateZExt(Len, Int64Ty); Len = Builder.CreateAdd(Len, One); // Final join. _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
