vikramRH created this revision.
vikramRH added reviewers: sameerds, b-sumner, yaxunl, arsenm.
Herald added subscribers: hoy, kerbowa, hiraditya, Anastasia, tpr, dstuttard,
jvesely, kzhuravl.
Herald added a project: All.
vikramRH requested review of this revision.
Herald added subscribers: llvm-commits, cfe-commits, MaskRay, wdng.
Herald added projects: clang, LLVM.
This is an alternative to currently existing hostcall implementation and uses
printf buffer similar to OpenCL,
The data stored in the buffer (i.e the data frame) for each printf call are as
follows,
1. Control DWord - contains info regarding stream, format string constness and
size of data frame
2. Hash of the format string (if constant) else the format string itself
3. Printf arguments (each aligned to 8 byte boundary)
The format string Hash is generated using LLVM's MD5 Message-Digest Algorithm
implementation and only low 64 bits are used.
The implementation still uses amdhsa metadata and hash is stored as part of
format string itself to ensure
minimal changes in runtime.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D150427
Files:
clang/include/clang/Basic/LangOptions.def
clang/include/clang/Basic/LangOptions.h
clang/include/clang/Driver/Options.td
clang/lib/CodeGen/CGGPUBuiltin.cpp
clang/lib/Driver/ToolChains/Clang.cpp
clang/test/CodeGenHIP/printf_nonhostcall.cpp
clang/test/Driver/hip-options.hip
llvm/include/llvm/Transforms/Utils/AMDGPUEmitPrintf.h
llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp
Index: llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp
===================================================================
--- llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp
+++ llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp
@@ -17,6 +17,8 @@
#include "llvm/Transforms/Utils/AMDGPUEmitPrintf.h"
#include "llvm/ADT/SparseBitVector.h"
#include "llvm/Analysis/ValueTracking.h"
+#include "llvm/Support/DataExtractor.h"
+#include "llvm/Support/MD5.h"
using namespace llvm;
@@ -179,9 +181,8 @@
// Scan the format string to locate all specifiers, and mark the ones that
// specify a string, i.e, the "%s" specifier with optional '*' characters.
-static void locateCStrings(SparseBitVector<8> &BV, Value *Fmt) {
- StringRef Str;
- if (!getConstantStringInfo(Fmt, Str) || Str.empty())
+static void locateCStrings(SparseBitVector<8> &BV, Value *Fmt, StringRef &FmtStr) {
+ if (!getConstantStringInfo(Fmt, FmtStr) || FmtStr.empty())
return;
static const char ConvSpecifiers[] = "diouxXfFeEgGaAcspn";
@@ -189,17 +190,17 @@
// Skip the first argument, the format string.
unsigned ArgIdx = 1;
- while ((SpecPos = Str.find_first_of('%', SpecPos)) != StringRef::npos) {
- if (Str[SpecPos + 1] == '%') {
+ while ((SpecPos = FmtStr.find_first_of('%', SpecPos)) != StringRef::npos) {
+ if (FmtStr[SpecPos + 1] == '%') {
SpecPos += 2;
continue;
}
- auto SpecEnd = Str.find_first_of(ConvSpecifiers, SpecPos);
+ auto SpecEnd = FmtStr.find_first_of(ConvSpecifiers, SpecPos);
if (SpecEnd == StringRef::npos)
return;
- auto Spec = Str.slice(SpecPos, SpecEnd + 1);
+ auto Spec = FmtStr.slice(SpecPos, SpecEnd + 1);
ArgIdx += Spec.count('*');
- if (Str[SpecEnd] == 's') {
+ if (FmtStr[SpecEnd] == 's') {
BV.set(ArgIdx);
}
SpecPos = SpecEnd + 1;
@@ -207,14 +208,312 @@
}
}
+// helper struct to package the string related data
+typedef struct S {
+ std::string Str;
+ bool isConst;
+ Value *RealSize;
+ Value *AlignedSize;
+
+ S(std::string str = "", bool IC = true, Value *RS = nullptr,
+ Value *AS = nullptr)
+ : Str(str), isConst(IC), RealSize(RS), AlignedSize(AS) {}
+} StringData;
+
+static size_t alignUp(size_t Value, uint alignment) {
+ return (Value + alignment - 1) & ~(alignment - 1);
+}
+
+// Calculates frame size required for current printf expansion and allocates
+// space on printf buffer. Printf frame includes following contents
+// [ ControlDWord , format string/Hash , Arguments (each aligned to 8 byte) ]
+static Value *callBufferedPrintfStart(IRBuilder<> &Builder,
+ ArrayRef<Value *> &Args, Value *Fmt,
+ StringRef &FmtStr,
+ SparseBitVector<8> &SpecIsCString,
+ SmallVector<StringData, 8> &StringContents,
+ Value *&ArgSize) {
+ Value *NonConstStrLen = nullptr;
+
+ // First 8 bytes to be reserved for control dword
+ size_t BufSize = 4;
+ if (!FmtStr.empty())
+ // First 8 bytes of MD5 hash
+ BufSize += 8;
+ else {
+ auto LenWithNull = getStrlenWithNull(Builder, Fmt);
+
+ // Align the computed length to next 8 byte boundary
+ auto TempAdd = Builder.CreateAdd(
+ LenWithNull, ConstantInt::get(LenWithNull->getType(), 7U));
+ NonConstStrLen = Builder.CreateAnd(
+ TempAdd, ConstantInt::get(LenWithNull->getType(), ~7U));
+
+ StringContents.push_back(
+ StringData("", false, LenWithNull, NonConstStrLen));
+ }
+
+ StringRef ArgStr;
+ for (size_t i = 1; i < Args.size(); i++) {
+ if (SpecIsCString.test(i)) {
+ // This is a tradeoff. we might end up taking more compile
+ // time to calculate string contents if possible, but the generated
+ // code would be better runtime wise.
+ if (getConstantStringInfo(Args[i], ArgStr)) {
+ auto alignedLen = alignUp(ArgStr.size() + 1, 8);
+ StringContents.push_back(StringData(ArgStr.str() + '\0'));
+ BufSize += alignedLen;
+ } else {
+ auto LenWithNull = getStrlenWithNull(Builder, Args[i]);
+
+ // Align the computed length to next 8 byte boundary
+ auto TempAdd = Builder.CreateAdd(
+ LenWithNull, ConstantInt::get(LenWithNull->getType(), 7U));
+ auto LenWithNullAligned = Builder.CreateAnd(
+ TempAdd, ConstantInt::get(LenWithNull->getType(), ~7U));
+
+ if (NonConstStrLen) {
+ auto Val = Builder.CreateAdd(LenWithNullAligned, NonConstStrLen,
+ "cumulativeAdd");
+ NonConstStrLen = Val;
+ } else
+ NonConstStrLen = LenWithNullAligned;
+
+ StringContents.push_back(
+ StringData("", false, LenWithNull, LenWithNullAligned));
+ }
+ } else
+ // We end up expanding non string arguments to 8 bytes
+ BufSize += 8;
+ }
+
+ // calculate final size value to be passed to printf_alloc
+ Value *SizeToReserve = ConstantInt::get(Builder.getInt64Ty(), BufSize, false);
+ SmallVector<Value *, 1> Alloc_args;
+ if (NonConstStrLen)
+ SizeToReserve = Builder.CreateAdd(NonConstStrLen, SizeToReserve);
+
+ ArgSize = Builder.CreateTrunc(SizeToReserve, Builder.getInt32Ty());
+ Alloc_args.push_back(ArgSize);
+
+ // call the printf_alloc function
+ AttributeList Attr = AttributeList::get(
+ Builder.getContext(), AttributeList::FunctionIndex, Attribute::NoUnwind);
+
+ Type *Tys_alloc[1] = {Builder.getInt32Ty()};
+ Type *I8Ptr = Builder.getInt8PtrTy(1);
+ FunctionType *FTy_alloc = FunctionType::get(I8Ptr, Tys_alloc, false);
+ auto M = Builder.GetInsertBlock()->getModule();
+ auto PrintfAllocFn =
+ M->getOrInsertFunction(StringRef("__printf_alloc"), FTy_alloc, Attr);
+
+ return Builder.CreateCall(PrintfAllocFn, Alloc_args, "printf_alloc_fn");
+}
+
+static void callBufferedPrintfArgPush(IRBuilder<> &Builder, ArrayRef<Value *> Args,
+ Value *PtrToStore,
+ SparseBitVector<8> &SpecIsCString,
+ SmallVector<StringData, 8> &StringContents,
+ bool isConstFmtStr) {
+ auto StrIt = StringContents.begin();
+ size_t i = isConstFmtStr ? 1 : 0;
+ for (; i < Args.size(); i++) {
+ StringRef Str;
+ SmallVector<Value *, 32> WhatToStore;
+ if ((i == 0) || SpecIsCString.test(i)) {
+ if ((*StrIt).isConst) {
+ Str = (*StrIt).Str;
+ const uint64_t ReadSize = 4;
+
+ DataExtractor Extractor(Str, /*IsLittleEndian=*/true, 8);
+ DataExtractor::Cursor Offset(0);
+ while (Offset && Offset.tell() < Str.size()) {
+ uint64_t ReadNow = std::min(ReadSize, Str.size() - Offset.tell());
+ uint64_t ReadBytes = 0;
+ switch (ReadNow) {
+ default:
+ llvm_unreachable("min(4, X) > 4?");
+ case 1:
+ ReadBytes = Extractor.getU8(Offset);
+ break;
+ case 2:
+ ReadBytes = Extractor.getU16(Offset);
+ break;
+ case 3:
+ ReadBytes = Extractor.getU24(Offset);
+ break;
+ case 4:
+ ReadBytes = Extractor.getU32(Offset);
+ break;
+ }
+ cantFail(Offset.takeError(),
+ "failed to read bytes from constant array");
+
+ APInt IntVal(8 * ReadSize, ReadBytes);
+
+ // TODO: Should not bothering aligning up.
+ if (ReadNow < ReadSize)
+ IntVal = IntVal.zext(8 * ReadSize);
+
+ Type *IntTy =
+ Type::getIntNTy(Builder.getContext(), IntVal.getBitWidth());
+ WhatToStore.push_back(ConstantInt::get(IntTy, IntVal));
+ }
+ // Additional padding for 8 byte alignment
+ int Rem = (Str.size() % 8);
+ if (Rem > 0 && Rem <= 4)
+ WhatToStore.push_back(ConstantInt::get(Builder.getInt32Ty(), 0));
+
+ } else {
+ auto val = (*StrIt).RealSize;
+ Type *Tys[] = {PtrToStore->getType(), Args[i]->getType(),
+ val->getType()};
+ Function *TheFn = Intrinsic::getDeclaration(
+ Builder.GetInsertBlock()->getModule(), Intrinsic::memcpy, Tys);
+ SmallVector<Value *, 1> BuffOffset;
+
+ Value *FnArgs[] = {
+ PtrToStore, Args[i], val,
+ ConstantInt::get(Type::getInt1Ty(Builder.getContext()), false)};
+
+ // 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, I feel this would be safe as long as runtime is sync with
+ // the offsets.
+ Builder.CreateCall(TheFn, FnArgs, "");
+
+ BuffOffset.push_back((*StrIt).AlignedSize);
+ PtrToStore = Builder.CreateGEP(Builder.getInt8Ty(), PtrToStore,
+ BuffOffset, "PrintBuffNextPtr");
+ LLVM_DEBUG(dbgs() << "inserting gep to the printf buffer:\n"
+ << *PtrToStore << '\n');
+
+ // done with current argument, move to next
+ continue;
+ }
+ StrIt++;
+ } else
+ WhatToStore.push_back(fitArgInto64Bits(Builder, Args[i]));
+
+ for (unsigned I = 0, E = WhatToStore.size(); I != E; ++I) {
+ Value *toStore = WhatToStore[I];
+ SmallVector<Value *, 1> BuffOffset;
+ uint offsetVal = toStore->getType()->getIntegerBitWidth() == 32 ? 4 : 8;
+ BuffOffset.push_back(ConstantInt::get(Builder.getInt32Ty(), offsetVal));
+
+ StoreInst *StBuff = Builder.CreateStore(toStore, PtrToStore);
+ LLVM_DEBUG(dbgs() << "inserting store to printf buffer:\n"
+ << *StBuff << '\n');
+ PtrToStore = Builder.CreateGEP(Builder.getInt8Ty(), PtrToStore,
+ BuffOffset, "PrintBuffNextPtr");
+ LLVM_DEBUG(dbgs() << "inserting gep to the printf buffer:\n"
+ << *PtrToStore << '\n');
+ }
+ }
+}
+
Value *llvm::emitAMDGPUPrintfCall(IRBuilder<> &Builder,
- ArrayRef<Value *> Args) {
+ ArrayRef<Value *> Args, bool isBuffered) {
auto NumOps = Args.size();
assert(NumOps >= 1);
auto Fmt = Args[0];
SparseBitVector<8> SpecIsCString;
- locateCStrings(SpecIsCString, Fmt);
+ StringRef FmtStr;
+ locateCStrings(SpecIsCString, Fmt, FmtStr);
+
+ if (isBuffered) {
+ SmallVector<StringData, 8> StringContents;
+ llvm::Module *M = Builder.GetInsertBlock()->getModule();
+ LLVMContext &Ctx = Builder.getContext();
+ auto Int1Ty = Builder.getInt1Ty();
+ auto Int8Ty = Builder.getInt8Ty();
+ auto Int32Ty = Builder.getInt32Ty();
+
+ Value *ArgSize = nullptr;
+ Value *Ptr = callBufferedPrintfStart(Builder, Args, Fmt, FmtStr,
+ SpecIsCString, StringContents, ArgSize);
+
+ // The buffered version still follows OpenCL printf standards for
+ // printf return value, i.e 0 on success, 1 on failure.
+ ConstantPointerNull *zeroIntPtr =
+ ConstantPointerNull::get(cast<PointerType>(Ptr->getType()));
+
+ auto *Cmp = cast<ICmpInst>(Builder.CreateICmpNE(Ptr, zeroIntPtr, ""));
+
+ BasicBlock *End = BasicBlock::Create(Ctx, "end.block",
+ Builder.GetInsertBlock()->getParent());
+ BasicBlock *ArgPush = BasicBlock::Create(
+ Ctx, "argpush.block", Builder.GetInsertBlock()->getParent());
+
+ BranchInst::Create(ArgPush, End, Cmp, Builder.GetInsertBlock());
+ Builder.SetInsertPoint(ArgPush);
+
+ // Create controlDWord and store as the first entry, format as follows
+ // Bit 0 (LSB) -> stream (1 if stderr, 0 if stdout)
+ // Bit 1 -> constant format string (1 if constant)
+ // Bits 2-31 -> size of printf data frame
+ auto CreateControlDWord = M->getOrInsertFunction(
+ StringRef("__ockl_create_control_dword"), Builder.getInt32Ty(),
+ Builder.getInt32Ty(), Int1Ty, Int1Ty);
+ auto valueToStore = Builder.CreateCall(
+ CreateControlDWord,
+ {ArgSize,
+ ConstantInt::get(Int1Ty, !FmtStr.empty() ? 1 : 0, false),
+ ConstantInt::get(Int1Ty, 0, false)});
+ Builder.CreateStore(valueToStore, Ptr);
+
+ Ptr = Builder.CreateGEP(Int8Ty, Ptr,
+ ConstantInt::get(Ctx, APInt(32, 4)));
+
+ // Create MD5 hash for costant format string, push low 64 bits of the
+ // same onto buffer.
+ if (!FmtStr.empty()) {
+ llvm::MD5 Hasher;
+ llvm::MD5::MD5Result Hash;
+ Hasher.update(FmtStr);
+ Hasher.final(Hash);
+
+ // Try sticking to llvm.printf.fmts format, although we are not going to
+ // use the ID and argument size fields while printing,
+ std::string MetadataStr =
+ "0:0:" + llvm::utohexstr(Hash.low(), /*LowerCase=*/true) + "," +
+ FmtStr.str();
+ NamedMDNode *metaD = M->getOrInsertNamedMetadata("llvm.printf.fmts");
+ MDString *fmtStrArray = MDString::get(Ctx, MetadataStr);
+ MDNode *myMD = MDNode::get(Ctx, fmtStrArray);
+ metaD->addOperand(myMD);
+
+ Builder.CreateStore(ConstantInt::get(Builder.getInt64Ty(), Hash.low()),
+ Ptr);
+ Ptr = Builder.CreateGEP(Int8Ty, Ptr,
+ {ConstantInt::get(Int32Ty, 8)});
+ }
+ else {
+ // Include a dummy metadata instance in case of only non constant
+ // format string usage, This might be an absurd usecase but needs to
+ // be done for completeness
+ NamedMDNode *metaD = M->getOrInsertNamedMetadata("llvm.printf.fmts");
+ if(0 == metaD->getNumOperands()) {
+ MDString *fmtStrArray = MDString::get(Ctx, "0:0:deadbeef,\"\"");
+ MDNode *myMD = MDNode::get(Ctx, fmtStrArray);
+ metaD->addOperand(myMD);
+ }
+ }
+
+ // Push The printf arguments onto buffer
+ callBufferedPrintfArgPush(Builder, Args, Ptr, SpecIsCString, StringContents,
+ !FmtStr.empty());
+
+ //End block, returns -1 on failure
+ BranchInst::Create(End, ArgPush);
+ Builder.SetInsertPoint(End);
+ auto toReturn = Builder.CreateSExt(Builder.CreateNot(Cmp),
+ Int32Ty, "printf_result");
+
+ return toReturn;
+ }
auto Desc = callPrintfBegin(Builder, Builder.getIntN(64, 0));
Desc = appendString(Builder, Desc, Fmt, NumOps == 1);
Index: llvm/include/llvm/Transforms/Utils/AMDGPUEmitPrintf.h
===================================================================
--- llvm/include/llvm/Transforms/Utils/AMDGPUEmitPrintf.h
+++ llvm/include/llvm/Transforms/Utils/AMDGPUEmitPrintf.h
@@ -18,7 +18,7 @@
namespace llvm {
-Value *emitAMDGPUPrintfCall(IRBuilder<> &Builder, ArrayRef<Value *> Args);
+Value *emitAMDGPUPrintfCall(IRBuilder<> &Builder, ArrayRef<Value *> Args, bool isBuffered);
} // end namespace llvm
Index: clang/test/Driver/hip-options.hip
===================================================================
--- clang/test/Driver/hip-options.hip
+++ clang/test/Driver/hip-options.hip
@@ -21,6 +21,22 @@
// PTH: "-cc1"{{.*}} "-E" {{.*}}"-fgpu-default-stream=per-thread"
// PTH: "-cc1"{{.*}} "-fgpu-default-stream=per-thread" {{.*}}"-x" "hip-cpp-output"
+// Check -famdgpu-printf-kind=hostcall
+// RUN: %clang -### -famdgpu-printf-kind=hostcall %s -save-temps 2>&1 | FileCheck -check-prefix=HOSTC %s
+// HOSTC: "-cc1" "-triple" "amdgcn-amd-amdhsa"{{.*}}"-famdgpu-printf-kind=hostcall" "-E" {{.*}}
+// HOSTC: "-cc1" "-triple" "amdgcn-amd-amdhsa"{{.*}} "-famdgpu-printf-kind=hostcall" {{.*}}"-x" "hip-cpp-output"
+// HOSTC: "-cc1" "-triple" "amdgcn-amd-amdhsa"{{.*}} "-famdgpu-printf-kind=hostcall" {{.*}}"-x" "ir"
+// HOSTC: "-cc1" "-triple" "x86_64-unknown-linux-gnu"{{.*}} "-E" {{.*}}
+// HOSTC: "-cc1" "-triple" "x86_64-unknown-linux-gnu"{{.*}}"-x" "ir"
+
+// Check -famdgpu-printf-kind=buffered
+// RUN: %clang -### -famdgpu-printf-kind=buffered %s -save-temps 2>&1 | FileCheck -check-prefix=BUFF %s
+// BUFF: "-cc1" "-triple" "amdgcn-amd-amdhsa"{{.*}}"-famdgpu-printf-kind=buffered" "-E" {{.*}}
+// BUFF: "-cc1" "-triple" "amdgcn-amd-amdhsa"{{.*}} "-famdgpu-printf-kind=buffered" {{.*}}"-x" "hip-cpp-output"
+// BUFF: "-cc1" "-triple" "amdgcn-amd-amdhsa"{{.*}} "-famdgpu-printf-kind=buffered" {{.*}}"-x" "ir"
+// BUFF: "-cc1" "-triple" "x86_64-unknown-linux-gnu"{{.*}} "-E" {{.*}}
+// BUFF: "-cc1" "-triple" "x86_64-unknown-linux-gnu"{{.*}}"-x" "ir"
+
// RUN: %clang -### -x hip --target=x86_64-pc-windows-msvc -fms-extensions \
// RUN: -mllvm -amdgpu-early-inline-all=true %s 2>&1 | \
// RUN: FileCheck -check-prefix=MLLVM %s
Index: clang/test/CodeGenHIP/printf_nonhostcall.cpp
===================================================================
--- /dev/null
+++ clang/test/CodeGenHIP/printf_nonhostcall.cpp
@@ -0,0 +1,122 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 2
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -emit-llvm -famdgpu-printf-kind=buffered -fcuda-is-device \
+// RUN: -o - %s | FileCheck --enable-var-scope %s
+
+#define __device__ __attribute__((device))
+
+extern "C" __device__ int printf(const char *format, ...);
+
+// CHECK-LABEL: define dso_local noundef i32 @_Z4foo1v
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: [[S:%.*]] = alloca ptr, align 8, addrspace(5)
+// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
+// CHECK-NEXT: [[S_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[S]] to ptr
+// CHECK-NEXT: store ptr addrspacecast (ptr addrspace(4) @.str to ptr), ptr [[S_ASCAST]], align 8
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[S_ASCAST]], align 8
+// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[S_ASCAST]], align 8
+// CHECK-NEXT: [[TMP2:%.*]] = icmp eq ptr [[TMP0]], null
+// CHECK-NEXT: br i1 [[TMP2]], label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]
+// CHECK: strlen.while:
+// CHECK-NEXT: [[TMP3:%.*]] = phi ptr [ [[TMP0]], [[ENTRY:%.*]] ], [ [[TMP4:%.*]], [[STRLEN_WHILE]] ]
+// CHECK-NEXT: [[TMP4]] = getelementptr i8, ptr [[TMP3]], i64 1
+// CHECK-NEXT: [[TMP5:%.*]] = load i8, ptr [[TMP3]], align 1
+// 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: [[TMP10:%.*]] = add i64 [[TMP9]], 1
+// CHECK-NEXT: br label [[STRLEN_JOIN]]
+// CHECK: strlen.join:
+// CHECK-NEXT: [[TMP11:%.*]] = phi i64 [ [[TMP10]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]
+// CHECK-NEXT: [[TMP12:%.*]] = add i64 [[TMP11]], 7
+// CHECK-NEXT: [[TMP13:%.*]] = and i64 [[TMP12]], 4294967288
+// CHECK-NEXT: [[TMP14:%.*]] = add i64 [[TMP13]], 52
+// CHECK-NEXT: [[TMP15:%.*]] = trunc i64 [[TMP14]] to i32
+// CHECK-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[TMP15]])
+// CHECK-NEXT: [[TMP16:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
+// CHECK-NEXT: br i1 [[TMP16]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]
+// CHECK: end.block:
+// CHECK-NEXT: [[TMP17:%.*]] = xor i1 [[TMP16]], true
+// CHECK-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP17]] to i32
+// CHECK-NEXT: ret i32 [[PRINTF_RESULT]]
+// CHECK: argpush.block:
+// CHECK-NEXT: [[TMP18:%.*]] = call i32 @__ockl_create_control_dword(i32 [[TMP15]], i1 true, i1 false)
+// CHECK-NEXT: store i32 [[TMP18]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
+// CHECK-NEXT: [[TMP19:%.*]] = getelementptr i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
+// CHECK-NEXT: store i64 1107004088646384690, ptr addrspace(1) [[TMP19]], align 8
+// CHECK-NEXT: [[TMP20:%.*]] = getelementptr i8, ptr addrspace(1) [[TMP19]], i32 8
+// CHECK-NEXT: store i64 8, ptr addrspace(1) [[TMP20]], align 8
+// CHECK-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr i8, ptr addrspace(1) [[TMP20]], i32 8
+// CHECK-NEXT: store i64 4614256650576692846, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8
+// CHECK-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8
+// CHECK-NEXT: store i64 8, ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], align 8
+// CHECK-NEXT: [[PRINTBUFFNEXTPTR2:%.*]] = getelementptr i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], i32 8
+// CHECK-NEXT: store i64 4, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], align 8
+// CHECK-NEXT: [[PRINTBUFFNEXTPTR3:%.*]] = getelementptr i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], i32 8
+// CHECK-NEXT: call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) [[PRINTBUFFNEXTPTR3]], ptr [[TMP0]], i64 [[TMP11]], i1 false)
+// CHECK-NEXT: [[PRINTBUFFNEXTPTR4:%.*]] = getelementptr i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR3]], i64 [[TMP13]]
+// CHECK-NEXT: [[TMP21:%.*]] = ptrtoint ptr [[TMP1]] to i64
+// CHECK-NEXT: store i64 [[TMP21]], ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], align 8
+// CHECK-NEXT: [[PRINTBUFFNEXTPTR5:%.*]] = getelementptr i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], i32 8
+// CHECK-NEXT: br label [[END_BLOCK]]
+//
+__device__ int foo1() {
+ const char *s = "hello world";
+ return printf("%.*f %*.*s %p\n", 8, 3.14159, 8, 4, s, s);
+}
+
+__device__ char *dstr;
+
+// CHECK-LABEL: define dso_local noundef i32 @_Z4foo2v
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspacecast (ptr addrspace(1) @dstr to ptr), align 8
+// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr addrspacecast (ptr addrspace(1) @dstr to ptr), align 8
+// CHECK-NEXT: [[TMP2:%.*]] = icmp eq ptr [[TMP0]], null
+// CHECK-NEXT: br i1 [[TMP2]], label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]
+// CHECK: strlen.while:
+// CHECK-NEXT: [[TMP3:%.*]] = phi ptr [ [[TMP0]], [[ENTRY:%.*]] ], [ [[TMP4:%.*]], [[STRLEN_WHILE]] ]
+// CHECK-NEXT: [[TMP4]] = getelementptr i8, ptr [[TMP3]], i64 1
+// CHECK-NEXT: [[TMP5:%.*]] = load i8, ptr [[TMP3]], align 1
+// 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: [[TMP10:%.*]] = add i64 [[TMP9]], 1
+// CHECK-NEXT: br label [[STRLEN_JOIN]]
+// CHECK: strlen.join:
+// CHECK-NEXT: [[TMP11:%.*]] = phi i64 [ [[TMP10]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]
+// CHECK-NEXT: [[TMP12:%.*]] = add i64 [[TMP11]], 7
+// CHECK-NEXT: [[TMP13:%.*]] = and i64 [[TMP12]], 4294967288
+// CHECK-NEXT: [[TMP14:%.*]] = add i64 [[TMP13]], 20
+// CHECK-NEXT: [[TMP15:%.*]] = trunc i64 [[TMP14]] to i32
+// CHECK-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[TMP15]])
+// CHECK-NEXT: [[TMP16:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
+// CHECK-NEXT: br i1 [[TMP16]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]
+// CHECK: end.block:
+// CHECK-NEXT: [[TMP17:%.*]] = xor i1 [[TMP16]], true
+// CHECK-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP17]] to i32
+// CHECK-NEXT: ret i32 [[PRINTF_RESULT]]
+// CHECK: argpush.block:
+// CHECK-NEXT: [[TMP18:%.*]] = call i32 @__ockl_create_control_dword(i32 [[TMP15]], i1 true, i1 false)
+// CHECK-NEXT: store i32 [[TMP18]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
+// CHECK-NEXT: [[TMP19:%.*]] = getelementptr i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
+// CHECK-NEXT: store i64 -9166875625161956257, ptr addrspace(1) [[TMP19]], align 8
+// CHECK-NEXT: [[TMP20:%.*]] = getelementptr i8, ptr addrspace(1) [[TMP19]], i32 8
+// CHECK-NEXT: call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) [[TMP20]], ptr [[TMP0]], i64 [[TMP11]], i1 false)
+// CHECK-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr i8, ptr addrspace(1) [[TMP20]], i64 [[TMP13]]
+// CHECK-NEXT: [[TMP21:%.*]] = ptrtoint ptr [[TMP1]] to i64
+// CHECK-NEXT: store i64 [[TMP21]], ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8
+// CHECK-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8
+// CHECK-NEXT: br label [[END_BLOCK]]
+//
+__device__ int foo2() {
+ return printf("%s %p\n", dstr, dstr);
+}
\ No newline at end of file
Index: clang/lib/Driver/ToolChains/Clang.cpp
===================================================================
--- clang/lib/Driver/ToolChains/Clang.cpp
+++ clang/lib/Driver/ToolChains/Clang.cpp
@@ -4659,7 +4659,22 @@
}
CmdArgs.push_back("-aux-triple");
CmdArgs.push_back(Args.MakeArgString(NormalizedTriple));
+
+ if (JA.isDeviceOffloading(Action::OFK_HIP) &&
+ (types::isHIP(Input.getType()) || types::isLLVMIR(Input.getType()))) {
+ // Device side compilation printf
+ if (Args.getLastArg(options::OPT_famdgpu_printf_kind))
+ CmdArgs.push_back(Args.MakeArgString(
+ "-famdgpu-printf-kind=" +
+ Args.getLastArgValue(options::OPT_famdgpu_printf_kind)));
+ }
}
+
+ // unconditionally claim the pritnf option now to avoid unused diagnostic.
+ // TODO: OpenCL targets will should use this option to switch between
+ // hostcall and buffered printf schemes.
+ if (const Arg *PF = Args.getLastArg(options::OPT_famdgpu_printf_kind))
+ PF->claim();
if (Args.hasFlag(options::OPT_fsycl, options::OPT_fno_sycl, false)) {
CmdArgs.push_back("-fsycl-is-device");
Index: clang/lib/CodeGen/CGGPUBuiltin.cpp
===================================================================
--- clang/lib/CodeGen/CGGPUBuiltin.cpp
+++ clang/lib/CodeGen/CGGPUBuiltin.cpp
@@ -202,7 +202,10 @@
llvm::IRBuilder<> IRB(Builder.GetInsertBlock(), Builder.GetInsertPoint());
IRB.SetCurrentDebugLocation(Builder.getCurrentDebugLocation());
- auto Printf = llvm::emitAMDGPUPrintfCall(IRB, Args);
+
+ bool isBuffered = (CGM.getLangOpts().getAMDGPUPrintfKindVal() ==
+ clang::LangOptions::AMDGPUPrintfKind::Buffered);
+ auto Printf = llvm::emitAMDGPUPrintfCall(IRB, Args, isBuffered);
Builder.SetInsertPoint(IRB.GetInsertBlock(), IRB.GetInsertPoint());
return RValue::get(Printf);
}
Index: clang/include/clang/Driver/Options.td
===================================================================
--- clang/include/clang/Driver/Options.td
+++ clang/include/clang/Driver/Options.td
@@ -1027,6 +1027,13 @@
TargetOpts<"NVPTXUseShortPointers">, DefaultFalse,
PosFlag<SetTrue, [CC1Option], "Use 32-bit pointers for accessing const/local/shared address spaces">,
NegFlag<SetFalse>>;
+def famdgpu_printf_kind : Joined<["-"], "famdgpu-printf-kind=">,
+ HelpText<"specify the printf lowering scheme, value depends on the language being compiled (currently HIP only)">,
+ Flags<[CC1Option]>,
+ Values<"hostcall,buffered">,
+ NormalizedValuesScope<"LangOptions::AMDGPUPrintfKind">,
+ NormalizedValues<["Hostcall", "Buffered"]>,
+ MarshallingInfoEnum<LangOpts<"AMDGPUPrintfKindVal">, "Hostcall">;
def fgpu_default_stream_EQ : Joined<["-"], "fgpu-default-stream=">,
HelpText<"Specify default stream. The default value is 'legacy'. (HIP only)">,
Flags<[CC1Option]>,
Index: clang/include/clang/Basic/LangOptions.h
===================================================================
--- clang/include/clang/Basic/LangOptions.h
+++ clang/include/clang/Basic/LangOptions.h
@@ -380,6 +380,16 @@
IncompleteOnly = 3,
};
+ enum class AMDGPUPrintfKind {
+ /// printf lowering scheme involving hostcalls, currently used by HIP
+ /// programs by default
+ Hostcall = 0,
+
+ /// pritnf lowering scheme involving implicit printf buffers, used by OpenCL
+ /// code by default
+ Buffered = 1,
+ };
+
public:
/// The used language standard.
LangStandard::Kind LangStd;
Index: clang/include/clang/Basic/LangOptions.def
===================================================================
--- clang/include/clang/Basic/LangOptions.def
+++ clang/include/clang/Basic/LangOptions.def
@@ -271,6 +271,7 @@
LANGOPT(GPUDeferDiag, 1, 0, "defer host/device related diagnostic messages for CUDA/HIP")
LANGOPT(GPUExcludeWrongSideOverloads, 1, 0, "always exclude wrong side overloads in overloading resolution for CUDA/HIP")
LANGOPT(OffloadingNewDriver, 1, 0, "use the new driver for generating offloading code.")
+ENUM_LANGOPT(AMDGPUPrintfKindVal, AMDGPUPrintfKind, 2, AMDGPUPrintfKind::Buffered, "printf lowering scheme to be used, hostcall or buffer based")
LANGOPT(SYCLIsDevice , 1, 0, "Generate code for SYCL device")
LANGOPT(SYCLIsHost , 1, 0, "SYCL host compilation")
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
