diff --git a/clang/include/clang/Basic/LangOptions.h b/clang/include/clang/Basic/LangOptions.h --- a/clang/include/clang/Basic/LangOptions.h +++ b/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; diff --git a/clang/include/clang/Basic/LangOptions.def b/clang/include/clang/Basic/LangOptions.def --- a/clang/include/clang/Basic/LangOptions.def +++ b/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") diff --git a/clang/include/clang/Driver/Options.td b/clang/include/clang/Driver/Options.td --- a/clang/include/clang/Driver/Options.td +++ b/clang/include/clang/Driver/Options.td @@ -1027,6 +1027,13 @@ TargetOpts<"NVPTXUseShortPointers">, DefaultFalse, PosFlag, NegFlag>; +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, "Hostcall">; def fgpu_default_stream_EQ : Joined<["-"], "fgpu-default-stream=">, HelpText<"Specify default stream. The default value is 'legacy'. (HIP only)">, Flags<[CC1Option]>, diff --git a/clang/lib/CodeGen/CGGPUBuiltin.cpp b/clang/lib/CodeGen/CGGPUBuiltin.cpp --- a/clang/lib/CodeGen/CGGPUBuiltin.cpp +++ b/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); } diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/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"); diff --git a/clang/test/CodeGenHIP/printf_nonhostcall.cpp b/clang/test/CodeGenHIP/printf_nonhostcall.cpp new file mode 100644 --- /dev/null +++ b/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 diff --git a/clang/test/Driver/hip-options.hip b/clang/test/Driver/hip-options.hip --- a/clang/test/Driver/hip-options.hip +++ b/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 diff --git a/llvm/include/llvm/Transforms/Utils/AMDGPUEmitPrintf.h b/llvm/include/llvm/Transforms/Utils/AMDGPUEmitPrintf.h --- a/llvm/include/llvm/Transforms/Utils/AMDGPUEmitPrintf.h +++ b/llvm/include/llvm/Transforms/Utils/AMDGPUEmitPrintf.h @@ -18,7 +18,7 @@ namespace llvm { -Value *emitAMDGPUPrintfCall(IRBuilder<> &Builder, ArrayRef Args); +Value *emitAMDGPUPrintfCall(IRBuilder<> &Builder, ArrayRef Args, bool isBuffered); } // end namespace llvm diff --git a/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp b/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp --- a/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp +++ b/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 &Args, Value *Fmt, + StringRef &FmtStr, + SparseBitVector<8> &SpecIsCString, + SmallVector &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 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 Args, + Value *PtrToStore, + SparseBitVector<8> &SpecIsCString, + SmallVector &StringContents, + bool isConstFmtStr) { + auto StrIt = StringContents.begin(); + size_t i = isConstFmtStr ? 1 : 0; + for (; i < Args.size(); i++) { + StringRef Str; + SmallVector 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 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 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 Args) { + ArrayRef 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 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(Ptr->getType())); + + auto *Cmp = cast(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);