llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang
Author: Steffen Larsen (steffenlarsen)
<details>
<summary>Changes</summary>
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
---
Patch is 61.80 KiB, truncated to 20.00 KiB below, full version:
https://github.com/llvm/llvm-project/pull/231967.diff
5 Files Affected:
- (modified) clang/lib/CodeGen/CGGPUBuiltin.cpp (+32)
- (modified) clang/lib/CodeGen/CodeGenModule.cpp (+23)
- (modified) clang/lib/CodeGen/CodeGenModule.h (+5)
- (modified) clang/test/CodeGenHIP/printf.cpp (+389-2)
- (modified) clang/test/CodeGenHIP/printf_nonhostcall.cpp (+326)
``````````diff
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]], ...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/231967
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
