https://github.com/steffenlarsen created https://github.com/llvm/llvm-project/pull/237178
The device runtime reads each non-string argument of a buffered printf from 8-byte slots, and the buffer reserves at least 8 bytes for each. The pointer to the next argument was advanced by the stored value's own alloc size, though, so a value stored in fewer than 8 bytes shifted every argument after it. Round both the reservation and the advance for a non-string argument up to a multiple of 8 bytes. The words of a constant string are still stored back to back, as the string as a whole is already padded to 8 bytes. Store non-string arguments in their own case so their advance is separate from the string words, which flattens the argument loop. Assisted-by: Claude Opus 5.5 >From ba9ba51b37c593de2bdc76f9672cf46cfe183f38 Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Fri, 30 Oct 2026 06:09:09 -0500 Subject: [PATCH] [AMDGPU] Fix 8-byte printf slot placement The device runtime reads each non-string argument of a buffered printf from 8-byte slots, and the buffer reserves at least 8 bytes for each. The pointer to the next argument was advanced by the stored value's own alloc size, though, so a value stored in fewer than 8 bytes shifted every argument after it. Round both the reservation and the advance for a non-string argument up to a multiple of 8 bytes. The words of a constant string are still stored back to back, as the string as a whole is already padded to 8 bytes. Store non-string arguments in their own case so their advance is separate from the string words, which flattens the argument loop. Assisted-by: Claude Opus 5.5 Signed-off-by: Steffen Holst Larsen <[email protected]> --- clang/test/CodeGenHIP/printf_nonhostcall.cpp | 73 +++++++++++++++++++ .../lib/Transforms/Utils/AMDGPUEmitPrintf.cpp | 72 ++++++++++-------- 2 files changed, 116 insertions(+), 29 deletions(-) diff --git a/clang/test/CodeGenHIP/printf_nonhostcall.cpp b/clang/test/CodeGenHIP/printf_nonhostcall.cpp index 7e8a468a335271..956d8950bcb18a 100644 --- a/clang/test/CodeGenHIP/printf_nonhostcall.cpp +++ b/clang/test/CodeGenHIP/printf_nonhostcall.cpp @@ -460,3 +460,76 @@ __device__ int foo4() { const char* s = "format str%d"; return printf(s, 10); } + +// A 32-bit pointer, here to LDS or private memory, is stored in 4 bytes but +// still takes an 8-byte slot, so the argument after it starts 8 bytes later. +// CHECK-LABEL: define dso_local noundef i32 @_Z4foo5PU3AS3vPU3AS5v +// CHECK-SAME: (ptr addrspace(3) noundef [[LDS:%.*]], ptr addrspace(5) noundef [[PRIV:%.*]]) #{{[0-9]+}} { +// CHECK-NEXT: entry: +// CHECK-NEXT: [[LDS_ADDR:%.*]] = alloca ptr addrspace(3), align 4, addrspace(5) +// CHECK-NEXT: [[PRIV_ADDR:%.*]] = alloca ptr addrspace(5), align 4, addrspace(5) +// CHECK-NEXT: [[LDS_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[LDS_ADDR]] to ptr +// CHECK-NEXT: [[PRIV_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[PRIV_ADDR]] to ptr +// CHECK-NEXT: store ptr addrspace(3) [[LDS]], ptr [[LDS_ADDR_ASCAST]], align 4 +// CHECK-NEXT: store ptr addrspace(5) [[PRIV]], ptr [[PRIV_ADDR_ASCAST]], align 4 +// CHECK-NEXT: [[TMP0:%.*]] = load ptr addrspace(3), ptr [[LDS_ADDR_ASCAST]], align 4 +// CHECK-NEXT: [[TMP1:%.*]] = load ptr addrspace(5), ptr [[PRIV_ADDR_ASCAST]], align 4 +// CHECK-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 44) +// CHECK-NEXT: [[TMP2:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null +// CHECK-NEXT: br i1 [[TMP2]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]] +// CHECK: end.block: +// CHECK-NEXT: [[TMP3:%.*]] = xor i1 [[TMP2]], true +// CHECK-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP3]] to i32 +// CHECK-NEXT: ret i32 [[PRINTF_RESULT]] +// CHECK: argpush.block: +// CHECK-NEXT: store i32 178, ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4 +// CHECK-NEXT: [[TMP4:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4 +// CHECK-NEXT: store i64 -2001953306523763028, ptr addrspace(1) [[TMP4]], align 8 +// CHECK-NEXT: [[TMP5:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP4]], i32 8 +// CHECK-NEXT: store ptr addrspace(3) [[TMP0]], ptr addrspace(1) [[TMP5]], align 4 +// CHECK-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP5]], i32 8 +// CHECK-NEXT: store i64 1, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8 +// CHECK-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8 +// CHECK-NEXT: store ptr addrspace(5) [[TMP1]], ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], align 4 +// CHECK-NEXT: [[PRINTBUFFNEXTPTR2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], i32 8 +// CHECK-NEXT: store i64 2, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], align 8 +// CHECK-NEXT: [[PRINTBUFFNEXTPTR3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], i32 8 +// CHECK-NEXT: br label [[END_BLOCK]] +// +// CHECK_CONSTRAINED-LABEL: define dso_local noundef i32 @_Z4foo5PU3AS3vPU3AS5v +// CHECK_CONSTRAINED-SAME: (ptr addrspace(3) noundef [[LDS:%.*]], ptr addrspace(5) noundef [[PRIV:%.*]]) #{{[0-9]+}} { +// CHECK_CONSTRAINED-NEXT: entry: +// CHECK_CONSTRAINED-NEXT: [[LDS_ADDR:%.*]] = alloca ptr addrspace(3), align 4, addrspace(5) +// CHECK_CONSTRAINED-NEXT: [[PRIV_ADDR:%.*]] = alloca ptr addrspace(5), align 4, addrspace(5) +// CHECK_CONSTRAINED-NEXT: [[LDS_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[LDS_ADDR]] to ptr +// CHECK_CONSTRAINED-NEXT: [[PRIV_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[PRIV_ADDR]] to ptr +// CHECK_CONSTRAINED-NEXT: store ptr addrspace(3) [[LDS]], ptr [[LDS_ADDR_ASCAST]], align 4 +// CHECK_CONSTRAINED-NEXT: store ptr addrspace(5) [[PRIV]], ptr [[PRIV_ADDR_ASCAST]], align 4 +// CHECK_CONSTRAINED-NEXT: [[TMP0:%.*]] = load ptr addrspace(3), ptr [[LDS_ADDR_ASCAST]], align 4 +// CHECK_CONSTRAINED-NEXT: [[TMP1:%.*]] = load ptr addrspace(5), ptr [[PRIV_ADDR_ASCAST]], align 4 +// CHECK_CONSTRAINED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 44) +// CHECK_CONSTRAINED-NEXT: [[TMP2:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null +// CHECK_CONSTRAINED-NEXT: br i1 [[TMP2]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]] +// CHECK_CONSTRAINED: end.block: +// CHECK_CONSTRAINED-NEXT: [[TMP3:%.*]] = xor i1 [[TMP2]], true +// CHECK_CONSTRAINED-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP3]] to i32 +// CHECK_CONSTRAINED-NEXT: ret i32 [[PRINTF_RESULT]] +// CHECK_CONSTRAINED: argpush.block: +// CHECK_CONSTRAINED-NEXT: store i32 178, ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4 +// CHECK_CONSTRAINED-NEXT: [[TMP4:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4 +// CHECK_CONSTRAINED-NEXT: store i64 -2001953306523763028, ptr addrspace(1) [[TMP4]], align 8 +// CHECK_CONSTRAINED-NEXT: [[TMP5:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP4]], i32 8 +// CHECK_CONSTRAINED-NEXT: store ptr addrspace(3) [[TMP0]], ptr addrspace(1) [[TMP5]], align 4 +// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP5]], i32 8 +// CHECK_CONSTRAINED-NEXT: store i64 1, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8 +// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8 +// CHECK_CONSTRAINED-NEXT: store ptr addrspace(5) [[TMP1]], ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], align 4 +// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], i32 8 +// CHECK_CONSTRAINED-NEXT: store i64 2, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], align 8 +// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], i32 8 +// CHECK_CONSTRAINED-NEXT: br label [[END_BLOCK]] +// +__device__ int foo5(__attribute__((address_space(3))) void *lds, + __attribute__((address_space(5))) void *priv) { + return printf("%p %d %p %d\n", lds, 1, priv, 2); +} diff --git a/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp b/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp index 11fb00f9b6a373..2c233f77b35f30 100644 --- a/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp +++ b/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp @@ -274,10 +274,11 @@ static Value *callBufferedPrintfStart( StringData(StringRef(), LenWithNull, LenWithNullAligned, false)); } } else { - int AllocSize = M->getDataLayout().getTypeAllocSize(Args[i]->getType()); + uint64_t AllocSize = + M->getDataLayout().getTypeAllocSize(Args[i]->getType()); // We end up expanding non string arguments to 8 bytes // (args smaller than 8 bytes) - BufSize += std::max(AllocSize, 8); + BufSize += alignTo(AllocSize, 8); } } @@ -377,35 +378,48 @@ callBufferedPrintfArgPush(IRBuilder<> &Builder, ArrayRef<Value *> Args, auto StrIt = StringContents.begin(); size_t i = IsConstFmtStr ? 1 : 0; for (; i < Args.size(); i++) { - SmallVector<Value *, 32> WhatToStore; - if ((i == 0) || SpecIsCString.test(i)) { - if (StrIt->IsConst) { - processConstantStringArg(StrIt, Builder, WhatToStore); - StrIt++; - } else { - // This copies the contents of the string, however the next offset - // is at aligned length, the extra space that might be created due - // to alignment padding is not populated with any specific value - // here. This would be safe as long as runtime is sync with - // the offsets. - Builder.CreateMemCpy(PtrToStore, /*DstAlign*/ Align(1), Args[i], - /*SrcAlign*/ Args[i]->getPointerAlignment(DL), - StrIt->RealSize); - - PtrToStore = - Builder.CreateInBoundsGEP(Builder.getInt8Ty(), PtrToStore, - {StrIt->AlignedSize}, "PrintBuffNextPtr"); - LLVM_DEBUG(dbgs() << "inserting gep to the printf buffer:" - << *PtrToStore << '\n'); - - // done with current argument, move to next - StrIt++; - continue; - } - } else { - WhatToStore.push_back(processNonStringArg(Args[i], Builder)); + if (i != 0 && !SpecIsCString.test(i)) { + // The runtime reads each non-string argument from a whole number of + // 8-byte slots, also when the stored value is smaller. + Value *ToStore = processNonStringArg(Args[i], Builder); + StoreInst *StBuff = Builder.CreateStore(ToStore, PtrToStore); + LLVM_DEBUG(dbgs() << "inserting store to printf buffer:" << *StBuff + << '\n'); + (void)StBuff; + PtrToStore = Builder.CreateConstInBoundsGEP1_32( + Builder.getInt8Ty(), PtrToStore, + alignTo(DL.getTypeAllocSize(ToStore->getType()), 8), + "PrintBuffNextPtr"); + LLVM_DEBUG(dbgs() << "inserting gep to the printf buffer:" << *PtrToStore + << '\n'); + continue; + } + + if (!StrIt->IsConst) { + // This copies the contents of the string, however the next offset + // is at aligned length, the extra space that might be created due + // to alignment padding is not populated with any specific value + // here. This would be safe as long as runtime is sync with + // the offsets. + Builder.CreateMemCpy(PtrToStore, /*DstAlign*/ Align(1), Args[i], + /*SrcAlign*/ Args[i]->getPointerAlignment(DL), + StrIt->RealSize); + + PtrToStore = + Builder.CreateInBoundsGEP(Builder.getInt8Ty(), PtrToStore, + {StrIt->AlignedSize}, "PrintBuffNextPtr"); + LLVM_DEBUG(dbgs() << "inserting gep to the printf buffer:" << *PtrToStore + << '\n'); + + // done with current argument, move to next + StrIt++; + continue; } + SmallVector<Value *, 32> WhatToStore; + processConstantStringArg(StrIt, Builder, WhatToStore); + StrIt++; + for (Value *toStore : WhatToStore) { StoreInst *StBuff = Builder.CreateStore(toStore, PtrToStore); LLVM_DEBUG(dbgs() << "inserting store to printf buffer:" << *StBuff _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
