https://github.com/steffenlarsen created
https://github.com/llvm/llvm-project/pull/231967
AMDGPU device printf is expanded while the calling function is emitted, and a
format or %s argument is only treated as a constant string if the expansion can
read the initializer of the global it points to. A constant global whose
definition is emitted on first use, such as
```c++
const char fmt[] = {'%', 's', '\n', 0};
__device__ void f() { printf(fmt, "x"); }
```
is still a declaration at that point. The format was therefore treated as
non-constant and its %s arguments were sent as pointer values, making the
hostcall runtime print garbage and buffered printf crash.
Before emitting the arguments, emit the definitions of the deferred constant
globals they refer to. Only definitions that would otherwise go through the
deferred queue are pulled forward, and EmitDeferred skips them later because
they are already defined. A definition is not pulled forward if it would
replace an existing declaration whose type differs from the initializer's,
which can happen for an array with a long zero tail, since values referring to
the old global may still be live.
Assisted-by: Claude Opus 5.5
>From fd5bed62e93abe36109e24bc38a18753639a8d35 Mon Sep 17 00:00:00 2001
From: Steffen Holst Larsen <[email protected]>
Date: Wed, 14 Oct 2026 09:49:01 -0500
Subject: [PATCH] [clang][AMDGPU] Emit deferred string globals before expanding
printf
AMDGPU device printf is expanded while the calling function is emitted,
and a format or %s argument is only treated as a constant string if the
expansion can read the initializer of the global it points to. A
constant global whose definition is emitted on first use, such as
```c++
const char fmt[] = {'%', 's', '\n', 0};
__device__ void f() { printf(fmt, "x"); }
```
is still a declaration at that point. The format was therefore treated
as non-constant and its %s arguments were sent as pointer values,
making the hostcall runtime print garbage and buffered printf crash.
Before emitting the arguments, emit the definitions of the deferred
constant globals they refer to. Only definitions that would otherwise go
through the deferred queue are pulled forward, and EmitDeferred skips
them later because they are already defined. A definition is not pulled
forward if it would replace an existing declaration whose type differs
from the initializer's, which can happen for an array with a long zero
tail, since values referring to the old global may still be live.
Assisted-by: Claude Opus 5.5
Signed-off-by: Steffen Holst Larsen <[email protected]>
---
clang/lib/CodeGen/CGGPUBuiltin.cpp | 32 ++
clang/lib/CodeGen/CodeGenModule.cpp | 23 ++
clang/lib/CodeGen/CodeGenModule.h | 5 +
clang/test/CodeGenHIP/printf.cpp | 391 ++++++++++++++++++-
clang/test/CodeGenHIP/printf_nonhostcall.cpp | 326 ++++++++++++++++
5 files changed, 775 insertions(+), 2 deletions(-)
diff --git a/clang/lib/CodeGen/CGGPUBuiltin.cpp
b/clang/lib/CodeGen/CGGPUBuiltin.cpp
index 327a317a1e78ed..ab5e58908b5caf 100644
--- a/clang/lib/CodeGen/CGGPUBuiltin.cpp
+++ b/clang/lib/CodeGen/CGGPUBuiltin.cpp
@@ -155,6 +155,29 @@ RValue
CodeGenFunction::EmitNVPTXDevicePrintfCallExpr(const CallExpr *E) {
E, this, GetVprintfDeclaration(CGM.getModule()), false);
}
+// Collects the constant global variables that \p S refers to.
+static void
+collectReferencedGlobalVars(const ASTContext &Ctx, const Stmt *S,
+ SmallVectorImpl<const VarDecl *> &Vars) {
+ if (!S)
+ return;
+ const ValueDecl *D = nullptr;
+ if (const auto *DRE = dyn_cast<DeclRefExpr>(S)) {
+ if (DRE->isNonOdrUse() == NOUR_None)
+ D = DRE->getDecl();
+ } else if (const auto *ME = dyn_cast<MemberExpr>(S)) {
+ if (ME->isNonOdrUse() == NOUR_None)
+ D = ME->getMemberDecl();
+ }
+ if (const auto *VD = dyn_cast_if_present<VarDecl>(D))
+ if (VD->hasGlobalStorage() && !VD->isStaticLocal() &&
+ VD->getType().isConstantStorage(Ctx, /*ExcludeCtor=*/true,
+ /*ExcludeDtor=*/true))
+ Vars.push_back(VD);
+ for (const Stmt *Child : S->children())
+ collectReferencedGlobalVars(Ctx, Child, Vars);
+}
+
RValue CodeGenFunction::EmitAMDGPUDevicePrintfCallExpr(const CallExpr *E) {
assert(getTarget().getTriple().isAMDGCN() ||
(getTarget().getTriple().isSPIRV() &&
@@ -163,6 +186,15 @@ RValue
CodeGenFunction::EmitAMDGPUDevicePrintfCallExpr(const CallExpr *E) {
E->getBuiltinCallee() == Builtin::BI__builtin_printf);
assert(E->getNumArgs() >= 1); // printf always has at least one arg.
+ // The expansion finds constant format and %s strings in the initializers of
+ // the globals they point to, so emit any that are still deferred before the
+ // arguments refer to them.
+ SmallVector<const VarDecl *, 4> Globals;
+ for (const Expr *Arg : E->arguments())
+ collectReferencedGlobalVars(getContext(), Arg, Globals);
+ for (const VarDecl *VD : Globals)
+ CGM.emitDeferredGlobalVarDefinition(VD);
+
CallArgList CallArgs;
EmitCallArgs(CallArgs,
E->getDirectCallee()->getType()->getAs<FunctionProtoType>(),
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp
b/clang/lib/CodeGen/CodeGenModule.cpp
index bf03e10f9fb422..04e7240e2254bb 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -8579,6 +8579,29 @@ bool CodeGenModule::lookupRepresentativeDecl(StringRef
MangledName,
return true;
}
+void CodeGenModule::emitDeferredGlobalVarDefinition(const VarDecl *VD) {
+ const VarDecl *Def = VD->getDefinition();
+ // Definitions that must be emitted can be emitted directly by EmitGlobal, so
+ // emitting them here could define them twice. Deferred ones are only emitted
+ // by EmitDeferred, which skips globals that are already defined.
+ if (!Def || MustBeEmitted(Def) || !Def->hasConstantInitialization())
+ return;
+ // Defining a global that is already declared with a different type or
+ // address space replaces it, which would leave the caller's values referring
+ // to the erased global.
+ if (llvm::GlobalValue *GV = GetGlobalValue(getMangledName(Def))) {
+ if (!GV->isDeclaration())
+ return;
+ llvm::Constant *Init =
+ ConstantEmitter(*this).tryEmitAbstractForInitializer(*Def);
+ if (!Init || GV->getValueType() != Init->getType() ||
+ GV->getAddressSpace() !=
+ getContext().getTargetAddressSpace(GetGlobalVarAddressSpace(Def)))
+ return;
+ }
+ EmitGlobalVarDefinition(Def);
+}
+
/// Emits metadata nodes associating all the global values in the
/// current module with the Decls they came from. This is useful for
/// projects using IR gen as a subroutine.
diff --git a/clang/lib/CodeGen/CodeGenModule.h
b/clang/lib/CodeGen/CodeGenModule.h
index 77ed7abc48a9ed..97a44457aea787 100644
--- a/clang/lib/CodeGen/CodeGenModule.h
+++ b/clang/lib/CodeGen/CodeGenModule.h
@@ -866,6 +866,11 @@ class CodeGenModule : public CodeGenTypeCache {
bool lookupRepresentativeDecl(StringRef MangledName,
GlobalDecl &Result) const;
+ /// If the constant-initialized definition of \p VD is deferred, emit it now
+ /// so that its initializer can be inspected before the end of the
+ /// translation unit. Does nothing if that would replace an existing global.
+ void emitDeferredGlobalVarDefinition(const VarDecl *VD);
+
llvm::Constant *getAtomicSetterHelperFnMap(QualType Ty) {
return AtomicSetterHelperFnMap[Ty];
}
diff --git a/clang/test/CodeGenHIP/printf.cpp b/clang/test/CodeGenHIP/printf.cpp
index 5cddf6c8c59c5f..e0f19c56db00d6 100644
--- a/clang/test/CodeGenHIP/printf.cpp
+++ b/clang/test/CodeGenHIP/printf.cpp
@@ -119,7 +119,7 @@ __device__ int foo1() {
__device__ char *dstr;
// AMDGCN-LABEL: define dso_local noundef i32 @_Z4foo2v(
-// AMDGCN-SAME: ) #[[ATTR0:[0-9]+]] {
+// AMDGCN-SAME: ) #[[ATTR0]] {
// AMDGCN-NEXT: [[ENTRY:.*]]:
// AMDGCN-NEXT: [[TMP0:%.*]] = load ptr, ptr addrspacecast (ptr
addrspace(1) @dstr to ptr), align 8
// AMDGCN-NEXT: [[TMP1:%.*]] = load ptr, ptr addrspacecast (ptr
addrspace(1) @dstr to ptr), align 8
@@ -163,7 +163,7 @@ __device__ char *dstr;
// AMDGCN-NEXT: ret i32 [[TMP26]]
//
// AMDGCNSPIRV-LABEL: define spir_func noundef i32 @_Z4foo2v(
-// AMDGCNSPIRV-SAME: ) addrspace(4) #[[ATTR0:[0-9]+]] {
+// AMDGCNSPIRV-SAME: ) addrspace(4) #[[ATTR0]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*]]:
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4)
addrspacecast (ptr addrspace(1) @dstr to ptr addrspace(4)), align 8
// AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = load ptr addrspace(4), ptr addrspace(4)
addrspacecast (ptr addrspace(1) @dstr to ptr addrspace(4)), align 8
@@ -209,3 +209,390 @@ __device__ char *dstr;
__device__ int foo2() {
return printf("%s %p\n", dstr, dstr);
}
+
+// Constant formats in globals whose definitions are deferred.
+const char fmt_arr[] = {'%', 's', '\n', 0};
+struct S { static constexpr char fmt[] = "%s!\n"; };
+
+// AMDGCN-LABEL: define dso_local noundef i32 @_Z9fmt_arrayv(
+// AMDGCN-SAME: ) #[[ATTR0]] {
+// AMDGCN-NEXT: [[ENTRY:.*]]:
+// AMDGCN-NEXT: [[TMP0:%.*]] = call i64 @__ockl_printf_begin(i64 0)
+// AMDGCN-NEXT: [[TMP1:%.*]] = icmp eq ptr addrspacecast (ptr addrspace(4)
@_ZL7fmt_arr to ptr), null
+// AMDGCN-NEXT: br i1 [[TMP1]], label %[[STRLEN_JOIN:.*]], label
%[[STRLEN_WHILE:.*]]
+// AMDGCN: [[STRLEN_WHILE]]:
+// AMDGCN-NEXT: [[TMP2:%.*]] = phi ptr [ addrspacecast (ptr addrspace(4)
@_ZL7fmt_arr to ptr), %[[ENTRY]] ], [ [[TMP3:%.*]], %[[STRLEN_WHILE]] ]
+// AMDGCN-NEXT: [[TMP3]] = getelementptr i8, ptr [[TMP2]], i64 1
+// AMDGCN-NEXT: [[TMP4:%.*]] = load i8, ptr [[TMP2]], align 1
+// AMDGCN-NEXT: [[TMP5:%.*]] = icmp eq i8 [[TMP4]], 0
+// AMDGCN-NEXT: br i1 [[TMP5]], label %[[STRLEN_WHILE_DONE:.*]], label
%[[STRLEN_WHILE]]
+// AMDGCN: [[STRLEN_WHILE_DONE]]:
+// AMDGCN-NEXT: [[TMP6:%.*]] = ptrtoaddr ptr [[TMP2]] to i64
+// AMDGCN-NEXT: [[TMP7:%.*]] = sub i64 [[TMP6]], ptrtoaddr (ptr
addrspacecast (ptr addrspace(4) @_ZL7fmt_arr to ptr) to i64)
+// AMDGCN-NEXT: [[TMP8:%.*]] = add i64 [[TMP7]], 1
+// AMDGCN-NEXT: br label %[[STRLEN_JOIN]]
+// AMDGCN: [[STRLEN_JOIN]]:
+// AMDGCN-NEXT: [[TMP9:%.*]] = phi i64 [ [[TMP8]], %[[STRLEN_WHILE_DONE]]
], [ 0, %[[ENTRY]] ]
+// AMDGCN-NEXT: [[TMP10:%.*]] = call i64 @__ockl_printf_append_string_n(i64
[[TMP0]], ptr addrspacecast (ptr addrspace(4) @_ZL7fmt_arr to ptr), i64
[[TMP9]], i32 0)
+// AMDGCN-NEXT: [[TMP11:%.*]] = icmp eq ptr addrspacecast (ptr addrspace(4)
@.str.3 to ptr), null
+// AMDGCN-NEXT: br i1 [[TMP11]], label %[[STRLEN_JOIN1:.*]], label
%[[STRLEN_WHILE2:.*]]
+// AMDGCN: [[STRLEN_WHILE2]]:
+// AMDGCN-NEXT: [[TMP12:%.*]] = phi ptr [ addrspacecast (ptr addrspace(4)
@.str.3 to ptr), %[[STRLEN_JOIN]] ], [ [[TMP13:%.*]], %[[STRLEN_WHILE2]] ]
+// AMDGCN-NEXT: [[TMP13]] = getelementptr i8, ptr [[TMP12]], i64 1
+// AMDGCN-NEXT: [[TMP14:%.*]] = load i8, ptr [[TMP12]], align 1
+// AMDGCN-NEXT: [[TMP15:%.*]] = icmp eq i8 [[TMP14]], 0
+// AMDGCN-NEXT: br i1 [[TMP15]], label %[[STRLEN_WHILE_DONE3:.*]], label
%[[STRLEN_WHILE2]]
+// AMDGCN: [[STRLEN_WHILE_DONE3]]:
+// AMDGCN-NEXT: [[TMP16:%.*]] = ptrtoaddr ptr [[TMP12]] to i64
+// AMDGCN-NEXT: [[TMP17:%.*]] = sub i64 [[TMP16]], ptrtoaddr (ptr
addrspacecast (ptr addrspace(4) @.str.3 to ptr) to i64)
+// AMDGCN-NEXT: [[TMP18:%.*]] = add i64 [[TMP17]], 1
+// AMDGCN-NEXT: br label %[[STRLEN_JOIN1]]
+// AMDGCN: [[STRLEN_JOIN1]]:
+// AMDGCN-NEXT: [[TMP19:%.*]] = phi i64 [ [[TMP18]],
%[[STRLEN_WHILE_DONE3]] ], [ 0, %[[STRLEN_JOIN]] ]
+// AMDGCN-NEXT: [[TMP20:%.*]] = call i64 @__ockl_printf_append_string_n(i64
[[TMP10]], ptr addrspacecast (ptr addrspace(4) @.str.3 to ptr), i64 [[TMP19]],
i32 1)
+// AMDGCN-NEXT: [[TMP21:%.*]] = trunc i64 [[TMP20]] to i32
+// AMDGCN-NEXT: ret i32 [[TMP21]]
+//
+// AMDGCNSPIRV-LABEL: define spir_func noundef i32 @_Z9fmt_arrayv(
+// AMDGCNSPIRV-SAME: ) addrspace(4) #[[ATTR0]] {
+// AMDGCNSPIRV-NEXT: [[ENTRY:.*]]:
+// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = call addrspace(4) i64
@__ockl_printf_begin(i64 0)
+// AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = icmp eq ptr addrspace(4) addrspacecast
(ptr addrspace(1) @_ZL7fmt_arr to ptr addrspace(4)), null
+// AMDGCNSPIRV-NEXT: br i1 [[TMP1]], label %[[STRLEN_JOIN:.*]], label
%[[STRLEN_WHILE:.*]]
+// AMDGCNSPIRV: [[STRLEN_WHILE]]:
+// AMDGCNSPIRV-NEXT: [[TMP2:%.*]] = phi ptr addrspace(4) [ addrspacecast
(ptr addrspace(1) @_ZL7fmt_arr to ptr addrspace(4)), %[[ENTRY]] ], [
[[TMP3:%.*]], %[[STRLEN_WHILE]] ]
+// AMDGCNSPIRV-NEXT: [[TMP3]] = getelementptr i8, ptr addrspace(4)
[[TMP2]], i64 1
+// AMDGCNSPIRV-NEXT: [[TMP4:%.*]] = load i8, ptr addrspace(4) [[TMP2]],
align 1
+// AMDGCNSPIRV-NEXT: [[TMP5:%.*]] = icmp eq i8 [[TMP4]], 0
+// AMDGCNSPIRV-NEXT: br i1 [[TMP5]], label %[[STRLEN_WHILE_DONE:.*]], label
%[[STRLEN_WHILE]]
+// AMDGCNSPIRV: [[STRLEN_WHILE_DONE]]:
+// AMDGCNSPIRV-NEXT: [[TMP6:%.*]] = ptrtoaddr ptr addrspace(4) [[TMP2]] to
i64
+// AMDGCNSPIRV-NEXT: [[TMP7:%.*]] = sub i64 [[TMP6]], ptrtoaddr (ptr
addrspace(4) addrspacecast (ptr addrspace(1) @_ZL7fmt_arr to ptr addrspace(4))
to i64)
+// AMDGCNSPIRV-NEXT: [[TMP8:%.*]] = add i64 [[TMP7]], 1
+// AMDGCNSPIRV-NEXT: br label %[[STRLEN_JOIN]]
+// AMDGCNSPIRV: [[STRLEN_JOIN]]:
+// AMDGCNSPIRV-NEXT: [[TMP9:%.*]] = phi i64 [ [[TMP8]],
%[[STRLEN_WHILE_DONE]] ], [ 0, %[[ENTRY]] ]
+// AMDGCNSPIRV-NEXT: [[TMP10:%.*]] = call addrspace(4) i64
@__ockl_printf_append_string_n(i64 [[TMP0]], ptr addrspace(4) addrspacecast
(ptr addrspace(1) @_ZL7fmt_arr to ptr addrspace(4)), i64 [[TMP9]], i32 0)
+// AMDGCNSPIRV-NEXT: [[TMP11:%.*]] = icmp eq ptr addrspace(4) addrspacecast
(ptr addrspace(1) @.str.3 to ptr addrspace(4)), null
+// AMDGCNSPIRV-NEXT: br i1 [[TMP11]], label %[[STRLEN_JOIN1:.*]], label
%[[STRLEN_WHILE2:.*]]
+// AMDGCNSPIRV: [[STRLEN_WHILE2]]:
+// AMDGCNSPIRV-NEXT: [[TMP12:%.*]] = phi ptr addrspace(4) [ addrspacecast
(ptr addrspace(1) @.str.3 to ptr addrspace(4)), %[[STRLEN_JOIN]] ], [
[[TMP13:%.*]], %[[STRLEN_WHILE2]] ]
+// AMDGCNSPIRV-NEXT: [[TMP13]] = getelementptr i8, ptr addrspace(4)
[[TMP12]], i64 1
+// AMDGCNSPIRV-NEXT: [[TMP14:%.*]] = load i8, ptr addrspace(4) [[TMP12]],
align 1
+// AMDGCNSPIRV-NEXT: [[TMP15:%.*]] = icmp eq i8 [[TMP14]], 0
+// AMDGCNSPIRV-NEXT: br i1 [[TMP15]], label %[[STRLEN_WHILE_DONE3:.*]],
label %[[STRLEN_WHILE2]]
+// AMDGCNSPIRV: [[STRLEN_WHILE_DONE3]]:
+// AMDGCNSPIRV-NEXT: [[TMP16:%.*]] = ptrtoaddr ptr addrspace(4) [[TMP12]]
to i64
+// AMDGCNSPIRV-NEXT: [[TMP17:%.*]] = sub i64 [[TMP16]], ptrtoaddr (ptr
addrspace(4) addrspacecast (ptr addrspace(1) @.str.3 to ptr addrspace(4)) to
i64)
+// AMDGCNSPIRV-NEXT: [[TMP18:%.*]] = add i64 [[TMP17]], 1
+// AMDGCNSPIRV-NEXT: br label %[[STRLEN_JOIN1]]
+// AMDGCNSPIRV: [[STRLEN_JOIN1]]:
+// AMDGCNSPIRV-NEXT: [[TMP19:%.*]] = phi i64 [ [[TMP18]],
%[[STRLEN_WHILE_DONE3]] ], [ 0, %[[STRLEN_JOIN]] ]
+// AMDGCNSPIRV-NEXT: [[TMP20:%.*]] = call addrspace(4) i64
@__ockl_printf_append_string_n(i64 [[TMP10]], ptr addrspace(4) addrspacecast
(ptr addrspace(1) @.str.3 to ptr addrspace(4)), i64 [[TMP19]], i32 1)
+// AMDGCNSPIRV-NEXT: [[TMP21:%.*]] = trunc i64 [[TMP20]] to i32
+// AMDGCNSPIRV-NEXT: ret i32 [[TMP21]]
+//
+__device__ int fmt_array() { return printf(fmt_arr, "x"); }
+
+// AMDGCN-LABEL: define dso_local noundef i32 @_Z10member_fmtv(
+// AMDGCN-SAME: ) #[[ATTR0]] {
+// AMDGCN-NEXT: [[ENTRY:.*]]:
+// AMDGCN-NEXT: [[TMP0:%.*]] = call i64 @__ockl_printf_begin(i64 0)
+// AMDGCN-NEXT: [[TMP1:%.*]] = icmp eq ptr addrspacecast (ptr addrspace(4)
@_ZN1S3fmtE to ptr), null
+// AMDGCN-NEXT: br i1 [[TMP1]], label %[[STRLEN_JOIN:.*]], label
%[[STRLEN_WHILE:.*]]
+// AMDGCN: [[STRLEN_WHILE]]:
+// AMDGCN-NEXT: [[TMP2:%.*]] = phi ptr [ addrspacecast (ptr addrspace(4)
@_ZN1S3fmtE to ptr), %[[ENTRY]] ], [ [[TMP3:%.*]], %[[STRLEN_WHILE]] ]
+// AMDGCN-NEXT: [[TMP3]] = getelementptr i8, ptr [[TMP2]], i64 1
+// AMDGCN-NEXT: [[TMP4:%.*]] = load i8, ptr [[TMP2]], align 1
+// AMDGCN-NEXT: [[TMP5:%.*]] = icmp eq i8 [[TMP4]], 0
+// AMDGCN-NEXT: br i1 [[TMP5]], label %[[STRLEN_WHILE_DONE:.*]], label
%[[STRLEN_WHILE]]
+// AMDGCN: [[STRLEN_WHILE_DONE]]:
+// AMDGCN-NEXT: [[TMP6:%.*]] = ptrtoaddr ptr [[TMP2]] to i64
+// AMDGCN-NEXT: [[TMP7:%.*]] = sub i64 [[TMP6]], ptrtoaddr (ptr
addrspacecast (ptr addrspace(4) @_ZN1S3fmtE to ptr) to i64)
+// AMDGCN-NEXT: [[TMP8:%.*]] = add i64 [[TMP7]], 1
+// AMDGCN-NEXT: br label %[[STRLEN_JOIN]]
+// AMDGCN: [[STRLEN_JOIN]]:
+// AMDGCN-NEXT: [[TMP9:%.*]] = phi i64 [ [[TMP8]], %[[STRLEN_WHILE_DONE]]
], [ 0, %[[ENTRY]] ]
+// AMDGCN-NEXT: [[TMP10:%.*]] = call i64 @__ockl_printf_append_string_n(i64
[[TMP0]], ptr addrspacecast (ptr addrspace(4) @_ZN1S3fmtE to ptr), i64
[[TMP9]], i32 0)
+// AMDGCN-NEXT: [[TMP11:%.*]] = icmp eq ptr addrspacecast (ptr addrspace(4)
@.str.4 to ptr), null
+// AMDGCN-NEXT: br i1 [[TMP11]], label %[[STRLEN_JOIN1:.*]], label
%[[STRLEN_WHILE2:.*]]
+// AMDGCN: [[STRLEN_WHILE2]]:
+// AMDGCN-NEXT: [[TMP12:%.*]] = phi ptr [ addrspacecast (ptr addrspace(4)
@.str.4 to ptr), %[[STRLEN_JOIN]] ], [ [[TMP13:%.*]], %[[STRLEN_WHILE2]] ]
+// AMDGCN-NEXT: [[TMP13]] = getelementptr i8, ptr [[TMP12]], i64 1
+// AMDGCN-NEXT: [[TMP14:%.*]] = load i8, ptr [[TMP12]], align 1
+// AMDGCN-NEXT: [[TMP15:%.*]] = icmp eq i8 [[TMP14]], 0
+// AMDGCN-NEXT: br i1 [[TMP15]], label %[[STRLEN_WHILE_DONE3:.*]], label
%[[STRLEN_WHILE2]]
+// AMDGCN: [[STRLEN_WHILE_DONE3]]:
+// AMDGCN-NEXT: [[TMP16:%.*]] = ptrtoaddr ptr [[TMP12]] to i64
+// AMDGCN-NEXT: [[TMP17:%.*]] = sub i64 [[TMP16]], ptrtoaddr (ptr
addrspacecast (ptr addrspace(4) @.str.4 to ptr) to i64)
+// AMDGCN-NEXT: [[TMP18:%.*]] = add i64 [[TMP17]], 1
+// AMDGCN-NEXT: br label %[[STRLEN_JOIN1]]
+// AMDGCN: [[STRLEN_JOIN1]]:
+// AMDGCN-NEXT: [[TMP19:%.*]] = phi i64 [ [[TMP18]],
%[[STRLEN_WHILE_DONE3]] ], [ 0, %[[STRLEN_JOIN]] ]
+// AMDGCN-NEXT: [[TMP20:%.*]] = call i64 @__ockl_printf_append_string_n(i64
[[TMP10]], ptr addrspacecast (ptr addrspace(4) @.str.4 to ptr), i64 [[TMP19]],
i32 1)
+// AMDGCN-NEXT: [[TMP21:%.*]] = trunc i64 [[TMP20]] to i32
+// AMDGCN-NEXT: ret i32 [[TMP21]]
+//
+// AMDGCNSPIRV-LABEL: define spir_func noundef i32 @_Z10member_fmtv(
+// AMDGCNSPIRV-SAME: ) addrspace(4) #[[ATTR0]] {
+// AMDGCNSPIRV-NEXT: [[ENTRY:.*]]:
+// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = call addrspace(4) i64
@__ockl_printf_begin(i64 0)
+// AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = icmp eq ptr addrspace(4) addrspacecast
(ptr addrspace(1) @_ZN1S3fmtE to ptr addrspace(4)), null
+// AMDGCNSPIRV-NEXT: br i1 [[TMP1]], label %[[STRLEN_JOIN:.*]], label
%[[STRLEN_WHILE:.*]]
+// AMDGCNSPIRV: [[STRLEN_WHILE]]:
+// AMDGCNSPIRV-NEXT: [[TMP2:%.*]] = phi ptr addrspace(4) [ addrspacecast
(ptr addrspace(1) @_ZN1S3fmtE to ptr addrspace(4)), %[[ENTRY]] ], [
[[TMP3:%.*]], %[[STRLEN_WHILE]] ]
+// AMDGCNSPIRV-NEXT: [[TMP3]] = getelementptr i8, ptr addrspace(4)
[[TMP2]], i64 1
+// AMDGCNSPIRV-NEXT: [[TMP4:%.*]] = load i8, ptr addrspace(4) [[TMP2]],
align 1
+// AMDGCNSPIRV-NEXT: [[TMP5:%.*]] = icmp eq i8 [[TMP4]], 0
+// AMDGCNSPIRV-NEXT: br i1 [[TMP5]], label %[[STRLEN_WHILE_DONE:.*]], label
%[[STRLEN_WHILE]]
+// AMDGCNSPIRV: [[STRLEN_WHILE_DONE]]:
+// AMDGCNSPIRV-NEXT: [[TMP6:%.*]] = ptrtoaddr ptr addrspace(4) [[TMP2]] to
i64
+// AMDGCNSPIRV-NEXT: [[TMP7:%.*]] = sub i64 [[TMP6]], ptrtoaddr (ptr
addrspace(4) addrspacecast (ptr addrspace(1) @_ZN1S3fmtE to ptr addrspace(4))
to i64)
+// AMDGCNSPIRV-NEXT: [[TMP8:%.*]] = add i64 [[TMP7]], 1
+// AMDGCNSPIRV-NEXT: br label %[[STRLEN_JOIN]]
+// AMDGCNSPIRV: [[STRLEN_JOIN]]:
+// AMDGCNSPIRV-NEXT: [[TMP9:%.*]] = phi i64 [ [[TMP8]],
%[[STRLEN_WHILE_DONE]] ], [ 0, %[[ENTRY]] ]
+// AMDGCNSPIRV-NEXT: [[TMP10:%.*]] = call addrspace(4) i64
@__ockl_printf_append_string_n(i64 [[TMP0]], ptr addrspace(4) addrspacecast
(ptr addrspace(1) @_ZN1S3fmtE to ptr addrspace(4)), i64 [[TMP9]], i32 0)
+// AMDGCNSPIRV-NEXT: [[TMP11:%.*]] = icmp eq ptr addrspace(4) addrspacecast
(ptr addrspace(1) @.str.4 to ptr addrspace(4)), null
+// AMDGCNSPIRV-NEXT: br i1 [[TMP11]], label %[[STRLEN_JOIN1:.*]], label
%[[STRLEN_WHILE2:.*]]
+// AMDGCNSPIRV: [[STRLEN_WHILE2]]:
+// AMDGCNSPIRV-NEXT: [[TMP12:%.*]] = phi ptr addrspace(4) [ addrspacecast
(ptr addrspace(1) @.str.4 to ptr addrspace(4)), %[[STRLEN_JOIN]] ], [
[[TMP13:%.*]], %[[STRLEN_WHILE2]] ]
+// AMDGCNSPIRV-NEXT: [[TMP13]] = getelementptr i8, ptr addrspace(4)
[[TMP12]], i64 1
+// AMDGCNSPIRV-NEXT: [[TMP14:%.*]] = load i8, ptr addrspace(4) [[TMP12]],
align 1
+// AMDGCNSPIRV-NEXT: [[TMP15:%.*]] = icmp eq i8 [[TMP14]], 0
+// AMDGCNSPIRV-NEXT: br i1 [[TMP15]], label %[[STRLEN_WHILE_DONE3:.*]],
label %[[STRLEN_WHILE2]]
+// AMDGCNSPIRV: [[STRLEN_WHILE_DONE3]]:
+// AMDGCNSPIRV-NEXT: [[TMP16:%.*]] = ptrtoaddr ptr addrspace(4) [[TMP12]]
to i64
+// AMDGCNSPIRV-NEXT: [[TMP17:%.*]] = sub i64 [[TMP16]], ptrtoaddr (ptr
addrspace(4) addrspacecast (ptr addrspace(1) @.str.4 to ptr addrspace(4)) to
i64)
+// AMDGCNSPIRV-NEXT: [[TMP18:%.*]] = add i64 [[TMP17]], 1
+// AMDGCNSPIRV-NEXT: br label %[[STRLEN_JOIN1]]
+// AMDGCNSPIRV: [[STRLEN_JOIN1]]:
+// AMDGCNSPIRV-NEXT: [[TMP19:%.*]] = phi i64 [ [[TMP18]],
%[[STRLEN_WHILE_DONE3]] ], [ 0, %[[STRLEN_JOIN]] ]
+// AMDGCNSPIRV-NEXT: [[TMP20:%.*]] = call addrspace(4) i64
@__ockl_printf_append_string_n(i64 [[TMP10]], ptr addrspace(4) addrspacecast
(ptr addrspace(1) @.str.4 to ptr addrspace(4)), i64 [[TMP19]], i32 1)
+// AMDGCNSPIRV-NEXT: [[TMP21:%.*]] = trunc i64 [[TMP20]] to i32
+// AMDGCNSPIRV-NEXT: ret i32 [[TMP21]]
+//
+__device__ int member_fmt() { return printf(S::fmt, "y"); }
+
+// The initializer's type differs from the array type because of the zero tail.
+const char long_fmt[64] = {'%', 's', '\n'};
+
+// AMDGCN-LABEL: define dso_local noundef i32 @_Z14long_fmt_arrayv(
+// AMDGCN-SAME: ) #[[ATTR0]] {
+// AMDGCN-NEXT: [[ENTRY:.*]]:
+// AMDGCN-NEXT: [[TMP0:%.*]] = call i64 @__ockl_printf_begin(i64 0)
+// AMDGCN-NEXT: [[TMP1:%.*]] = icmp eq ptr addrspacecast (ptr addrspace(4)
@_ZL8long_fmt to ptr), null
+// AMDGCN-NEXT: br i1 [[TMP1]], label %[[STRLEN_JOIN:.*]], label
%[[STRLEN_WHILE:.*]]
+// AMDGCN: [[STRLEN_WHILE]]:
+// AMDGCN-NEXT: [[TMP2:%.*]] = phi ptr [ addrspacecast (ptr addrspace(4)
@_ZL8long_fmt to ptr), %[[ENTRY]] ], [ [[TMP3:%.*]], %[[STRLEN_WHILE]] ]
+// AMDGCN-NEXT: [[TMP3]] = getelementptr i8, ptr [[TMP2]], i64 1
+// AMDGCN-NEXT: [[TMP4:%.*]] = load i8, ptr [[TMP2]], align 1
+// AMDGCN-NEXT: [[TMP5:%.*]] = icmp eq i8 [[TMP4]], 0
+// AMDGCN-NEXT: br i1 [[TMP5]], label %[[STRLEN_WHILE_DONE:.*]], label
%[[STRLEN_WHILE]]
+// AMDGCN: [[STRLEN_WHILE_DONE]]:
+// AMDGCN-NEXT: [[TMP6:%.*]] = ptrtoaddr ptr [[TMP2]] to i64
+// AMDGCN-NEXT: [[TMP7:%.*]] = sub i64 [[TMP6]], ptrtoaddr (ptr
addrspacecast (ptr addrspace(4) @_ZL8long_fmt to ptr) to i64)
+// AMDGCN-NEXT: [[TMP8:%.*]] = add i64 [[TMP7]], 1
+// AMDGCN-NEXT: br label %[[STRLEN_JOIN]]
+// AMDGCN: [[STRLEN_JOIN]]:
+// AMDGCN-NEXT: [[TMP9:%.*]] = phi i64 [ [[TMP8]], %[[STRLEN_WHILE_DONE]]
], [ 0, %[[ENTRY]] ]
+// AMDGCN-NEXT: [[TMP10:%.*]] = call i64 @__ockl_printf_append_string_n(i64
[[TMP0]], ptr addrspacecast (ptr addrspace(4) @_ZL8long_fmt to ptr), i64
[[TMP9]], i32 0)
+// AMDGCN-NEXT: [[TMP11:%.*]] = icmp eq ptr addrspacecast (ptr addrspace(4)
@.str.5 to ptr), null
+// AMDGCN-NEXT: br i1 [[TMP11]], label %[[STRLEN_JOIN1:.*]], label
%[[STRLEN_WHILE2:.*]]
+// AMDGCN: [[STRLEN_WHILE2]]:
+// AMDGCN-NEXT: [[TMP12:%.*]] = phi ptr [ addrspacecast (ptr addrspace(4)
@.str.5 to ptr), %[[STRLEN_JOIN]] ], [ [[TMP13:%.*]], %[[STRLEN_WHILE2]] ]
+// AMDGCN-NEXT: [[TMP13]] = getelementptr i8, ptr [[TMP12]], i64 1
+// AMDGCN-NEXT: [[TMP14:%.*]] = load i8, ptr [[TMP12]], align 1
+// AMDGCN-NEXT: [[TMP15:%.*]] = icmp eq i8 [[TMP14]], 0
+// AMDGCN-NEXT: br i1 [[TMP15]], label %[[STRLEN_WHILE_DONE3:.*]], label
%[[STRLEN_WHILE2]]
+// AMDGCN: [[STRLEN_WHILE_DONE3]]:
+// AMDGCN-NEXT: [[TMP16:%.*]] = ptrtoaddr ptr [[TMP12]] to i64
+// AMDGCN-NEXT: [[TMP17:%.*]] = sub i64 [[TMP16]], ptrtoaddr (ptr
addrspacecast (ptr addrspace(4) @.str.5 to ptr) to i64)
+// AMDGCN-NEXT: [[TMP18:%.*]] = add i64 [[TMP17]], 1
+// AMDGCN-NEXT: br label %[[STRLEN_JOIN1]]
+// AMDGCN: [[STRLEN_JOIN1]]:
+// AMDGCN-NEXT: [[TMP19:%.*]] = phi i64 [ [[TMP18]],
%[[STRLEN_WHILE_DONE3]] ], [ 0, %[[STRLEN_JOIN]] ]
+// AMDGCN-NEXT: [[TMP20:%.*]] = call i64 @__ockl_printf_append_string_n(i64
[[TMP10]], ptr addrspacecast (ptr addrspace(4) @.str.5 to ptr), i64 [[TMP19]],
i32 1)
+// AMDGCN-NEXT: [[TMP21:%.*]] = trunc i64 [[TMP20]] to i32
+// AMDGCN-NEXT: ret i32 [[TMP21]]
+//
+// AMDGCNSPIRV-LABEL: define spir_func noundef i32 @_Z14long_fmt_arrayv(
+// AMDGCNSPIRV-SAME: ) addrspace(4) #[[ATTR0]] {
+// AMDGCNSPIRV-NEXT: [[ENTRY:.*]]:
+// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = call addrspace(4) i64
@__ockl_printf_begin(i64 0)
+// AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = icmp eq ptr addrspace(4) addrspacecast
(ptr addrspace(1) @_ZL8long_fmt to ptr addrspace(4)), null
+// AMDGCNSPIRV-NEXT: br i1 [[TMP1]], label %[[STRLEN_JOIN:.*]], label
%[[STRLEN_WHILE:.*]]
+// AMDGCNSPIRV: [[STRLEN_WHILE]]:
+// AMDGCNSPIRV-NEXT: [[TMP2:%.*]] = phi ptr addrspace(4) [ addrspacecast
(ptr addrspace(1) @_ZL8long_fmt to ptr addrspace(4)), %[[ENTRY]] ], [
[[TMP3:%.*]], %[[STRLEN_WHILE]] ]
+// AMDGCNSPIRV-NEXT: [[TMP3]] = getelementptr i8, ptr addrspace(4)
[[TMP2]], i64 1
+// AMDGCNSPIRV-NEXT: [[TMP4:%.*]] = load i8, ptr addrspace(4) [[TMP2]],
align 1
+// AMDGCNSPIRV-NEXT: [[TMP5:%.*]] = icmp eq i8 [[TMP4]], 0
+// AMDGCNSPIRV-NEXT: br i1 [[TMP5]], label %[[STRLEN_WHILE_DONE:.*]], label
%[[STRLEN_WHILE]]
+// AMDGCNSPIRV: [[STRLEN_WHILE_DONE]]:
+// AMDGCNSPIRV-NEXT: [[TMP6:%.*]] = ptrtoaddr ptr addrspace(4) [[TMP2]] to
i64
+// AMDGCNSPIRV-NEXT: [[TMP7:%.*]] = sub i64 [[TMP6]], ptrtoaddr (ptr
addrspace(4) addrspacecast (ptr addrspace(1) @_ZL8long_fmt to ptr addrspace(4))
to i64)
+// AMDGCNSPIRV-NEXT: [[TMP8:%.*]] = add i64 [[TMP7]], 1
+// AMDGCNSPIRV-NEXT: br label %[[STRLEN_JOIN]]
+// AMDGCNSPIRV: [[STRLEN_JOIN]]:
+// AMDGCNSPIRV-NEXT: [[TMP9:%.*]] = phi i64 [ [[TMP8]],
%[[STRLEN_WHILE_DONE]] ], [ 0, %[[ENTRY]] ]
+// AMDGCNSPIRV-NEXT: [[TMP10:%.*]] = call addrspace(4) i64
@__ockl_printf_append_string_n(i64 [[TMP0]], ptr addrspace(4) addrspacecast
(ptr addrspace(1) @_ZL8long_fmt to ptr addrspace(4)), i64 [[TMP9]], i32 0)
+// AMDGCNSPIRV-NEXT: [[TMP11:%.*]] = icmp eq ptr addrspace(4) addrspacecast
(ptr addrspace(1) @.str.5 to ptr addrspace(4)), null
+// AMDGCNSPIRV-NEXT: br i1 [[TMP11]], label %[[STRLEN_JOIN1:.*]], label
%[[STRLEN_WHILE2:.*]]
+// AMDGCNSPIRV: [[STRLEN_WHILE2]]:
+// AMDGCNSPIRV-NEXT: [[TMP12:%.*]] = phi ptr addrspace(4) [ addrspacecast
(ptr addrspace(1) @.str.5 to ptr addrspace(4)), %[[STRLEN_JOIN]] ], [
[[TMP13:%.*]], %[[STRLEN_WHILE2]] ]
+// AMDGCNSPIRV-NEXT: [[TMP13]] = getelementptr i8, ptr addrspace(4)
[[TMP12]], i64 1
+// AMDGCNSPIRV-NEXT: [[TMP14:%.*]] = load i8, ptr addrspace(4) [[TMP12]],
align 1
+// AMDGCNSPIRV-NEXT: [[TMP15:%.*]] = icmp eq i8 [[TMP14]], 0
+// AMDGCNSPIRV-NEXT: br i1 [[TMP15]], label %[[STRLEN_WHILE_DONE3:.*]],
label %[[STRLEN_WHILE2]]
+// AMDGCNSPIRV: [[STRLEN_WHILE_DONE3]]:
+// AMDGCNSPIRV-NEXT: [[TMP16:%.*]] = ptrtoaddr ptr addrspace(4) [[TMP12]]
to i64
+// AMDGCNSPIRV-NEXT: [[TMP17:%.*]] = sub i64 [[TMP16]], ptrtoaddr (ptr
addrspace(4) addrspacecast (ptr addrspace(1) @.str.5 to ptr addrspace(4)) to
i64)
+// AMDGCNSPIRV-NEXT: [[TMP18:%.*]] = add i64 [[TMP17]], 1
+// AMDGCNSPIRV-NEXT: br label %[[STRLEN_JOIN1]]
+// AMDGCNSPIRV: [[STRLEN_JOIN1]]:
+// AMDGCNSPIRV-NEXT: [[TMP19:%.*]] = phi i64 [ [[TMP18]],
%[[STRLEN_WHILE_DONE3]] ], [ 0, %[[STRLEN_JOIN]] ]
+// AMDGCNSPIRV-NEXT: [[TMP20:%.*]] = call addrspace(4) i64
@__ockl_printf_append_string_n(i64 [[TMP10]], ptr addrspace(4) addrspacecast
(ptr addrspace(1) @.str.5 to ptr addrspace(4)), i64 [[TMP19]], i32 1)
+// AMDGCNSPIRV-NEXT: [[TMP21:%.*]] = trunc i64 [[TMP20]] to i32
+// AMDGCNSPIRV-NEXT: ret i32 [[TMP21]]
+//
+__device__ int long_fmt_array() { return printf(long_fmt, "z"); }
+
+template <class T> struct TS { static constexpr char fmt[] = "%s\n"; };
+
+// AMDGCN-LABEL: define dso_local noundef i32 @_Z15tmpl_member_fmtv(
+// AMDGCN-SAME: ) #[[ATTR0]] {
+// AMDGCN-NEXT: [[ENTRY:.*]]:
+// AMDGCN-NEXT: [[TMP0:%.*]] = call i64 @__ockl_printf_begin(i64 0)
+// AMDGCN-NEXT: [[TMP1:%.*]] = icmp eq ptr addrspacecast (ptr addrspace(4)
@_ZN2TSIiE3fmtE to ptr), null
+// AMDGCN-NEXT: br i1 [[TMP1]], label %[[STRLEN_JOIN:.*]], label
%[[STRLEN_WHILE:.*]]
+// AMDGCN: [[STRLEN_WHILE]]:
+// AMDGCN-NEXT: [[TMP2:%.*]] = phi ptr [ addrspacecast (ptr addrspace(4)
@_ZN2TSIiE3fmtE to ptr), %[[ENTRY]] ], [ [[TMP3:%.*]], %[[STRLEN_WHILE]] ]
+// AMDGCN-NEXT: [[TMP3]] = getelementptr i8, ptr [[TMP2]], i64 1
+// AMDGCN-NEXT: [[TMP4:%.*]] = load i8, ptr [[TMP2]], align 1
+// AMDGCN-NEXT: [[TMP5:%.*]] = icmp eq i8 [[TMP4]], 0
+// AMDGCN-NEXT: br i1 [[TMP5]], label %[[STRLEN_WHILE_DONE:.*]], label
%[[STRLEN_WHILE]]
+// AMDGCN: [[STRLEN_WHILE_DONE]]:
+// AMDGCN-NEXT: [[TMP6:%.*]] = ptrtoaddr ptr [[TMP2]] to i64
+// AMDGCN-NEXT: [[TMP7:%.*]] = sub i64 [[TMP6]], ptrtoaddr (ptr
addrspacecast (ptr addrspace(4) @_ZN2TSIiE3fmtE to ptr) to i64)
+// AMDGCN-NEXT: [[TMP8:%.*]] = add i64 [[TMP7]], 1
+// AMDGCN-NEXT: br label %[[STRLEN_JOIN]]
+// AMDGCN: [[STRLEN_JOIN]]:
+// AMDGCN-NEXT: [[TMP9:%.*]] = phi i64 [ [[TMP8]], %[[STRLEN_WHILE_DONE]]
], [ 0, %[[ENTRY]] ]
+// AMDGCN-NEXT: [[TMP10:%.*]] = call i64 @__ockl_printf_append_string_n(i64
[[TMP0]], ptr addrspacecast (ptr addrspace(4) @_ZN2TSIiE3fmtE to ptr), i64
[[TMP9]], i32 0)
+// AMDGCN-NEXT: [[TMP11:%.*]] = icmp eq ptr addrspacecast (ptr addrspace(4)
@.str.6 to ptr), null
+// AMDGCN-NEXT: br i1 [[TMP11]], label %[[STRLEN_JOIN1:.*]], label
%[[STRLEN_WHILE2:.*]]
+// AMDGCN: [[STRLEN_WHILE2]]:
+// AMDGCN-NEXT: [[TMP12:%.*]] = phi ptr [ addrspacecast (ptr addrspace(4)
@.str.6 to ptr), %[[STRLEN_JOIN]] ], [ [[TMP13:%.*]], %[[STRLEN_WHILE2]] ]
+// AMDGCN-NEXT: [[TMP13]] = getelementptr i8, ptr [[TMP12]], i64 1
+// AMDGCN-NEXT: [[TMP14:%.*]] = load i8, ptr [[TMP12]], align 1
+// AMDGCN-NEXT: [[TMP15:%.*]] = icmp eq i8 [[TMP14]], 0
+// AMDGCN-NEXT: br i1 [[TMP15]], label %[[STRLEN_WHILE_DONE3:.*]], label
%[[STRLEN_WHILE2]]
+// AMDGCN: [[STRLEN_WHILE_DONE3]]:
+// AMDGCN-NEXT: [[TMP16:%.*]] = ptrtoaddr ptr [[TMP12]] to i64
+// AMDGCN-NEXT: [[TMP17:%.*]] = sub i64 [[TMP16]], ptrtoaddr (ptr
addrspacecast (ptr addrspace(4) @.str.6 to ptr) to i64)
+// AMDGCN-NEXT: [[TMP18:%.*]] = add i64 [[TMP17]], 1
+// AMDGCN-NEXT: br label %[[STRLEN_JOIN1]]
+// AMDGCN: [[STRLEN_JOIN1]]:
+// AMDGCN-NEXT: [[TMP19:%.*]] = phi i64 [ [[TMP18]],
%[[STRLEN_WHILE_DONE3]] ], [ 0, %[[STRLEN_JOIN]] ]
+// AMDGCN-NEXT: [[TMP20:%.*]] = call i64 @__ockl_printf_append_string_n(i64
[[TMP10]], ptr addrspacecast (ptr addrspace(4) @.str.6 to ptr), i64 [[TMP19]],
i32 1)
+// AMDGCN-NEXT: [[TMP21:%.*]] = trunc i64 [[TMP20]] to i32
+// AMDGCN-NEXT: ret i32 [[TMP21]]
+//
+// AMDGCNSPIRV-LABEL: define spir_func noundef i32 @_Z15tmpl_member_fmtv(
+// AMDGCNSPIRV-SAME: ) addrspace(4) #[[ATTR0]] {
+// AMDGCNSPIRV-NEXT: [[ENTRY:.*]]:
+// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = call addrspace(4) i64
@__ockl_printf_begin(i64 0)
+// AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = icmp eq ptr addrspace(4) addrspacecast
(ptr addrspace(1) @_ZN2TSIiE3fmtE to ptr addrspace(4)), null
+// AMDGCNSPIRV-NEXT: br i1 [[TMP1]], label %[[STRLEN_JOIN:.*]], label
%[[STRLEN_WHILE:.*]]
+// AMDGCNSPIRV: [[STRLEN_WHILE]]:
+// AMDGCNSPIRV-NEXT: [[TMP2:%.*]] = phi ptr addrspace(4) [ addrspacecast
(ptr addrspace(1) @_ZN2TSIiE3fmtE to ptr addrspace(4)), %[[ENTRY]] ], [
[[TMP3:%.*]], %[[STRLEN_WHILE]] ]
+// AMDGCNSPIRV-NEXT: [[TMP3]] = getelementptr i8, ptr addrspace(4)
[[TMP2]], i64 1
+// AMDGCNSPIRV-NEXT: [[TMP4:%.*]] = load i8, ptr addrspace(4) [[TMP2]],
align 1
+// AMDGCNSPIRV-NEXT: [[TMP5:%.*]] = icmp eq i8 [[TMP4]], 0
+// AMDGCNSPIRV-NEXT: br i1 [[TMP5]], label %[[STRLEN_WHILE_DONE:.*]], label
%[[STRLEN_WHILE]]
+// AMDGCNSPIRV: [[STRLEN_WHILE_DONE]]:
+// AMDGCNSPIRV-NEXT: [[TMP6:%.*]] = ptrtoaddr ptr addrspace(4) [[TMP2]] to
i64
+// AMDGCNSPIRV-NEXT: [[TMP7:%.*]] = sub i64 [[TMP6]], ptrtoaddr (ptr
addrspace(4) addrspacecast (ptr addrspace(1) @_ZN2TSIiE3fmtE to ptr
addrspace(4)) to i64)
+// AMDGCNSPIRV-NEXT: [[TMP8:%.*]] = add i64 [[TMP7]], 1
+// AMDGCNSPIRV-NEXT: br label %[[STRLEN_JOIN]]
+// AMDGCNSPIRV: [[STRLEN_JOIN]]:
+// AMDGCNSPIRV-NEXT: [[TMP9:%.*]] = phi i64 [ [[TMP8]],
%[[STRLEN_WHILE_DONE]] ], [ 0, %[[ENTRY]] ]
+// AMDGCNSPIRV-NEXT: [[TMP10:%.*]] = call addrspace(4) i64
@__ockl_printf_append_string_n(i64 [[TMP0]], ptr addrspace(4) addrspacecast
(ptr addrspace(1) @_ZN2TSIiE3fmtE to ptr addrspace(4)), i64 [[TMP9]], i32 0)
+// AMDGCNSPIRV-NEXT: [[TMP11:%.*]] = icmp eq ptr addrspace(4) addrspacecast
(ptr addrspace(1) @.str.6 to ptr addrspace(4)), null
+// AMDGCNSPIRV-NEXT: br i1 [[TMP11]], label %[[STRLEN_JOIN1:.*]], label
%[[STRLEN_WHILE2:.*]]
+// AMDGCNSPIRV: [[STRLEN_WHILE2]]:
+// AMDGCNSPIRV-NEXT: [[TMP12:%.*]] = phi ptr addrspace(4) [ addrspacecast
(ptr addrspace(1) @.str.6 to ptr addrspace(4)), %[[STRLEN_JOIN]] ], [
[[TMP13:%.*]], %[[STRLEN_WHILE2]] ]
+// AMDGCNSPIRV-NEXT: [[TMP13]] = getelementptr i8, ptr addrspace(4)
[[TMP12]], i64 1
+// AMDGCNSPIRV-NEXT: [[TMP14:%.*]] = load i8, ptr addrspace(4) [[TMP12]],
align 1
+// AMDGCNSPIRV-NEXT: [[TMP15:%.*]] = icmp eq i8 [[TMP14]], 0
+// AMDGCNSPIRV-NEXT: br i1 [[TMP15]], label %[[STRLEN_WHILE_DONE3:.*]],
label %[[STRLEN_WHILE2]]
+// AMDGCNSPIRV: [[STRLEN_WHILE_DONE3]]:
+// AMDGCNSPIRV-NEXT: [[TMP16:%.*]] = ptrtoaddr ptr addrspace(4) [[TMP12]]
to i64
+// AMDGCNSPIRV-NEXT: [[TMP17:%.*]] = sub i64 [[TMP16]], ptrtoaddr (ptr
addrspace(4) addrspacecast (ptr addrspace(1) @.str.6 to ptr addrspace(4)) to
i64)
+// AMDGCNSPIRV-NEXT: [[TMP18:%.*]] = add i64 [[TMP17]], 1
+// AMDGCNSPIRV-NEXT: br label %[[STRLEN_JOIN1]]
+// AMDGCNSPIRV: [[STRLEN_JOIN1]]:
+// AMDGCNSPIRV-NEXT: [[TMP19:%.*]] = phi i64 [ [[TMP18]],
%[[STRLEN_WHILE_DONE3]] ], [ 0, %[[STRLEN_JOIN]] ]
+// AMDGCNSPIRV-NEXT: [[TMP20:%.*]] = call addrspace(4) i64
@__ockl_printf_append_string_n(i64 [[TMP10]], ptr addrspace(4) addrspacecast
(ptr addrspace(1) @.str.6 to ptr addrspace(4)), i64 [[TMP19]], i32 1)
+// AMDGCNSPIRV-NEXT: [[TMP21:%.*]] = trunc i64 [[TMP20]] to i32
+// AMDGCNSPIRV-NEXT: ret i32 [[TMP21]]
+//
+__device__ int tmpl_member_fmt() { return printf(TS<int>::fmt, "t"); }
+
+// Already declared with a type different from its initializer's when printf
+// is emitted, so its definition is not pulled forward and the format is
+// treated as non-constant.
+__device__ int sink(const char *, int);
+const char decl_fmt[64] = {'%', 's', '\n'};
+
+// AMDGCN-LABEL: define dso_local noundef i32 @_Z15predeclared_fmtv(
+// AMDGCN-SAME: ) #[[ATTR0]] {
+// AMDGCN-NEXT: [[ENTRY:.*]]:
+// AMDGCN-NEXT: [[TMP0:%.*]] = call i64 @__ockl_printf_begin(i64 0)
+// AMDGCN-NEXT: [[TMP1:%.*]] = icmp eq ptr addrspacecast (ptr addrspace(4)
@_ZL8decl_fmt to ptr), null
+// AMDGCN-NEXT: br i1 [[TMP1]], label %[[STRLEN_JOIN:.*]], label
%[[STRLEN_WHILE:.*]]
+// AMDGCN: [[STRLEN_WHILE]]:
+// AMDGCN-NEXT: [[TMP2:%.*]] = phi ptr [ addrspacecast (ptr addrspace(4)
@_ZL8decl_fmt to ptr), %[[ENTRY]] ], [ [[TMP3:%.*]], %[[STRLEN_WHILE]] ]
+// AMDGCN-NEXT: [[TMP3]] = getelementptr i8, ptr [[TMP2]], i64 1
+// AMDGCN-NEXT: [[TMP4:%.*]] = load i8, ptr [[TMP2]], align 1
+// AMDGCN-NEXT: [[TMP5:%.*]] = icmp eq i8 [[TMP4]], 0
+// AMDGCN-NEXT: br i1 [[TMP5]], label %[[STRLEN_WHILE_DONE:.*]], label
%[[STRLEN_WHILE]]
+// AMDGCN: [[STRLEN_WHILE_DONE]]:
+// AMDGCN-NEXT: [[TMP6:%.*]] = ptrtoaddr ptr [[TMP2]] to i64
+// AMDGCN-NEXT: [[TMP7:%.*]] = sub i64 [[TMP6]], ptrtoaddr (ptr
addrspacecast (ptr addrspace(4) @_ZL8decl_fmt to ptr) to i64)
+// AMDGCN-NEXT: [[TMP8:%.*]] = add i64 [[TMP7]], 1
+// AMDGCN-NEXT: br label %[[STRLEN_JOIN]]
+// AMDGCN: [[STRLEN_JOIN]]:
+// AMDGCN-NEXT: [[TMP9:%.*]] = phi i64 [ [[TMP8]], %[[STRLEN_WHILE_DONE]]
], [ 0, %[[ENTRY]] ]
+// AMDGCN-NEXT: [[TMP10:%.*]] = call i64 @__ockl_printf_append_string_n(i64
[[TMP0]], ptr addrspacecast (ptr addrspace(4) @_ZL8decl_fmt to ptr), i64
[[TMP9]], i32 0)
+// AMDGCN-NEXT: [[TMP11:%.*]] = call i64 @__ockl_printf_append_args(i64
[[TMP10]], i32 1, i64 ptrtoint (ptr addrspacecast (ptr addrspace(4) @.str.7 to
ptr) to i64), i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 1)
+// AMDGCN-NEXT: [[TMP12:%.*]] = trunc i64 [[TMP11]] to i32
+// AMDGCN-NEXT: [[CALL:%.*]] = call noundef i32 @_Z4sinkPKci(ptr noundef
addrspacecast (ptr addrspace(4) @_ZL8decl_fmt to ptr), i32 noundef [[TMP12]])
#[[ATTR2:[0-9]+]]
+// AMDGCN-NEXT: ret i32 [[CALL]]
+//
+// AMDGCNSPIRV-LABEL: define spir_func noundef i32 @_Z15predeclared_fmtv(
+// AMDGCNSPIRV-SAME: ) addrspace(4) #[[ATTR0]] {
+// AMDGCNSPIRV-NEXT: [[ENTRY:.*]]:
+// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = call addrspace(4) i64
@__ockl_printf_begin(i64 0)
+// AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = icmp eq ptr addrspace(4) addrspacecast
(ptr addrspace(1) @_ZL8decl_fmt to ptr addrspace(4)), null
+// AMDGCNSPIRV-NEXT: br i1 [[TMP1]], label %[[STRLEN_JOIN:.*]], label
%[[STRLEN_WHILE:.*]]
+// AMDGCNSPIRV: [[STRLEN_WHILE]]:
+// AMDGCNSPIRV-NEXT: [[TMP2:%.*]] = phi ptr addrspace(4) [ addrspacecast
(ptr addrspace(1) @_ZL8decl_fmt to ptr addrspace(4)), %[[ENTRY]] ], [
[[TMP3:%.*]], %[[STRLEN_WHILE]] ]
+// AMDGCNSPIRV-NEXT: [[TMP3]] = getelementptr i8, ptr addrspace(4)
[[TMP2]], i64 1
+// AMDGCNSPIRV-NEXT: [[TMP4:%.*]] = load i8, ptr addrspace(4) [[TMP2]],
align 1
+// AMDGCNSPIRV-NEXT: [[TMP5:%.*]] = icmp eq i8 [[TMP4]], 0
+// AMDGCNSPIRV-NEXT: br i1 [[TMP5]], label %[[STRLEN_WHILE_DONE:.*]], label
%[[STRLEN_WHILE]]
+// AMDGCNSPIRV: [[STRLEN_WHILE_DONE]]:
+// AMDGCNSPIRV-NEXT: [[TMP6:%.*]] = ptrtoaddr ptr addrspace(4) [[TMP2]] to
i64
+// AMDGCNSPIRV-NEXT: [[TMP7:%.*]] = sub i64 [[TMP6]], ptrtoaddr (ptr
addrspace(4) addrspacecast (ptr addrspace(1) @_ZL8decl_fmt to ptr addrspace(4))
to i64)
+// AMDGCNSPIRV-NEXT: [[TMP8:%.*]] = add i64 [[TMP7]], 1
+// AMDGCNSPIRV-NEXT: br label %[[STRLEN_JOIN]]
+// AMDGCNSPIRV: [[STRLEN_JOIN]]:
+// AMDGCNSPIRV-NEXT: [[TMP9:%.*]] = phi i64 [ [[TMP8]],
%[[STRLEN_WHILE_DONE]] ], [ 0, %[[ENTRY]] ]
+// AMDGCNSPIRV-NEXT: [[TMP10:%.*]] = call addrspace(4) i64
@__ockl_printf_append_string_n(i64 [[TMP0]], ptr addrspace(4) addrspacecast
(ptr addrspace(1) @_ZL8decl_fmt to ptr addrspace(4)), i64 [[TMP9]], i32 0)
+// AMDGCNSPIRV-NEXT: [[TMP11:%.*]] = call addrspace(4) i64
@__ockl_printf_append_args(i64 [[TMP10]], i32 1, i64 ptrtoint (ptr addrspace(4)
addrspacecast (ptr addrspace(1) @.str.7 to ptr addrspace(4)) to i64), i64 0,
i64 0, i64 0, i64 0, i64 0, i64 0, i32 1)
+// AMDGCNSPIRV-NEXT: [[TMP12:%.*]] = trunc i64 [[TMP11]] to i32
+// AMDGCNSPIRV-NEXT: [[CALL:%.*]] = call spir_func noundef addrspace(4) i32
@_Z4sinkPKci(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1)
@_ZL8decl_fmt to ptr addrspace(4)), i32 noundef [[TMP12]]) #[[ATTR2:[0-9]+]]
+// AMDGCNSPIRV-NEXT: ret i32 [[CALL]]
+//
+__device__ int predeclared_fmt() { return sink(decl_fmt, printf(decl_fmt,
"w")); }
diff --git a/clang/test/CodeGenHIP/printf_nonhostcall.cpp
b/clang/test/CodeGenHIP/printf_nonhostcall.cpp
index 7e8a468a335271..941299741f9f1d 100644
--- a/clang/test/CodeGenHIP/printf_nonhostcall.cpp
+++ b/clang/test/CodeGenHIP/printf_nonhostcall.cpp
@@ -12,6 +12,7 @@
extern "C" __device__ int printf(const char *format, ...);
// CHECK-LABEL: define dso_local noundef i32 @_Z4foo1v
+// CHECK-SAME: () #[[ATTR0:[0-9]+]] {
// CHECK-NEXT: entry:
// CHECK-NEXT: [[S:%.*]] = alloca ptr, align 8, addrspace(5)
// CHECK-NEXT: [[S_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[S]] to
ptr
@@ -67,6 +68,7 @@ extern "C" __device__ int printf(const char *format, ...);
// CHECK-NEXT: br label [[END_BLOCK]]
//
// CHECK_CONSTRAINED-LABEL: define dso_local noundef i32 @_Z4foo1v
+// CHECK_CONSTRAINED-SAME: () #[[ATTR0:[0-9]+]] {
// CHECK_CONSTRAINED-NEXT: entry:
// CHECK_CONSTRAINED-NEXT: [[S:%.*]] = alloca ptr, align 8, addrspace(5)
// CHECK_CONSTRAINED-NEXT: [[S_ASCAST:%.*]] = addrspacecast ptr
addrspace(5) [[S]] to ptr
@@ -129,6 +131,7 @@ __device__ int foo1() {
__device__ char *dstr;
__device__ const
// CHECK-LABEL: define dso_local noundef i32 @_Z4foo2v
+// CHECK-SAME: () #[[ATTR0]] {
// CHECK-NEXT: entry:
// CHECK-NEXT: [[LCVAL:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK-NEXT: [[LCVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[LCVAL]] to ptr
@@ -179,6 +182,7 @@ __device__ const
// CHECK-NEXT: br label [[END_BLOCK]]
//
// CHECK_CONSTRAINED-LABEL: define dso_local noundef i32 @_Z4foo2v
+// CHECK_CONSTRAINED-SAME: () #[[ATTR0]] {
// CHECK_CONSTRAINED-NEXT: entry:
// CHECK_CONSTRAINED-NEXT: [[LCVAL:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK_CONSTRAINED-NEXT: [[LCVAL_ASCAST:%.*]] = addrspacecast ptr
addrspace(5) [[LCVAL]] to ptr
@@ -246,6 +250,7 @@ __device__ _BitInt(44) Int44 = 312;
__device__ _BitInt(128) Int128 = 45637;
// CHECK-LABEL: define dso_local noundef i32 @_Z4foo3v
+// CHECK-SAME: () #[[ATTR0]] {
// CHECK-NEXT: entry:
// 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
@@ -305,6 +310,7 @@ __device__ _BitInt(128) Int128 = 45637;
// CHECK-NEXT: br label [[END_BLOCK]]
//
// CHECK_CONSTRAINED-LABEL: define dso_local noundef i32 @_Z4foo3v
+// CHECK_CONSTRAINED-SAME: () #[[ATTR0]] {
// CHECK_CONSTRAINED-NEXT: entry:
// CHECK_CONSTRAINED-NEXT: store i32 25, ptr addrspacecast (ptr
addrspace(3) @_ZZ4foo3vE1s to ptr), align 4
// CHECK_CONSTRAINED-NEXT: [[TMP0:%.*]] = load i32, ptr addrspacecast (ptr
addrspace(3) @_ZZ4foo3vE1s to ptr), align 4
@@ -371,6 +377,7 @@ __device__ int foo3() {
//A non trivial case,
// CHECK-LABEL: define dso_local noundef i32 @_Z4foo4v
+// CHECK-SAME: () #[[ATTR0]] {
// CHECK-NEXT: entry:
// CHECK-NEXT: [[S:%.*]] = alloca ptr, align 8, addrspace(5)
// CHECK-NEXT: [[S_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[S]] to
ptr
@@ -414,6 +421,7 @@ __device__ int foo3() {
// CHECK-NEXT: br label [[END_BLOCK]]
//
// CHECK_CONSTRAINED-LABEL: define dso_local noundef i32 @_Z4foo4v
+// CHECK_CONSTRAINED-SAME: () #[[ATTR0]] {
// CHECK_CONSTRAINED-NEXT: entry:
// CHECK_CONSTRAINED-NEXT: [[S:%.*]] = alloca ptr, align 8, addrspace(5)
// CHECK_CONSTRAINED-NEXT: [[S_ASCAST:%.*]] = addrspacecast ptr
addrspace(5) [[S]] to ptr
@@ -460,3 +468,321 @@ __device__ int foo4() {
const char* s = "format str%d";
return printf(s, 10);
}
+
+// Constant format and %s strings in globals whose definitions are deferred.
+const char fmt_arr[] = {'%', 's', '\n', 0};
+const char arg_arr[] = {'a', 'r', 'g', 0};
+struct S { static constexpr char fmt[] = "%s!\n"; };
+
+// CHECK-LABEL: define dso_local noundef i32 @_Z9fmt_arrayv
+// CHECK-SAME: () #[[ATTR0]] {
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1)
@__printf_alloc(i32 20)
+// CHECK-NEXT: [[TMP0:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]],
null
+// CHECK-NEXT: br i1 [[TMP0]], label [[ARGPUSH_BLOCK:%.*]], label
[[END_BLOCK:%.*]]
+// CHECK: end.block:
+// CHECK-NEXT: [[TMP1:%.*]] = xor i1 [[TMP0]], true
+// CHECK-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP1]] to i32
+// CHECK-NEXT: ret i32 [[PRINTF_RESULT]]
+// CHECK: argpush.block:
+// CHECK-NEXT: store i32 82, ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
+// CHECK-NEXT: [[TMP2:%.*]] = getelementptr inbounds i8, ptr addrspace(1)
[[PRINTF_ALLOC_FN]], i32 4
+// CHECK-NEXT: store i64 -88937716022153584, ptr addrspace(1) [[TMP2]],
align 8
+// CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8, ptr addrspace(1)
[[TMP2]], i32 8
+// CHECK-NEXT: store i32 120, ptr addrspace(1) [[TMP3]], align 4
+// CHECK-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr
addrspace(1) [[TMP3]], i32 4
+// CHECK-NEXT: store i32 0, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 4
+// CHECK-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr
addrspace(1) [[PRINTBUFFNEXTPTR]], i32 4
+// CHECK-NEXT: br label [[END_BLOCK]]
+//
+// CHECK_CONSTRAINED-LABEL: define dso_local noundef i32 @_Z9fmt_arrayv
+// CHECK_CONSTRAINED-SAME: () #[[ATTR0]] {
+// CHECK_CONSTRAINED-NEXT: entry:
+// CHECK_CONSTRAINED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1)
@__printf_alloc(i32 20)
+// CHECK_CONSTRAINED-NEXT: [[TMP0:%.*]] = icmp ne ptr addrspace(1)
[[PRINTF_ALLOC_FN]], null
+// CHECK_CONSTRAINED-NEXT: br i1 [[TMP0]], label [[ARGPUSH_BLOCK:%.*]],
label [[END_BLOCK:%.*]]
+// CHECK_CONSTRAINED: end.block:
+// CHECK_CONSTRAINED-NEXT: [[TMP1:%.*]] = xor i1 [[TMP0]], true
+// CHECK_CONSTRAINED-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP1]] to i32
+// CHECK_CONSTRAINED-NEXT: ret i32 [[PRINTF_RESULT]]
+// CHECK_CONSTRAINED: argpush.block:
+// CHECK_CONSTRAINED-NEXT: store i32 82, ptr addrspace(1)
[[PRINTF_ALLOC_FN]], align 4
+// CHECK_CONSTRAINED-NEXT: [[TMP2:%.*]] = getelementptr inbounds i8, ptr
addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
+// CHECK_CONSTRAINED-NEXT: store i64 -88937716022153584, ptr addrspace(1)
[[TMP2]], align 8
+// CHECK_CONSTRAINED-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8, ptr
addrspace(1) [[TMP2]], i32 8
+// CHECK_CONSTRAINED-NEXT: store i32 120, ptr addrspace(1) [[TMP3]], align 4
+// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr
inbounds i8, ptr addrspace(1) [[TMP3]], i32 4
+// CHECK_CONSTRAINED-NEXT: store i32 0, ptr addrspace(1)
[[PRINTBUFFNEXTPTR]], align 4
+// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr
inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 4
+// CHECK_CONSTRAINED-NEXT: br label [[END_BLOCK]]
+//
+__device__ int fmt_array() { return printf(fmt_arr, "x"); }
+
+// CHECK-LABEL: define dso_local noundef i32 @_Z9arg_arrayv
+// CHECK-SAME: () #[[ATTR0]] {
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1)
@__printf_alloc(i32 20)
+// CHECK-NEXT: [[TMP0:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]],
null
+// CHECK-NEXT: br i1 [[TMP0]], label [[ARGPUSH_BLOCK:%.*]], label
[[END_BLOCK:%.*]]
+// CHECK: end.block:
+// CHECK-NEXT: [[TMP1:%.*]] = xor i1 [[TMP0]], true
+// CHECK-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP1]] to i32
+// CHECK-NEXT: ret i32 [[PRINTF_RESULT]]
+// CHECK: argpush.block:
+// CHECK-NEXT: store i32 82, ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
+// CHECK-NEXT: [[TMP2:%.*]] = getelementptr inbounds i8, ptr addrspace(1)
[[PRINTF_ALLOC_FN]], i32 4
+// CHECK-NEXT: store i64 -88937716022153584, ptr addrspace(1) [[TMP2]],
align 8
+// CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8, ptr addrspace(1)
[[TMP2]], i32 8
+// CHECK-NEXT: store i32 6779489, ptr addrspace(1) [[TMP3]], align 4
+// CHECK-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr
addrspace(1) [[TMP3]], i32 4
+// CHECK-NEXT: store i32 0, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 4
+// CHECK-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr
addrspace(1) [[PRINTBUFFNEXTPTR]], i32 4
+// CHECK-NEXT: br label [[END_BLOCK]]
+//
+// CHECK_CONSTRAINED-LABEL: define dso_local noundef i32 @_Z9arg_arrayv
+// CHECK_CONSTRAINED-SAME: () #[[ATTR0]] {
+// CHECK_CONSTRAINED-NEXT: entry:
+// CHECK_CONSTRAINED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1)
@__printf_alloc(i32 20)
+// CHECK_CONSTRAINED-NEXT: [[TMP0:%.*]] = icmp ne ptr addrspace(1)
[[PRINTF_ALLOC_FN]], null
+// CHECK_CONSTRAINED-NEXT: br i1 [[TMP0]], label [[ARGPUSH_BLOCK:%.*]],
label [[END_BLOCK:%.*]]
+// CHECK_CONSTRAINED: end.block:
+// CHECK_CONSTRAINED-NEXT: [[TMP1:%.*]] = xor i1 [[TMP0]], true
+// CHECK_CONSTRAINED-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP1]] to i32
+// CHECK_CONSTRAINED-NEXT: ret i32 [[PRINTF_RESULT]]
+// CHECK_CONSTRAINED: argpush.block:
+// CHECK_CONSTRAINED-NEXT: store i32 82, ptr addrspace(1)
[[PRINTF_ALLOC_FN]], align 4
+// CHECK_CONSTRAINED-NEXT: [[TMP2:%.*]] = getelementptr inbounds i8, ptr
addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
+// CHECK_CONSTRAINED-NEXT: store i64 -88937716022153584, ptr addrspace(1)
[[TMP2]], align 8
+// CHECK_CONSTRAINED-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8, ptr
addrspace(1) [[TMP2]], i32 8
+// CHECK_CONSTRAINED-NEXT: store i32 6779489, ptr addrspace(1) [[TMP3]],
align 4
+// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr
inbounds i8, ptr addrspace(1) [[TMP3]], i32 4
+// CHECK_CONSTRAINED-NEXT: store i32 0, ptr addrspace(1)
[[PRINTBUFFNEXTPTR]], align 4
+// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr
inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 4
+// CHECK_CONSTRAINED-NEXT: br label [[END_BLOCK]]
+//
+__device__ int arg_array() { return printf("%s\n", arg_arr); }
+
+// CHECK-LABEL: define dso_local noundef i32 @_Z10member_fmtv
+// CHECK-SAME: () #[[ATTR0]] {
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1)
@__printf_alloc(i32 20)
+// CHECK-NEXT: [[TMP0:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]],
null
+// CHECK-NEXT: br i1 [[TMP0]], label [[ARGPUSH_BLOCK:%.*]], label
[[END_BLOCK:%.*]]
+// CHECK: end.block:
+// CHECK-NEXT: [[TMP1:%.*]] = xor i1 [[TMP0]], true
+// CHECK-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP1]] to i32
+// CHECK-NEXT: ret i32 [[PRINTF_RESULT]]
+// CHECK: argpush.block:
+// CHECK-NEXT: store i32 82, ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
+// CHECK-NEXT: [[TMP2:%.*]] = getelementptr inbounds i8, ptr addrspace(1)
[[PRINTF_ALLOC_FN]], i32 4
+// CHECK-NEXT: store i64 354905338204825358, ptr addrspace(1) [[TMP2]],
align 8
+// CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8, ptr addrspace(1)
[[TMP2]], i32 8
+// CHECK-NEXT: store i32 121, ptr addrspace(1) [[TMP3]], align 4
+// CHECK-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr
addrspace(1) [[TMP3]], i32 4
+// CHECK-NEXT: store i32 0, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 4
+// CHECK-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr
addrspace(1) [[PRINTBUFFNEXTPTR]], i32 4
+// CHECK-NEXT: br label [[END_BLOCK]]
+//
+// CHECK_CONSTRAINED-LABEL: define dso_local noundef i32 @_Z10member_fmtv
+// CHECK_CONSTRAINED-SAME: () #[[ATTR0]] {
+// CHECK_CONSTRAINED-NEXT: entry:
+// CHECK_CONSTRAINED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1)
@__printf_alloc(i32 20)
+// CHECK_CONSTRAINED-NEXT: [[TMP0:%.*]] = icmp ne ptr addrspace(1)
[[PRINTF_ALLOC_FN]], null
+// CHECK_CONSTRAINED-NEXT: br i1 [[TMP0]], label [[ARGPUSH_BLOCK:%.*]],
label [[END_BLOCK:%.*]]
+// CHECK_CONSTRAINED: end.block:
+// CHECK_CONSTRAINED-NEXT: [[TMP1:%.*]] = xor i1 [[TMP0]], true
+// CHECK_CONSTRAINED-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP1]] to i32
+// CHECK_CONSTRAINED-NEXT: ret i32 [[PRINTF_RESULT]]
+// CHECK_CONSTRAINED: argpush.block:
+// CHECK_CONSTRAINED-NEXT: store i32 82, ptr addrspace(1)
[[PRINTF_ALLOC_FN]], align 4
+// CHECK_CONSTRAINED-NEXT: [[TMP2:%.*]] = getelementptr inbounds i8, ptr
addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
+// CHECK_CONSTRAINED-NEXT: store i64 354905338204825358, ptr addrspace(1)
[[TMP2]], align 8
+// CHECK_CONSTRAINED-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8, ptr
addrspace(1) [[TMP2]], i32 8
+// CHECK_CONSTRAINED-NEXT: store i32 121, ptr addrspace(1) [[TMP3]], align 4
+// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr
inbounds i8, ptr addrspace(1) [[TMP3]], i32 4
+// CHECK_CONSTRAINED-NEXT: store i32 0, ptr addrspace(1)
[[PRINTBUFFNEXTPTR]], align 4
+// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr
inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 4
+// CHECK_CONSTRAINED-NEXT: br label [[END_BLOCK]]
+//
+__device__ int member_fmt() { return printf(S::fmt, "y"); }
+
+// The initializer's type differs from the array type because of the zero tail.
+const char long_fmt[64] = {'%', 's', '\n'};
+
+// CHECK-LABEL: define dso_local noundef i32 @_Z14long_fmt_arrayv
+// CHECK-SAME: () #[[ATTR0]] {
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1)
@__printf_alloc(i32 20)
+// CHECK-NEXT: [[TMP0:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]],
null
+// CHECK-NEXT: br i1 [[TMP0]], label [[ARGPUSH_BLOCK:%.*]], label
[[END_BLOCK:%.*]]
+// CHECK: end.block:
+// CHECK-NEXT: [[TMP1:%.*]] = xor i1 [[TMP0]], true
+// CHECK-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP1]] to i32
+// CHECK-NEXT: ret i32 [[PRINTF_RESULT]]
+// CHECK: argpush.block:
+// CHECK-NEXT: store i32 82, ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
+// CHECK-NEXT: [[TMP2:%.*]] = getelementptr inbounds i8, ptr addrspace(1)
[[PRINTF_ALLOC_FN]], i32 4
+// CHECK-NEXT: store i64 -88937716022153584, ptr addrspace(1) [[TMP2]],
align 8
+// CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8, ptr addrspace(1)
[[TMP2]], i32 8
+// CHECK-NEXT: store i32 122, ptr addrspace(1) [[TMP3]], align 4
+// CHECK-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr
addrspace(1) [[TMP3]], i32 4
+// CHECK-NEXT: store i32 0, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 4
+// CHECK-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr
addrspace(1) [[PRINTBUFFNEXTPTR]], i32 4
+// CHECK-NEXT: br label [[END_BLOCK]]
+//
+// CHECK_CONSTRAINED-LABEL: define dso_local noundef i32 @_Z14long_fmt_arrayv
+// CHECK_CONSTRAINED-SAME: () #[[ATTR0]] {
+// CHECK_CONSTRAINED-NEXT: entry:
+// CHECK_CONSTRAINED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1)
@__printf_alloc(i32 20)
+// CHECK_CONSTRAINED-NEXT: [[TMP0:%.*]] = icmp ne ptr addrspace(1)
[[PRINTF_ALLOC_FN]], null
+// CHECK_CONSTRAINED-NEXT: br i1 [[TMP0]], label [[ARGPUSH_BLOCK:%.*]],
label [[END_BLOCK:%.*]]
+// CHECK_CONSTRAINED: end.block:
+// CHECK_CONSTRAINED-NEXT: [[TMP1:%.*]] = xor i1 [[TMP0]], true
+// CHECK_CONSTRAINED-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP1]] to i32
+// CHECK_CONSTRAINED-NEXT: ret i32 [[PRINTF_RESULT]]
+// CHECK_CONSTRAINED: argpush.block:
+// CHECK_CONSTRAINED-NEXT: store i32 82, ptr addrspace(1)
[[PRINTF_ALLOC_FN]], align 4
+// CHECK_CONSTRAINED-NEXT: [[TMP2:%.*]] = getelementptr inbounds i8, ptr
addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
+// CHECK_CONSTRAINED-NEXT: store i64 -88937716022153584, ptr addrspace(1)
[[TMP2]], align 8
+// CHECK_CONSTRAINED-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8, ptr
addrspace(1) [[TMP2]], i32 8
+// CHECK_CONSTRAINED-NEXT: store i32 122, ptr addrspace(1) [[TMP3]], align 4
+// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr
inbounds i8, ptr addrspace(1) [[TMP3]], i32 4
+// CHECK_CONSTRAINED-NEXT: store i32 0, ptr addrspace(1)
[[PRINTBUFFNEXTPTR]], align 4
+// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr
inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 4
+// CHECK_CONSTRAINED-NEXT: br label [[END_BLOCK]]
+//
+__device__ int long_fmt_array() { return printf(long_fmt, "z"); }
+
+template <class T> struct TS { static constexpr char fmt[] = "%s\n"; };
+
+// CHECK-LABEL: define dso_local noundef i32 @_Z15tmpl_member_fmtv
+// CHECK-SAME: () #[[ATTR0]] {
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1)
@__printf_alloc(i32 20)
+// CHECK-NEXT: [[TMP0:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]],
null
+// CHECK-NEXT: br i1 [[TMP0]], label [[ARGPUSH_BLOCK:%.*]], label
[[END_BLOCK:%.*]]
+// CHECK: end.block:
+// CHECK-NEXT: [[TMP1:%.*]] = xor i1 [[TMP0]], true
+// CHECK-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP1]] to i32
+// CHECK-NEXT: ret i32 [[PRINTF_RESULT]]
+// CHECK: argpush.block:
+// CHECK-NEXT: store i32 82, ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
+// CHECK-NEXT: [[TMP2:%.*]] = getelementptr inbounds i8, ptr addrspace(1)
[[PRINTF_ALLOC_FN]], i32 4
+// CHECK-NEXT: store i64 -88937716022153584, ptr addrspace(1) [[TMP2]],
align 8
+// CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8, ptr addrspace(1)
[[TMP2]], i32 8
+// CHECK-NEXT: store i32 116, ptr addrspace(1) [[TMP3]], align 4
+// CHECK-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr
addrspace(1) [[TMP3]], i32 4
+// CHECK-NEXT: store i32 0, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 4
+// CHECK-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr
addrspace(1) [[PRINTBUFFNEXTPTR]], i32 4
+// CHECK-NEXT: br label [[END_BLOCK]]
+//
+// CHECK_CONSTRAINED-LABEL: define dso_local noundef i32 @_Z15tmpl_member_fmtv
+// CHECK_CONSTRAINED-SAME: () #[[ATTR0]] {
+// CHECK_CONSTRAINED-NEXT: entry:
+// CHECK_CONSTRAINED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1)
@__printf_alloc(i32 20)
+// CHECK_CONSTRAINED-NEXT: [[TMP0:%.*]] = icmp ne ptr addrspace(1)
[[PRINTF_ALLOC_FN]], null
+// CHECK_CONSTRAINED-NEXT: br i1 [[TMP0]], label [[ARGPUSH_BLOCK:%.*]],
label [[END_BLOCK:%.*]]
+// CHECK_CONSTRAINED: end.block:
+// CHECK_CONSTRAINED-NEXT: [[TMP1:%.*]] = xor i1 [[TMP0]], true
+// CHECK_CONSTRAINED-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP1]] to i32
+// CHECK_CONSTRAINED-NEXT: ret i32 [[PRINTF_RESULT]]
+// CHECK_CONSTRAINED: argpush.block:
+// CHECK_CONSTRAINED-NEXT: store i32 82, ptr addrspace(1)
[[PRINTF_ALLOC_FN]], align 4
+// CHECK_CONSTRAINED-NEXT: [[TMP2:%.*]] = getelementptr inbounds i8, ptr
addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
+// CHECK_CONSTRAINED-NEXT: store i64 -88937716022153584, ptr addrspace(1)
[[TMP2]], align 8
+// CHECK_CONSTRAINED-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8, ptr
addrspace(1) [[TMP2]], i32 8
+// CHECK_CONSTRAINED-NEXT: store i32 116, ptr addrspace(1) [[TMP3]], align 4
+// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr
inbounds i8, ptr addrspace(1) [[TMP3]], i32 4
+// CHECK_CONSTRAINED-NEXT: store i32 0, ptr addrspace(1)
[[PRINTBUFFNEXTPTR]], align 4
+// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr
inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 4
+// CHECK_CONSTRAINED-NEXT: br label [[END_BLOCK]]
+//
+__device__ int tmpl_member_fmt() { return printf(TS<int>::fmt, "t"); }
+
+// Already declared with a type different from its initializer's when printf
+// is emitted, so its definition is not pulled forward and the format is
+// treated as non-constant.
+__device__ int sink(const char *, int);
+const char decl_fmt[64] = {'%', 's', '\n'};
+
+// CHECK-LABEL: define dso_local noundef i32 @_Z15predeclared_fmtv
+// CHECK-SAME: () #[[ATTR0]] {
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[TMP0:%.*]] = icmp eq ptr addrspacecast (ptr addrspace(4)
@_ZL8decl_fmt to ptr), null
+// CHECK-NEXT: br i1 [[TMP0]], label [[STRLEN_JOIN:%.*]], label
[[STRLEN_WHILE:%.*]]
+// CHECK: strlen.while:
+// CHECK-NEXT: [[TMP1:%.*]] = phi ptr [ addrspacecast (ptr addrspace(4)
@_ZL8decl_fmt to ptr), [[ENTRY:%.*]] ], [ [[TMP2:%.*]], [[STRLEN_WHILE]] ]
+// CHECK-NEXT: [[TMP2]] = getelementptr i8, ptr [[TMP1]], i64 1
+// CHECK-NEXT: [[TMP3:%.*]] = load i8, ptr [[TMP1]], align 1
+// CHECK-NEXT: [[TMP4:%.*]] = icmp eq i8 [[TMP3]], 0
+// CHECK-NEXT: br i1 [[TMP4]], label [[STRLEN_WHILE_DONE:%.*]], label
[[STRLEN_WHILE]]
+// CHECK: strlen.while.done:
+// CHECK-NEXT: [[TMP5:%.*]] = ptrtoaddr ptr [[TMP1]] to i64
+// CHECK-NEXT: [[TMP6:%.*]] = sub i64 [[TMP5]], ptrtoaddr (ptr
addrspacecast (ptr addrspace(4) @_ZL8decl_fmt to ptr) to i64)
+// CHECK-NEXT: [[TMP7:%.*]] = add i64 [[TMP6]], 1
+// CHECK-NEXT: br label [[STRLEN_JOIN]]
+// CHECK: strlen.join:
+// CHECK-NEXT: [[TMP8:%.*]] = phi i64 [ [[TMP7]], [[STRLEN_WHILE_DONE]] ],
[ 0, [[ENTRY]] ]
+// CHECK-NEXT: [[TMP9:%.*]] = add i64 [[TMP8]], 7
+// CHECK-NEXT: [[TMP10:%.*]] = and i64 [[TMP9]], 4294967288
+// CHECK-NEXT: [[TMP11:%.*]] = add i64 [[TMP10]], 12
+// CHECK-NEXT: [[TMP12:%.*]] = trunc i64 [[TMP11]] to i32
+// CHECK-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1)
@__printf_alloc(i32 [[TMP12]])
+// CHECK-NEXT: [[TMP13:%.*]] = icmp ne ptr addrspace(1)
[[PRINTF_ALLOC_FN]], null
+// CHECK-NEXT: br i1 [[TMP13]], label [[ARGPUSH_BLOCK:%.*]], label
[[END_BLOCK:%.*]]
+// CHECK: end.block:
+// CHECK-NEXT: [[TMP14:%.*]] = xor i1 [[TMP13]], true
+// CHECK-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP14]] to i32
+// CHECK-NEXT: [[CALL:%.*]] = call noundef i32 @_Z4sinkPKci(ptr noundef
addrspacecast (ptr addrspace(4) @_ZL8decl_fmt to ptr), i32 noundef
[[PRINTF_RESULT]]) #[[ATTR4:[0-9]+]]
+// CHECK-NEXT: ret i32 [[CALL]]
+// CHECK: argpush.block:
+// CHECK-NEXT: [[TMP15:%.*]] = shl i32 [[TMP12]], 2
+// CHECK-NEXT: store i32 [[TMP15]], ptr addrspace(1) [[PRINTF_ALLOC_FN]],
align 4
+// CHECK-NEXT: [[TMP16:%.*]] = getelementptr inbounds i8, ptr addrspace(1)
[[PRINTF_ALLOC_FN]], i32 4
+// CHECK-NEXT: call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 1
[[TMP16]], ptr align 1 addrspacecast (ptr addrspace(4) @_ZL8decl_fmt to ptr),
i64 [[TMP8]], i1 false)
+// CHECK-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr
addrspace(1) [[TMP16]], i64 [[TMP10]]
+// CHECK-NEXT: store ptr addrspacecast (ptr addrspace(4) @.str.10 to ptr),
ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8
+// CHECK-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr
addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8
+// CHECK-NEXT: br label [[END_BLOCK]]
+//
+// CHECK_CONSTRAINED-LABEL: define dso_local noundef i32 @_Z15predeclared_fmtv
+// CHECK_CONSTRAINED-SAME: () #[[ATTR0]] {
+// CHECK_CONSTRAINED-NEXT: entry:
+// CHECK_CONSTRAINED-NEXT: [[TMP0:%.*]] = icmp eq ptr addrspacecast (ptr
addrspace(4) @_ZL8decl_fmt to ptr), null
+// CHECK_CONSTRAINED-NEXT: br i1 [[TMP0]], label [[STRLEN_JOIN:%.*]], label
[[STRLEN_WHILE:%.*]]
+// CHECK_CONSTRAINED: strlen.while:
+// CHECK_CONSTRAINED-NEXT: [[TMP1:%.*]] = phi ptr [ addrspacecast (ptr
addrspace(4) @_ZL8decl_fmt to ptr), [[ENTRY:%.*]] ], [ [[TMP2:%.*]],
[[STRLEN_WHILE]] ]
+// CHECK_CONSTRAINED-NEXT: [[TMP2]] = getelementptr i8, ptr [[TMP1]], i64 1
+// CHECK_CONSTRAINED-NEXT: [[TMP3:%.*]] = load i8, ptr [[TMP1]], align 1
+// CHECK_CONSTRAINED-NEXT: [[TMP4:%.*]] = icmp eq i8 [[TMP3]], 0
+// CHECK_CONSTRAINED-NEXT: br i1 [[TMP4]], label [[STRLEN_WHILE_DONE:%.*]],
label [[STRLEN_WHILE]]
+// CHECK_CONSTRAINED: strlen.while.done:
+// CHECK_CONSTRAINED-NEXT: [[TMP5:%.*]] = ptrtoaddr ptr [[TMP1]] to i64
+// CHECK_CONSTRAINED-NEXT: [[TMP6:%.*]] = sub i64 [[TMP5]], ptrtoaddr (ptr
addrspacecast (ptr addrspace(4) @_ZL8decl_fmt to ptr) to i64)
+// CHECK_CONSTRAINED-NEXT: [[TMP7:%.*]] = add i64 [[TMP6]], 1
+// CHECK_CONSTRAINED-NEXT: br label [[STRLEN_JOIN]]
+// CHECK_CONSTRAINED: strlen.join:
+// CHECK_CONSTRAINED-NEXT: [[TMP8:%.*]] = phi i64 [ [[TMP7]],
[[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]
+// CHECK_CONSTRAINED-NEXT: [[TMP9:%.*]] = add i64 [[TMP8]], 7
+// CHECK_CONSTRAINED-NEXT: [[TMP10:%.*]] = and i64 [[TMP9]], 4294967288
+// CHECK_CONSTRAINED-NEXT: [[TMP11:%.*]] = add i64 [[TMP10]], 12
+// CHECK_CONSTRAINED-NEXT: [[TMP12:%.*]] = trunc i64 [[TMP11]] to i32
+// CHECK_CONSTRAINED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1)
@__printf_alloc(i32 [[TMP12]])
+// CHECK_CONSTRAINED-NEXT: [[TMP13:%.*]] = icmp ne ptr addrspace(1)
[[PRINTF_ALLOC_FN]], null
+// CHECK_CONSTRAINED-NEXT: br i1 [[TMP13]], label [[ARGPUSH_BLOCK:%.*]],
label [[END_BLOCK:%.*]]
+// CHECK_CONSTRAINED: end.block:
+// CHECK_CONSTRAINED-NEXT: [[TMP14:%.*]] = xor i1 [[TMP13]], true
+// CHECK_CONSTRAINED-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP14]] to i32
+// CHECK_CONSTRAINED-NEXT: [[CALL:%.*]] = call noundef i32 @_Z4sinkPKci(ptr
noundef addrspacecast (ptr addrspace(4) @_ZL8decl_fmt to ptr), i32 noundef
[[PRINTF_RESULT]]) #[[ATTR4:[0-9]+]]
+// CHECK_CONSTRAINED-NEXT: ret i32 [[CALL]]
+// CHECK_CONSTRAINED: argpush.block:
+// CHECK_CONSTRAINED-NEXT: [[TMP15:%.*]] = shl i32 [[TMP12]], 2
+// CHECK_CONSTRAINED-NEXT: store i32 [[TMP15]], ptr addrspace(1)
[[PRINTF_ALLOC_FN]], align 4
+// CHECK_CONSTRAINED-NEXT: [[TMP16:%.*]] = getelementptr inbounds i8, ptr
addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
+// CHECK_CONSTRAINED-NEXT: call void @llvm.memcpy.p1.p0.i64(ptr
addrspace(1) align 1 [[TMP16]], ptr align 1 addrspacecast (ptr addrspace(4)
@_ZL8decl_fmt to ptr), i64 [[TMP8]], i1 false)
+// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr
inbounds i8, ptr addrspace(1) [[TMP16]], i64 [[TMP10]]
+// CHECK_CONSTRAINED-NEXT: store ptr addrspacecast (ptr addrspace(4)
@.str.10 to ptr), ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8
+// CHECK_CONSTRAINED-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr
inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8
+// CHECK_CONSTRAINED-NEXT: br label [[END_BLOCK]]
+//
+__device__ int predeclared_fmt() { return sink(decl_fmt, printf(decl_fmt,
"w")); }
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
