[AMDGPU] Non hostcall printf support for HIP

[vikramRH]

Summary

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.

[yaxunl]

clang/include/clang/Basic/LangOptions.def

274

This should be a target option like https://reviews.llvm.org/D91546 instead of a language option since it is target specific.

clang/lib/Driver/ToolChains/Clang.cpp

4663–4664

You can check whether target triple is amdgpu instead.

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

416

Please add some comments about the transformations done for buffered and non-buffered cases. Better provide a simple example.

[yaxunl]

clang/include/clang/Driver/Options.td

1036

as a target option, usually it starts with -m and does not contain target name, e.g. -mprintf=.

Better describe allowed values and their effect from users point of view, e.g. hostcall outputs immediately during kernel execution but requires some hardware feature.

[jhuber6]

Where does the runtime implementation of this live? I'm not very familiar with the HIP / hostcall ecosystem.

[arsenm]

clang/include/clang/Basic/LangOptions.def

274

Should be -m option

clang/lib/Driver/ToolChains/Clang.cpp

4673

Capitalize, typo 'pritnf'

4674

Typo 'will should'

Don't really understand the TODO, this should trigger for OpenCL as it is

clang/test/CodeGenHIP/printf_nonhostcall.cpp

5

Do we need a test with -fno-builtin?

69

Doesn't cover the full range of printable types. need some other non-string pointers and different address spaces, some FP promotions, 16 and 64 bit integers

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

259–261

Don't understand the point of the comment, I would assume anything involving analysis of a constant string has ignorable compile time

371–384

Builder.CreateMemCpy

402

Just use the type store size

403

You don't need a SmallVector to push back a single entry

428

Don't need llvm::

430

No auto

458

Do we really need another ockl control variable for this? Why isn't it a parameter? printf=stdout always

473

Don't need llvm::

498

Backwards conditional

[vikramRH]

Handled review comments

[vikramRH]

In D150427#4338334, @jhuber6 wrote:

Where does the runtime implementation of this live? I'm not very familiar with the HIP / hostcall ecosystem.

Currently there is a rocclr review up internally for the runtime side support

[vikramRH]

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

458

There are certain HIP API's such as "hip_assert" that output to stderr. currently such API's are supported via hostcalls. Although this implementation does not currently support the API's ,its kept as an option.

[sameerds]

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

212

you can just say "struct StringData"

218–219

Every argument has a default value. Why not simply assign them as member initializers?

238

Should this say 4?

256

move this closer to first use

313

function argument should use SmallVectorImpl, and not SmallVector

321

This "if (...) { ... }" has a really long body. Can it be moved into a separate function for readability?

324

Try to move declarations to the smallest scope that they are used in. For ReadSize, this seems to be the while loop later.

352

typo: bother

438–439

So we cannot use a buffered printf in HIP?

463

FmtStr.empty() gets checked in multiple places. Can this be captured as one or more boolean variables with meaningful names?

499

Probably not a good choice if it shows up where the user can see it. It's also not very informative. A more descriptive string will be better.

[arsenm]

clang/include/clang/Basic/TargetOptions.h

99 ↗

(On Diff #522954)

Typo pritnf

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

458

Right but the way to handle that would be a parameter for where to output, not an externally set global

[sameerds]

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

184–186

Seems like the only purpose of this change is to make the new StringRef available outside. It will be less confusing to just move the call to getConstantStringInfo() out of this function, and pass the StringRef as an input parameter instead of output. Also, try not to change existing variable names unless it is really relevant. That reduces the number of lines affected by mostly non-functional changes.

223

Needs consistent capitalization in the argument names. Since the rest of the code seems to prefer TitleCase, maybe alignment should start with "A"?

369

Avoid first person comments.

[vikramRH]

Handled Review comments from @sameerds

[vikramRH]

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

438–439

I guess hostcall version returns number of bytes printed as the printf return value, This is not done in Buffered case since we do not print stuff in "real time", instead I have decided to stick with OpenCL standard (return 0 on success and -1 on failure)

[vikramRH]

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

458

I am not clear here, you expect additional inputs to device lib function ?

[sameerds]

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

186

Where is Fmt used in this function? Also, StringRef is passed by value. It's already a lightweight reference, as demonstrated by its name.

212

typedef is not required for this in C++

458

@arsenm, this "control word" written into the buffer. In that sense, it is indeed a parameter passed from device to host as part of the printf packet. It is not yet another global variable.

495–496

else should be on the same line as the closing brace. Please do run clang-format on the entire change.

[vikramRH]

Handled review comments and rebase

[arsenm]

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

428

MathExtras already has alignTo

467

There's a getFalse

492

CreateConstGEP

599

Remove this newline

601

You can use one of the CreateConstGEPs. I also suspect this can be inbounds

607

Remove this newline

[arsenm]

clang/test/CodeGenHIP/printf_nonhostcall.cpp

171–172

Just store the double, don't need the bitcast?

[vikramRH]

Handled review comments and rebase

[vikramRH]

rebase and ping, apologies for the short frequency.

[arsenm]

clang/include/clang/Driver/Options.td

1030

I'm a bit worried this is introducing a side-ABI option not captured in the triple or at least module flags

clang/lib/Driver/ToolChains/Clang.cpp

4666

Braces

clang/test/CodeGenHIP/printf_nonhostcall.cpp

65

Can directly store the pointer, don't ptrtoint? Should have some other pointer address spaces tested, the spec is quiet on how %p is supposed to work with different sized pointers

229

Need some vector tests, especially 3 x vectors

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

369

You have better alignment info than Align(1)

419

Don't need = ""

420

Move last

432

ArrayRef should be passed by value

443

No auto

443–446

No auto

458

I'm not following why this part requires introducing another call into device libs. It's expanding this not-really defined API. I'd expect this to be a trivial function we can expand inline here and write directly into the parameter buffer?

502

I suppose you can't access the AMDGPU enum from Utils. You could use DL.getDefaultGlobalsAddrSpace

589

Don't see why isDoubleTy would be special cased here and not part of fitArgInto64Bits

603

This should probably be getTypeAllocSize

[arsenm]

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

419

Can't you just use the raw StringRef out of getConstantStringInfo

[vikramRH]

Handled most review comments

[vikramRH]

clang/test/CodeGenHIP/printf_nonhostcall.cpp

65

Added few more address space cases in foo2(), do you think we need more ?

229

vectors are not yet supported as in the hostcall case since the code currently runs for HIP only. I plan to port the OpenCL lowering also here along with hostcall support . we would need to extend hostcall and this implementation to support vectors as part of the porting activity.

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

419

The structure just caches the string calculated from getConstantStringInfo

458

I started the implementation keeping in mind the device side asserts currenlty supported via hostcall. They use explicit calls to such device lib functions in their definitions. I could not support them now but these would be required if I ever decide to. also I never ended up inlining this call. do you think its really necessary ?

589

The redundant cast instructions and this case were due to function "fitArgInto64Bits", which I reused from hostcall Implementation and did not want to modify. This dependency has been removed now and I have inlined the code

[vikramRH]

clang/include/clang/Driver/Options.td

1030

I did not understand the concern here. I agree The usage of option itself is not the most robust solution and it would have been better if we could do this without user intervention. but I do not see ways to do that now (atleast not immediately).

[vikramRH]

Few additional changes,

1. reflect printf-kind in module flags metadata
2. Test cases for the change

[vikramRH]

new line at the end of test file

[arsenm]

clang/include/clang/Driver/Options.td

1030

This introduces a way to produce binaries that will silently not work correctly depending on the host system. Without somehow tracking this property in the binary, there's no way to diagnose trying to load a binary on an incompatible system.

clang/test/CodeGenHIP/printf_nonhostcall.cpp

229

Something still needs to happen for vectors, it can't just crash. The different types of vectors may appear regardless of language

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

419

But you can just use returned StringRef for that. The actual constant contents persists in the LLVMContext, you don't need to copy it

458

But that implementation is in the library itself? What does the function actually do?

I think we're better off the more defined/documented ABI we have in terms of size-and-offset from base pointer than call into unstable library calls.

Can we get some documentation on the ABI in AMDGPUUsage?

490

Builder.getInt64(Hash.low())

[arsenm]

Should also be mentioned in the release notes

[vikramRH]

Handled few more review comments

[vikramRH]

clang/test/CodeGenHIP/printf_nonhostcall.cpp

229

I have forced "-Werror=format-invalid-specifier" with printf option which should be a sufficient condition for vectors I guess ? (this causes clang to break out earlier with an error)

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

458

The library function itself just handles creation of controlDWord via minor bit manipulations. nevertheless Inlining is pretty simple and im fine either way. Call has been expanded.

[arsenm]

clang/test/CodeGenHIP/printf_nonhostcall.cpp

138

Also half

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

579

I don't follow this alignment logic, but I guess the optimizer will figure it out later anyway so I suppose the Align(1) was fine.

Also should use Align.

581

The source align is known

594–595

isIntegerTy(32).

I also do not understand why only 32-bit integers would be promoted to 64-bit (or why this would be zext). This doesn't match varargs ABI handling, where everything smaller would be promoted to i32.

[arsenm]

Also should be noted in the release notes

[vikramRH]

Further review comments

[vikramRH]

clang/test/CodeGenHIP/printf_nonhostcall.cpp

138

C++ default arg promotions does not seem to list float16 case and consequently clang does not handle it. I have handled this case by promoting to double currenlty. Clang however still emits a warning when used with a %f conv specifier

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

594–595

Right, But default promotions would have happened already, this additional promotion is due to runtime processing requirement

[vikramRH]

updated formating

[arsenm]

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

594–595

Is this documented somewhere? If it's promote everything to i64, I'd prefer to handle all types rather than special casing 32

[vikramRH]

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

594–595

That's precisely what the function "fitArgInto64Bits()" did. We eliminated it due to some redundant instructions generated. int32 was the only case (apart from float16 now) that required this special casting.

[sameerds]

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

594–595

I think the problem is that the check for "isIntergerTy(32)" makes it look like it is some sort of special case. What we really mean is that if the type is an integer with width less than 64 bits, then zext it with padding to fit the 64-bit store into the buffer. We can do that with just "isIntegerTy()".

Separately, we know from auto promotion of varargs that the integer can only be of size 32 or 64. We can assert that instead of assuming it in the code.

[vikramRH]

Reafactored non string arg handling into a seperate function with additional asserts

[arsenm]

clang/docs/ReleaseNotes.rst

589 ↗

(On Diff #529551)

I thought the default was buffered, such that it would always work. Defaults that always work are better

Also, amdgpu-arch should probably learn to report if pcie atomics work to improve the autodetect

clang/lib/Driver/ToolChains/Clang.cpp

4670–4672

I don't see why we would special case this

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

569

It costs nothing to handle all types < 64. For larger just return the original value

576

Should avoid special casing this too, there's also bfloat to consider at a minimum

579

Safest to just return the original value

621

Range loop, index is unused

[arsenm]

clang/test/CodeGenHIP/printf_nonhostcall.cpp

241

Test _BitInt for small and odd types, plus i128

[vikramRH]

Handled last set of review comments from @arsenm, would be willing to handle any new findings/concerns via additional patches

[arsenm]

lgtm with nit

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

488

BufSize += std::max(getTypeAllocSize(), 8)

577

maybe add a constrained fp run line to a test

clang/include/clang/Basic/LangOptions.h

 //===- LangOptions.h - C Language Family Language Options -------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
enum class StrictFlexArraysLevelKind {
     /// Any trailing array member of undefined, 0, or 1 size is a FAM.
     OneZeroOrIncomplete = 1,
     /// Any trailing array member of undefined or 0 size is a FAM.
     ZeroOrIncomplete = 2,
     /// Any trailing array member of undefined size is a FAM.
     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;
 
   /// Set of enabled sanitizers.
   SanitizerSet Sanitize;
   /// Is at least one coverage instrumentation type enabled.
   bool SanitizeCoverage = false;

clang/include/clang/Basic/LangOptions.def

 LANGOPT(CUDAHostDeviceConstexpr, 1, 1, "treating unattributed constexpr functions as __host__ __device__")
 LANGOPT(CUDADeviceApproxTranscendentals, 1, 0, "using approximate transcendental functions")
 LANGOPT(GPURelocatableDeviceCode, 1, 0, "generate relocatable device code")
 LANGOPT(GPUAllowDeviceInit, 1, 0, "allowing device side global init functions for HIP")
 LANGOPT(GPUMaxThreadsPerBlock, 32, 1024, "default max threads per block for kernel launch bounds for HIP")
 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")
 ENUM_LANGOPT(SYCLVersion  , SYCLMajorVersion, 2, SYCL_None, "Version of the SYCL standard used")
 
 LANGOPT(HIPUseNewLaunchAPI, 1, 0, "Use new kernel launching API for HIP")
 
 LANGOPT(SizedDeallocation , 1, 0, "sized deallocation")

clang/include/clang/Driver/Options.td

defm gpu_rdc : BoolFOption<"gpu-rdc",
   PosFlag<SetTrue, [CC1Option], "Generate relocatable device code, also known as separate compilation mode">,
   NegFlag<SetFalse>>;
 def : Flag<["-"], "fcuda-rdc">, Alias<fgpu_rdc>;
 def : Flag<["-"], "fno-cuda-rdc">, Alias<fno_gpu_rdc>;
 defm cuda_short_ptr : BoolFOption<"cuda-short-ptr",
   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]>,
   Values<"legacy,per-thread">,
   NormalizedValuesScope<"LangOptions::GPUDefaultStreamKind">,
   NormalizedValues<["Legacy", "PerThread"]>,
   MarshallingInfoEnum<LangOpts<"GPUDefaultStream">, "Legacy">;
 def rocm_path_EQ : Joined<["--"], "rocm-path=">, Group<i_Group>,

clang/lib/CodeGen/CGGPUBuiltin.cpp

 //===------ CGGPUBuiltin.cpp - Codegen for GPU builtins -------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
for (const auto &A : CallArgs) {
     }
 
     llvm::Value *Arg = A.getRValue(*this).getScalarVal();
     Args.push_back(Arg);
   }
 
   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);
 }
 
 RValue CodeGenFunction::EmitOpenMPDevicePrintfCallExpr(const CallExpr *E) {
   assert(getTarget().getTriple().isNVPTX() ||
          getTarget().getTriple().isAMDGCN());
   return EmitDevicePrintfCallExpr(E, this, GetOpenMPVprintfDeclaration(CGM),
                                   true);
 }

clang/lib/Driver/ToolChains/Clang.cpp

 //===-- Clang.cpp - Clang+LLVM ToolChain Implementations --------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
else {
         if (CTC && CTC->CudaInstallation.version() != CudaVersion::UNKNOWN)
           CmdArgs.push_back(Args.MakeArgString(
               Twine("-target-sdk-version=") +
               CudaVersionToString(CTC->CudaInstallation.version())));
       }
     }
     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");
 
     if (Arg *A = Args.getLastArg(options::OPT_sycl_std_EQ)) {
       A->render(Args, CmdArgs);
     } else {
       // Ensure the default version in SYCL mode is 2020.

clang/test/CodeGenHIP/printf_nonhostcall.cpp

+// 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

clang/test/Driver/hip-options.hip

 // Check -fgpu-default-stream=per-thread.
 // RUN: %clang -### -nogpuinc -nogpulib -fgpu-default-stream=per-thread \
 // RUN:   %s -save-temps 2>&1 | FileCheck -check-prefix=PTH %s
 // PTH: "-cc1"{{.*}} "-E" {{.*}}"-fgpu-default-stream=per-thread"
 // PTH: "-cc1"{{.*}} "-fgpu-default-stream=per-thread" {{.*}}"-x" "hip-cpp-output"
 // 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
 // MLLVM-NOT: "-mllvm"{{.*}}"-amdgpu-early-inline-all=true"{{.*}}"-mllvm"{{.*}}"-amdgpu-early-inline-all=true"
 
 // RUN: %clang -### -Xarch_device -g -nogpulib --cuda-gpu-arch=gfx900 \
 // RUN:   -Xarch_device -fcf-protection=branch -Xarch_device -mllvm=--inline-threshold=100 \
 // RUN:   --cuda-gpu-arch=gfx906  %s 2>&1 | FileCheck -check-prefix=DEV %s

llvm/include/llvm/Transforms/Utils/AMDGPUEmitPrintf.h

 //===- AMDGPUEmitPrintf.h ---------------------------------------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // Utility function to lower a printf call into a series of device
 // library calls on the AMDGPU target.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef LLVM_TRANSFORMS_UTILS_AMDGPUEMITPRINTF_H
 #define LLVM_TRANSFORMS_UTILS_AMDGPUEMITPRINTF_H
 
 #include "llvm/IR/IRBuilder.h"
 
 namespace llvm {
 
-Value *emitAMDGPUPrintfCall(IRBuilder<> &Builder, ArrayRef<Value *> Args);
+Value *emitAMDGPUPrintfCall(IRBuilder<> &Builder, ArrayRef<Value *> Args, bool isBuffered);
 
 } // end namespace llvm
 
 #endif // LLVM_TRANSFORMS_UTILS_AMDGPUEMITPRINTF_H

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

 //===- AMDGPUEmitPrintf.cpp -----------------------------------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // Utility function to lower a printf call into a series of device
 // library calls on the AMDGPU target.
 //
 // WARNING: This file knows about certain library functions. It recognizes them
 // by name, and hardwires knowledge of their semantics.
 //
 //===----------------------------------------------------------------------===//
 
 #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;
 
 #define DEBUG_TYPE "amdgpu-emit-printf"
 
 static Value *fitArgInto64Bits(IRBuilder<> &Builder, Value *Arg) {
   auto Int64Ty = Builder.getInt64Ty();
   auto Ty = Arg->getType();
static Value *processArg(IRBuilder<> &Builder, Value *Desc, Value *Arg,
   // If the format specifies a string but the argument is not, the frontend will
   // have printed a warning. We just rely on undefined behaviour and send the
   // argument anyway.
   return appendArg(Builder, Desc, Arg, IsLast);
 }
 
 // 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";
   size_t SpecPos = 0;
   // 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;
     ++ArgIdx;
   }
 }
 
+// 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);
 
   // FIXME: This invokes hostcall once for each argument. We can pack up to
   // seven scalar printf arguments in a single hostcall. See the signature of
   // callAppendArgs().
   for (unsigned int i = 1; i != NumOps; ++i) {
     bool IsLast = i == NumOps - 1;
     bool IsCString = SpecIsCString.test(i);
     Desc = processArg(Builder, Desc, Args[i], IsCString, IsLast);
   }
 
   return Builder.CreateTrunc(Desc, Builder.getInt32Ty());
 }