[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 ↗

(On Diff #521570)

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

4700–4701

You can check whether target triple is amdgpu instead.

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

208–209

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

1042

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 ↗

(On Diff #521570)

Should be -m option

clang/lib/Driver/ToolChains/Clang.cpp

4711

Capitalize, typo 'pritnf'

4712

Typo 'will should'

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

clang/test/CodeGenHIP/printf_nonhostcall.cpp

6

Do we need a test with -fno-builtin?

70

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

257–259

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

369–382

Builder.CreateMemCpy

400

Just use the type store size

401

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

422

Don't need llvm::

424

No auto

452

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

467

Don't need llvm::

492

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

452

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

210

you can just say "struct StringData"

216–217

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

236

Should this say 4?

254

move this closer to first use

311

function argument should use SmallVectorImpl, and not SmallVector

319

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

322

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

350

typo: bother

432–433

So we cannot use a buffered printf in HIP?

457

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

493

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

Typo pritnf

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

452

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–185

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.

221

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

367

Avoid first person comments.

[vikramRH]

Handled Review comments from @sameerds

[vikramRH]

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

432–433

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

452

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

[sameerds]

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

184–185

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

210

typedef is not required for this in C++

452

@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.

489–490

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

220

MathExtras already has alignTo

391

Remove this newline

393

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

399

Remove this newline

461

There's a getFalse

486

CreateConstGEP

[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

1036

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

4703

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

211

Don't need = ""

212

Move last

224

ArrayRef should be passed by value

235

No auto

235–238

No auto

294

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

367

You have better alignment info than Align(1)

381

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

395

This should probably be getTypeAllocSize

452

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?

[arsenm]

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

211

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

211

The structure just caches the string calculated from getConstantStringInfo

381

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

452

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 ?

[vikramRH]

clang/include/clang/Driver/Options.td

1036

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

1036

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

211

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

452

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?

484

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

452

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

371

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.

373

The source align is known

386–387

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

386–387

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

386–387

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

386–387

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

386–387

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

4707–4709

I don't see why we would special case this

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

361

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

368

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

371

Safest to just return the original value

413

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

280

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

369

maybe add a constrained fp run line to a test

clang/include/clang/Basic/TargetOptions.h

 //===--- TargetOptions.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
 //
 //===----------------------------------------------------------------------===//
 ///
enum CodeObjectVersionKind {
     COV_2 = 200,
     COV_3 = 300,
     COV_4 = 400,
     COV_5 = 500,
   };
   /// \brief Code object version for AMDGPU.
   CodeObjectVersionKind CodeObjectVersion = CodeObjectVersionKind::COV_None;
 
+  /// \brief Enumeration values for AMDGPU printf lowering scheme
+  enum class AMDGPUPrintfKind {
+    /// printf lowering scheme involving hostcalls, currently used by HIP
+    /// programs by default
+    Hostcall = 0,
+
+    /// printf lowering scheme involving implicit printf buffers,
+    Buffered = 1,
+  };
+
+  /// \brief AMDGPU Printf lowering scheme
+  AMDGPUPrintfKind AMDGPUPrintfKindVal = AMDGPUPrintfKind::Hostcall;
+
   // The code model to be used as specified by the user. Corresponds to
   // CodeModel::Model enum defined in include/llvm/Support/CodeGen.h, plus
   // "default" for the case when the user has not explicitly specified a
   // code model.
   std::string CodeModel;
 
   /// The version of the SDK which was used during the compilation.
   /// The option is used for two different purposes:

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 mprintf_kind_EQ : Joined<["-"], "mprintf-kind=">, Group<m_Group>,
+  HelpText<"Specify the printf lowering scheme (AMDGPU only), allowed values are "
+  "\"hostcall\"(printing happens during kernel execution, this scheme "
+  "relies on hostcalls which require system to support pcie atomics) "
+  "and \"buffered\"(printing happens after all kernel threads exit"
+  "this uses a printf buffer and does not rely on pcie atomic support)">,
+  Flags<[CC1Option]>,
+  Values<"hostcall,buffered">,
+  NormalizedValuesScope<"TargetOptions::AMDGPUPrintfKind">,
+  NormalizedValues<["Hostcall", "Buffered"]>,
+  MarshallingInfoEnum<TargetOpts<"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.getTarget().getTargetOpts().AMDGPUPrintfKindVal ==
+                     clang::TargetOptions::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 {
         // recent CUDA headers rely on that.
         // https://github.com/llvm/llvm-project/issues/58410
         if (CTC->CudaInstallation.version() >= CudaVersion::CUDA_90)
           CmdArgs.push_back("-fcuda-allow-variadic-functions");
       }
     }
     CmdArgs.push_back("-aux-triple");
     CmdArgs.push_back(Args.MakeArgString(NormalizedTriple));
+
+    if (JA.isDeviceOffloading(Action::OFK_HIP) &&
+        getToolChain().getTriple().isAMDGPU()) {
+      // Device side compilation printf
+      if (Args.getLastArg(options::OPT_mprintf_kind_EQ)) {
+        CmdArgs.push_back(Args.MakeArgString(
+            "-mprintf-kind=" +
+            Args.getLastArgValue(options::OPT_mprintf_kind_EQ)));
+      }
     }
+  }
+
+  // Unconditionally claim the printf option now to avoid unused diagnostic.
+  if (const Arg *PF = Args.getLastArg(options::OPT_mprintf_kind_EQ))
+    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 -mprintf-kind=buffered -fcuda-is-device \
+// RUN:   -o - %s | FileCheck --enable-var-scope %s
+
+#define __device__ __attribute__((device))
+#define __shared__ __attribute__((shared))
+#define __constant__ __attribute__((constant))
+
+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 @__printf_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 inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
+// CHECK-NEXT:    store i64 1107004088646384690, ptr addrspace(1) [[TMP19]], align 8
+// CHECK-NEXT:    [[TMP20:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP19]], i32 8
+// CHECK-NEXT:    store i64 8, ptr addrspace(1) [[TMP20]], align 8
+// CHECK-NEXT:    [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP20]], i32 8
+// CHECK-NEXT:    store double 3.141590e+00, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8
+// CHECK-NEXT:    [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8
+// CHECK-NEXT:    store i64 8, ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], align 8
+// CHECK-NEXT:    [[PRINTBUFFNEXTPTR2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], i32 8
+// CHECK-NEXT:    store i64 4, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], align 8
+// CHECK-NEXT:    [[PRINTBUFFNEXTPTR3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], i32 8
+// CHECK-NEXT:    call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 8 [[PRINTBUFFNEXTPTR3]], ptr align 1 [[TMP0]], i64 [[TMP11]], i1 false)
+// CHECK-NEXT:    [[PRINTBUFFNEXTPTR4:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR3]], i64 [[TMP13]]
+// CHECK-NEXT:    store ptr [[TMP1]], ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], align 8
+// CHECK-NEXT:    [[PRINTBUFFNEXTPTR5:%.*]] = getelementptr inbounds 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;
+__device__ const
+// CHECK-LABEL: define dso_local noundef i32 @_Z4foo2v
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT:    [[LCVAL:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT:    [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
+// CHECK-NEXT:    [[LCVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[LCVAL]] 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]], 36
+// 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 @__printf_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 inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
+// CHECK-NEXT:    store i64 7257695813269076350, ptr addrspace(1) [[TMP19]], align 8
+// CHECK-NEXT:    [[TMP20:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP19]], i32 8
+// CHECK-NEXT:    call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 8 [[TMP20]], ptr align 1 [[TMP0]], i64 [[TMP11]], i1 false)
+// CHECK-NEXT:    [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP20]], i64 [[TMP13]]
+// CHECK-NEXT:    store ptr [[TMP1]], ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8
+// CHECK-NEXT:    [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8
+// CHECK-NEXT:    store ptr addrspacecast (ptr addrspace(3) @_ZZ4foo2vE5shval to ptr), ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], align 8
+// CHECK-NEXT:    [[PRINTBUFFNEXTPTR2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], i32 8
+// CHECK-NEXT:    store ptr [[LCVAL_ASCAST]], ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], align 8
+// CHECK-NEXT:    [[PRINTBUFFNEXTPTR3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], i32 8
+// CHECK-NEXT:    br label [[END_BLOCK]]
+//
+__device__ int foo2() {
+  __shared__ int shval;
+  int lcval;
+  return printf("%s %p %p %p\n", dstr, dstr, &shval, &lcval);
+}
+
+__device__ unsigned short g = 30;
+__device__ unsigned long n = 30;
+
+__device__ float f1 = 3.14f;
+__device__ double f2 = 2.71828;
+
+// CHECK-LABEL: define dso_local noundef i32 @_Z4foo3v
+// 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:    store i32 25, ptr addrspacecast (ptr addrspace(3) @_ZZ4foo3vE1s to ptr), align 4
+// CHECK-NEXT:    [[TMP0:%.*]] = load i32, ptr addrspacecast (ptr addrspace(3) @_ZZ4foo3vE1s to ptr), align 4
+// CHECK-NEXT:    [[TMP1:%.*]] = load i16, ptr addrspacecast (ptr addrspace(1) @g to ptr), align 2
+// CHECK-NEXT:    [[CONV:%.*]] = zext i16 [[TMP1]] to i32
+// CHECK-NEXT:    [[TMP2:%.*]] = load i64, ptr addrspacecast (ptr addrspace(1) @n to ptr), align 8
+// CHECK-NEXT:    [[TMP3:%.*]] = load float, ptr addrspacecast (ptr addrspace(1) @f1 to ptr), align 4
+// CHECK-NEXT:    [[CONV1:%.*]] = fpext float [[TMP3]] to double
+// CHECK-NEXT:    [[TMP4:%.*]] = load double, ptr addrspacecast (ptr addrspace(1) @f2 to ptr), align 8
+// CHECK-NEXT:    [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 60)
+// CHECK-NEXT:    [[TMP5:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
+// CHECK-NEXT:    br i1 [[TMP5]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]
+// CHECK:       end.block:
+// CHECK-NEXT:    [[TMP6:%.*]] = xor i1 [[TMP5]], true
+// CHECK-NEXT:    [[PRINTF_RESULT:%.*]] = sext i1 [[TMP6]] to i32
+// CHECK-NEXT:    ret i32 [[PRINTF_RESULT]]
+// CHECK:       argpush.block:
+// CHECK-NEXT:    [[TMP7:%.*]] = call i32 @__printf_control_dword(i32 60, i1 true, i1 false)
+// CHECK-NEXT:    store i32 [[TMP7]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
+// CHECK-NEXT:    [[TMP8:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
+// CHECK-NEXT:    store i64 2197983583858494848, ptr addrspace(1) [[TMP8]], align 8
+// CHECK-NEXT:    [[TMP9:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP8]], i32 8
+// CHECK-NEXT:    [[TMP10:%.*]] = zext i32 [[TMP0]] to i64
+// CHECK-NEXT:    store i64 [[TMP10]], ptr addrspace(1) [[TMP9]], align 8
+// CHECK-NEXT:    [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP9]], i32 8
+// CHECK-NEXT:    store ptr addrspacecast (ptr addrspace(3) @_ZZ4foo3vE1s to ptr), ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8
+// CHECK-NEXT:    [[PRINTBUFFNEXTPTR2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8
+// CHECK-NEXT:    [[TMP11:%.*]] = zext i32 [[CONV]] to i64
+// CHECK-NEXT:    store i64 [[TMP11]], ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], align 8
+// CHECK-NEXT:    [[PRINTBUFFNEXTPTR3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], i32 8
+// CHECK-NEXT:    store i64 [[TMP2]], ptr addrspace(1) [[PRINTBUFFNEXTPTR3]], align 8
+// CHECK-NEXT:    [[PRINTBUFFNEXTPTR4:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR3]], i32 8
+// CHECK-NEXT:    store double [[CONV1]], ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], align 8
+// CHECK-NEXT:    [[PRINTBUFFNEXTPTR5:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR4]], i32 8
+// CHECK-NEXT:    store double [[TMP4]], ptr addrspace(1) [[PRINTBUFFNEXTPTR5]], align 8
+// CHECK-NEXT:    [[PRINTBUFFNEXTPTR6:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR5]], i32 8
+// CHECK-NEXT:    br label [[END_BLOCK]]
+//
+__device__ int foo3() {
+  __shared__ int s;
+  s = 25;
+  return printf("Random values: %d,%p,%hd,%ld,%f,%f\n",s, &s, g, n, f1, f2);
+}
+
+//A non trivial case,
+// CHECK-LABEL: define dso_local noundef i32 @_Z4foo4v
+// 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.4 to ptr), ptr [[S_ASCAST]], align 8
+// CHECK-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[S_ASCAST]], align 8
+// CHECK-NEXT:    [[TMP1:%.*]] = icmp eq ptr [[TMP0]], null
+// CHECK-NEXT:    br i1 [[TMP1]], label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]
+// CHECK:       strlen.while:
+// CHECK-NEXT:    [[TMP2:%.*]] = phi ptr [ [[TMP0]], [[ENTRY:%.*]] ], [ [[TMP3:%.*]], [[STRLEN_WHILE]] ]
+// CHECK-NEXT:    [[TMP3]] = getelementptr i8, ptr [[TMP2]], i64 1
+// CHECK-NEXT:    [[TMP4:%.*]] = load i8, ptr [[TMP2]], align 1
+// CHECK-NEXT:    [[TMP5:%.*]] = icmp eq i8 [[TMP4]], 0
+// CHECK-NEXT:    br i1 [[TMP5]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]]
+// CHECK:       strlen.while.done:
+// CHECK-NEXT:    [[TMP6:%.*]] = ptrtoint ptr [[TMP0]] to i64
+// CHECK-NEXT:    [[TMP7:%.*]] = ptrtoint ptr [[TMP2]] to i64
+// CHECK-NEXT:    [[TMP8:%.*]] = sub i64 [[TMP7]], [[TMP6]]
+// CHECK-NEXT:    [[TMP9:%.*]] = add i64 [[TMP8]], 1
+// CHECK-NEXT:    br label [[STRLEN_JOIN]]
+// CHECK:       strlen.join:
+// CHECK-NEXT:    [[TMP10:%.*]] = phi i64 [ [[TMP9]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]
+// CHECK-NEXT:    [[TMP11:%.*]] = add i64 [[TMP10]], 7
+// CHECK-NEXT:    [[TMP12:%.*]] = and i64 [[TMP11]], 4294967288
+// CHECK-NEXT:    [[TMP13:%.*]] = add i64 [[TMP12]], 12
+// CHECK-NEXT:    [[TMP14:%.*]] = trunc i64 [[TMP13]] to i32
+// CHECK-NEXT:    [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[TMP14]])
+// CHECK-NEXT:    [[TMP15:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
+// CHECK-NEXT:    br i1 [[TMP15]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]
+// CHECK:       end.block:
+// CHECK-NEXT:    [[TMP16:%.*]] = xor i1 [[TMP15]], true
+// CHECK-NEXT:    [[PRINTF_RESULT:%.*]] = sext i1 [[TMP16]] to i32
+// CHECK-NEXT:    ret i32 [[PRINTF_RESULT]]
+// CHECK:       argpush.block:
+// CHECK-NEXT:    [[TMP17:%.*]] = call i32 @__printf_control_dword(i32 [[TMP14]], i1 false, i1 false)
+// CHECK-NEXT:    store i32 [[TMP17]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
+// CHECK-NEXT:    [[TMP18:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
+// CHECK-NEXT:    call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 4 [[TMP18]], ptr align 1 [[TMP0]], i64 [[TMP10]], i1 false)
+// CHECK-NEXT:    [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP18]], i64 [[TMP12]]
+// CHECK-NEXT:    store i64 10, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8
+// CHECK-NEXT:    [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8
+// CHECK-NEXT:    br label [[END_BLOCK]]
+//
+__device__ int foo4() {
+  const char* s = "format str%d";
+  return printf(s, 10);
+}

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 -mprintf-kind=hostcall
+// RUN: %clang -### -mprintf-kind=hostcall  %s -save-temps 2>&1 | FileCheck -check-prefix=HOSTC %s
+// HOSTC: "-cc1" "-triple" "amdgcn-amd-amdhsa"{{.*}}"-mprintf-kind=hostcall" "-E" {{.*}}
+// HOSTC: "-cc1" "-triple" "amdgcn-amd-amdhsa"{{.*}} "-mprintf-kind=hostcall" {{.*}}"-x" "hip-cpp-output"
+// HOSTC: "-cc1" "-triple" "amdgcn-amd-amdhsa"{{.*}} "-mprintf-kind=hostcall" {{.*}}"-x" "ir"
+// HOSTC: "-cc1" "-triple" "x86_64-unknown-linux-gnu"{{.*}} "-E" {{.*}}
+// HOSTC: "-cc1" "-triple" "x86_64-unknown-linux-gnu"{{.*}}"-x" "ir"
+
+// Check -mprintf-kind=buffered
+// RUN: %clang -### -mprintf-kind=buffered  %s -save-temps 2>&1 | FileCheck -check-prefix=BUFF %s
+// BUFF: "-cc1" "-triple" "amdgcn-amd-amdhsa"{{.*}}"-mprintf-kind=buffered" "-E" {{.*}}
+// BUFF: "-cc1" "-triple" "amdgcn-amd-amdhsa"{{.*}} "-mprintf-kind=buffered" {{.*}}"-x" "hip-cpp-output"
+// BUFF: "-cc1" "-triple" "amdgcn-amd-amdhsa"{{.*}} "-mprintf-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"
+#include "llvm/Support/MathExtras.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())
-    return;
-
+static void locateCStrings(SparseBitVector<8> &BV, StringRef Str) {
   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] == '%') {
       SpecPos += 2;
       continue;
     }
     auto SpecEnd = Str.find_first_of(ConvSpecifiers, SpecPos);
     if (SpecEnd == StringRef::npos)
       return;
     auto Spec = Str.slice(SpecPos, SpecEnd + 1);
     ArgIdx += Spec.count('*');
     if (Str[SpecEnd] == 's') {
       BV.set(ArgIdx);
     }
     SpecPos = SpecEnd + 1;
     ++ArgIdx;
   }
 }
 
-Value *llvm::emitAMDGPUPrintfCall(IRBuilder<> &Builder,
-                                  ArrayRef<Value *> Args) {
+// helper struct to package the string related data
+struct StringData {
+  std::string Str;
+  Value *RealSize = nullptr;
+  Value *AlignedSize = nullptr;
+  bool isConst = true;
+
+  StringData(std::string str, Value *RS, Value *AS, bool IC)
+      : Str(str), RealSize(RS), AlignedSize(AS), isConst(IC) {}
+};
+
+// 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,
+    bool isConstFmtStr, SparseBitVector<8> &SpecIsCString,
+    SmallVectorImpl<StringData> &StringContents, Value *&ArgSize) {
+  Module *M = Builder.GetInsertBlock()->getModule();
+  Value *NonConstStrLen = nullptr;
+  Value *LenWithNull = nullptr;
+  Value *LenWithNullAligned = nullptr;
+  Value *TempAdd = nullptr;
+
+  // First 4 bytes to be reserved for control dword
+  size_t BufSize = 4;
+  if (isConstFmtStr)
+    // First 8 bytes of MD5 hash
+    BufSize += 8;
+  else {
+    LenWithNull = getStrlenWithNull(Builder, Fmt);
+
+    // Align the computed length to next 8 byte boundary
+    TempAdd = Builder.CreateAdd(LenWithNull,
+                                ConstantInt::get(LenWithNull->getType(), 7U));
+    NonConstStrLen = Builder.CreateAnd(
+        TempAdd, ConstantInt::get(LenWithNull->getType(), ~7U));
+
+    StringContents.push_back(
+        StringData("", LenWithNull, NonConstStrLen, false));
+  }
+
+  for (size_t i = 1; i < Args.size(); i++) {
+    if (SpecIsCString.test(i)) {
+      StringRef ArgStr;
+      if (getConstantStringInfo(Args[i], ArgStr)) {
+        auto alignedLen = alignTo(ArgStr.size() + 1, 8);
+        StringContents.push_back(StringData(
+            (ArgStr.str() + '\0'),
+            /*RealSize*/ nullptr, /*AlignedSize*/ nullptr, /*isConst*/ true));
+        BufSize += alignedLen;
+      } else {
+        LenWithNull = getStrlenWithNull(Builder, Args[i]);
+
+        // Align the computed length to next 8 byte boundary
+        TempAdd = Builder.CreateAdd(
+            LenWithNull, ConstantInt::get(LenWithNull->getType(), 7U));
+        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("", LenWithNull, LenWithNullAligned, false));
+      }
+    } 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(M->getDataLayout().getDefaultGlobalsAddressSpace());
+  FunctionType *FTy_alloc = FunctionType::get(I8Ptr, Tys_alloc, false);
+  auto PrintfAllocFn =
+      M->getOrInsertFunction(StringRef("__printf_alloc"), FTy_alloc, Attr);
+
+  return Builder.CreateCall(PrintfAllocFn, Alloc_args, "printf_alloc_fn");
+}
+
+// Prepare constant string argument to push onto the buffer
+static void processConstantStringArg(StringData *SD, IRBuilder<> &Builder,
+                                     SmallVectorImpl<Value *> &WhatToStore) {
+  StringRef Str = SD->Str;
+
+  DataExtractor Extractor(Str, /*IsLittleEndian=*/true, 8);
+  DataExtractor::Cursor Offset(0);
+  while (Offset && Offset.tell() < Str.size()) {
+    const uint64_t ReadSize = 4;
+    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 bother 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));
+}
+
+static void
+callBufferedPrintfArgPush(IRBuilder<> &Builder, ArrayRef<Value *> Args,
+                          Value *PtrToStore, SparseBitVector<8> &SpecIsCString,
+                          SmallVectorImpl<StringData> &StringContents,
+                          bool IsConstFmtStr) {
+  Module *M = Builder.GetInsertBlock()->getModule();
+  auto StrIt = StringContents.begin();
+  size_t i = IsConstFmtStr ? 1 : 0;
+  for (; i < Args.size(); i++) {
+    SmallVector<Value *, 32> WhatToStore;
+    if ((i == 0) || SpecIsCString.test(i)) {
+      if (StrIt->isConst) {
+        processConstantStringArg(StrIt, Builder, WhatToStore);
+        StrIt++;
+      } else {
+        // This copies the contents of the string, however the next offset
+        // is at aligned length, the extra space that might be created due
+        // to alignment padding is not populated with any specific value
+        // here. This would be safe as long as runtime is sync with
+        // the offsets.
+        uint64_t DstAlign = (i == 0) ? 4 : 8;
+        Builder.CreateMemCpy(PtrToStore, /*DstAlign*/ Align(DstAlign), Args[i],
+                             /*SrcAlign*/ Align(1), StrIt->RealSize);
+
+        PtrToStore =
+            Builder.CreateInBoundsGEP(Builder.getInt8Ty(), PtrToStore,
+                                      {StrIt->AlignedSize}, "PrintBuffNextPtr");
+        LLVM_DEBUG(dbgs() << "inserting gep to the printf buffer:"
+                          << *PtrToStore << '\n');
+
+        // done with current argument, move to next
+        StrIt++;
+        continue;
+      }
+    } else {
+      auto IntTy = dyn_cast<IntegerType>(Args[i]->getType());
+      if (IntTy && IntTy->getBitWidth() == 32)
+        WhatToStore.push_back(
+            Builder.CreateZExt(Args[i], Builder.getInt64Ty()));
+      else
+        WhatToStore.push_back(Args[i]);
+    }
+
+    for (unsigned I = 0, E = WhatToStore.size(); I != E; ++I) {
+      Value *toStore = WhatToStore[I];
+
+      StoreInst *StBuff = Builder.CreateStore(toStore, PtrToStore);
+      LLVM_DEBUG(dbgs() << "inserting store to printf buffer:" << *StBuff
+                        << '\n');
+      PtrToStore = Builder.CreateConstInBoundsGEP1_32(
+          Builder.getInt8Ty(), PtrToStore,
+          M->getDataLayout().getTypeAllocSize(toStore->getType()),
+          "PrintBuffNextPtr");
+      LLVM_DEBUG(dbgs() << "inserting gep to the printf buffer:" << *PtrToStore
+                        << '\n');
+    }
+  }
+}
+
+Value *llvm::emitAMDGPUPrintfCall(IRBuilder<> &Builder, ArrayRef<Value *> Args,
+                                  bool IsBuffered) {
   auto NumOps = Args.size();
   assert(NumOps >= 1);
 
   auto Fmt = Args[0];
   SparseBitVector<8> SpecIsCString;
-  locateCStrings(SpecIsCString, Fmt);
+  StringRef FmtStr;
+
+  if (getConstantStringInfo(Fmt, FmtStr))
+    locateCStrings(SpecIsCString, FmtStr);
+
+  if (IsBuffered) {
+    SmallVector<StringData, 8> StringContents;
+    Module *M = Builder.GetInsertBlock()->getModule();
+    LLVMContext &Ctx = Builder.getContext();
+    auto Int1Ty = Builder.getInt1Ty();
+    auto Int8Ty = Builder.getInt8Ty();
+    auto Int32Ty = Builder.getInt32Ty();
+    bool IsConstFmtStr = !FmtStr.empty();
+
+    Value *ArgSize = nullptr;
+    Value *Ptr =
+        callBufferedPrintfStart(Builder, Args, Fmt, IsConstFmtStr,
+                                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("__printf_control_dword"), Builder.getInt32Ty(),
+        Builder.getInt32Ty(), Int1Ty, Int1Ty);
+    auto valueToStore = Builder.CreateCall(
+        CreateControlDWord,
+        {ArgSize, ConstantInt::get(Int1Ty, IsConstFmtStr ? 1 : 0, false),
+         Builder.getFalse()});
+    Builder.CreateStore(valueToStore, Ptr);
+
+    Ptr = Builder.CreateConstInBoundsGEP1_32(Int8Ty, Ptr, 4);
+
+    // Create MD5 hash for costant format string, push low 64 bits of the
+    // same onto buffer and metadata.
+    NamedMDNode *metaD = M->getOrInsertNamedMetadata("llvm.printf.fmts");
+    if (IsConstFmtStr) {
+      MD5 Hasher;
+      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();
+      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.CreateConstInBoundsGEP1_32(Int8Ty, Ptr, 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
+      if (metaD->getNumOperands() == 0) {
+        MDString *fmtStrArray =
+            MDString::get(Ctx, "0:0:ffffffff,\"Non const format string\"");
+        MDNode *myMD = MDNode::get(Ctx, fmtStrArray);
+        metaD->addOperand(myMD);
+      }
+    }
+
+    // Push The printf arguments onto buffer
+    callBufferedPrintfArgPush(Builder, Args, Ptr, SpecIsCString, StringContents,
+                              IsConstFmtStr);
+
+    // End block, returns -1 on failure
+    BranchInst::Create(End, ArgPush);
+    Builder.SetInsertPoint(End);
+    return Builder.CreateSExt(Builder.CreateNot(Cmp), Int32Ty, "printf_result");
+  }
 
   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());
 }