[llvm] [clang] [AMDGPU] Enable OpenCL hostcall printf (WIP) (PR #72556)
Vikram Hegde via cfe-commits
cfe-commits at lists.llvm.org
Sun Dec 3 22:38:03 PST 2023
https://github.com/vikramRH updated https://github.com/llvm/llvm-project/pull/72556
>From e6ed0c02edc27805fe008332e1e879d558742e94 Mon Sep 17 00:00:00 2001
From: Vikram <Vikram.Hegde at amd.com>
Date: Fri, 10 Nov 2023 09:39:41 +0000
Subject: [PATCH 1/5] [AMDGPU] Treat printf as builtin for OpenCL
---
clang/include/clang/Basic/BuiltinsAMDGPU.def | 8 ++++++++
clang/lib/AST/Decl.cpp | 7 +++++++
clang/lib/Basic/Targets/AMDGPU.cpp | 2 ++
clang/lib/CodeGen/CGBuiltin.cpp | 5 +++++
4 files changed, 22 insertions(+)
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index a19c8bd5f219e..1799c72806bfd 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -21,6 +21,10 @@
#if defined(BUILTIN) && !defined(TARGET_BUILTIN)
# define TARGET_BUILTIN(ID, TYPE, ATTRS, FEATURE) BUILTIN(ID, TYPE, ATTRS)
#endif
+
+#if defined(BUILTIN) && !defined(LANGBUILTIN)
+#define LANGBUILTIN(ID, TYPE, ATTRS, BUILTIN_LANG) BUILTIN(ID, TYPE, ATTRS)
+#endif
//===----------------------------------------------------------------------===//
// SI+ only builtins.
//===----------------------------------------------------------------------===//
@@ -406,5 +410,9 @@ TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_fp8_f32, "iffiIb", "nc", "fp8-insts")
TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_bf8_f32, "ifiiIi", "nc", "fp8-insts")
TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_fp8_f32, "ifiiIi", "nc", "fp8-insts")
+// OpenCL
+LANGBUILTIN(printf, "icC*4.", "fp:0:", ALL_OCL_LANGUAGES)
+
#undef BUILTIN
#undef TARGET_BUILTIN
+#undef LANGBUILTIN
diff --git a/clang/lib/AST/Decl.cpp b/clang/lib/AST/Decl.cpp
index c5c2edf1bfe3a..2597422bdd521 100644
--- a/clang/lib/AST/Decl.cpp
+++ b/clang/lib/AST/Decl.cpp
@@ -49,6 +49,7 @@
#include "clang/Basic/SourceLocation.h"
#include "clang/Basic/SourceManager.h"
#include "clang/Basic/Specifiers.h"
+#include "clang/Basic/TargetBuiltins.h"
#include "clang/Basic/TargetCXXABI.h"
#include "clang/Basic/TargetInfo.h"
#include "clang/Basic/Visibility.h"
@@ -3598,6 +3599,12 @@ unsigned FunctionDecl::getBuiltinID(bool ConsiderWrapperFunctions) const {
if (!ConsiderWrapperFunctions && getStorageClass() == SC_Static)
return 0;
+ // AMDGCN implementation supports printf as a builtin
+ // for OpenCL
+ if (Context.getTargetInfo().getTriple().isAMDGCN() &&
+ Context.getLangOpts().OpenCL && BuiltinID == AMDGPU::BIprintf)
+ return BuiltinID;
+
// OpenCL v1.2 s6.9.f - The library functions defined in
// the C99 standard headers are not available.
if (Context.getLangOpts().OpenCL &&
diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp
index 409ae32ab4242..307cfa49f54e9 100644
--- a/clang/lib/Basic/Targets/AMDGPU.cpp
+++ b/clang/lib/Basic/Targets/AMDGPU.cpp
@@ -91,6 +91,8 @@ static constexpr Builtin::Info BuiltinInfo[] = {
{#ID, TYPE, ATTRS, nullptr, HeaderDesc::NO_HEADER, ALL_LANGUAGES},
#define TARGET_BUILTIN(ID, TYPE, ATTRS, FEATURE) \
{#ID, TYPE, ATTRS, FEATURE, HeaderDesc::NO_HEADER, ALL_LANGUAGES},
+#define LANGBUILTIN(ID, TYPE, ATTRS, LANG) \
+ {#ID, TYPE, ATTRS, nullptr, HeaderDesc::NO_HEADER, LANG},
#include "clang/Basic/BuiltinsAMDGPU.def"
};
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index 65d9862621061..b724a381ae2fb 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -2481,6 +2481,11 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
&getTarget().getLongDoubleFormat() == &llvm::APFloat::IEEEquad())
BuiltinID = mutateLongDoubleBuiltin(BuiltinID);
+ // Mutate the printf builtin ID so that we use the same CodeGen path for
+ // HIP and OpenCL with AMDGPU targets.
+ if (getTarget().getTriple().isAMDGCN() && BuiltinID == AMDGPU::BIprintf)
+ BuiltinID = Builtin::BIprintf;
+
// If the builtin has been declared explicitly with an assembler label,
// disable the specialized emitting below. Ideally we should communicate the
// rename in IR, or at least avoid generating the intrinsic calls that are
>From ba0d4e599800f1a42559fd67e585d161bcf53f22 Mon Sep 17 00:00:00 2001
From: Vikram <Vikram.Hegde at amd.com>
Date: Mon, 20 Nov 2023 05:26:27 +0000
Subject: [PATCH 2/5] [AMDGPU] Enable OpenCL printf expansion at clang CodeGen
---
clang/lib/CodeGen/CGBuiltin.cpp | 3 ++-
clang/lib/CodeGen/CGGPUBuiltin.cpp | 25 +++++++++++++++++++------
clang/lib/Driver/ToolChains/Clang.cpp | 10 ++++++++++
3 files changed, 31 insertions(+), 7 deletions(-)
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index b724a381ae2fb..95f5d477f0146 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -5645,7 +5645,8 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
return EmitOpenMPDevicePrintfCallExpr(E);
if (getTarget().getTriple().isNVPTX())
return EmitNVPTXDevicePrintfCallExpr(E);
- if (getTarget().getTriple().isAMDGCN() && getLangOpts().HIP)
+ if (getTarget().getTriple().isAMDGCN() &&
+ (getLangOpts().HIP || getLangOpts().OpenCL))
return EmitAMDGPUDevicePrintfCallExpr(E);
}
diff --git a/clang/lib/CodeGen/CGGPUBuiltin.cpp b/clang/lib/CodeGen/CGGPUBuiltin.cpp
index e465789a003eb..32b457812af60 100644
--- a/clang/lib/CodeGen/CGGPUBuiltin.cpp
+++ b/clang/lib/CodeGen/CGGPUBuiltin.cpp
@@ -13,6 +13,7 @@
#include "CodeGenFunction.h"
#include "clang/Basic/Builtins.h"
+#include "clang/Basic/TargetBuiltins.h"
#include "llvm/IR/DataLayout.h"
#include "llvm/IR/Instruction.h"
#include "llvm/Support/MathExtras.h"
@@ -177,10 +178,20 @@ RValue CodeGenFunction::EmitNVPTXDevicePrintfCallExpr(const CallExpr *E) {
E, this, GetVprintfDeclaration(CGM.getModule()), false);
}
+// Deterimines if an argument is a string
+static bool isString(const clang::Type *argXTy) {
+ if ((argXTy->isPointerType() || argXTy->isConstantArrayType()) &&
+ argXTy->getPointeeOrArrayElementType()->isCharType())
+ return true;
+ else
+ return false;
+}
+
RValue CodeGenFunction::EmitAMDGPUDevicePrintfCallExpr(const CallExpr *E) {
assert(getTarget().getTriple().getArch() == llvm::Triple::amdgcn);
assert(E->getBuiltinCallee() == Builtin::BIprintf ||
- E->getBuiltinCallee() == Builtin::BI__builtin_printf);
+ E->getBuiltinCallee() == Builtin::BI__builtin_printf ||
+ E->getBuiltinCallee() == AMDGPU::BIprintf);
assert(E->getNumArgs() >= 1); // printf always has at least one arg.
CallArgList CallArgs;
@@ -188,6 +199,8 @@ RValue CodeGenFunction::EmitAMDGPUDevicePrintfCallExpr(const CallExpr *E) {
E->getDirectCallee()->getType()->getAs<FunctionProtoType>(),
E->arguments(), E->getDirectCallee(),
/* ParamsToSkip = */ 0);
+ llvm::IRBuilder<> IRB(Builder.GetInsertBlock(), Builder.GetInsertPoint());
+ IRB.SetCurrentDebugLocation(Builder.getCurrentDebugLocation());
SmallVector<llvm::Value *, 8> Args;
for (const auto &A : CallArgs) {
@@ -198,14 +211,14 @@ RValue CodeGenFunction::EmitAMDGPUDevicePrintfCallExpr(const CallExpr *E) {
}
llvm::Value *Arg = A.getRValue(*this).getScalarVal();
+ if (isString(A.getType().getTypePtr()) && CGM.getLangOpts().OpenCL)
+ Arg = Builder.CreateAddrSpaceCast(Arg, CGM.Int8PtrTy);
Args.push_back(Arg);
}
- llvm::IRBuilder<> IRB(Builder.GetInsertBlock(), Builder.GetInsertPoint());
- IRB.SetCurrentDebugLocation(Builder.getCurrentDebugLocation());
-
- bool isBuffered = (CGM.getTarget().getTargetOpts().AMDGPUPrintfKindVal ==
- clang::TargetOptions::AMDGPUPrintfKind::Buffered);
+ auto PFK = CGM.getTarget().getTargetOpts().AMDGPUPrintfKindVal;
+ bool isBuffered =
+ (PFK == clang::TargetOptions::AMDGPUPrintfKind::Buffered);
auto Printf = llvm::emitAMDGPUPrintfCall(IRB, Args, isBuffered);
Builder.SetInsertPoint(IRB.GetInsertBlock(), IRB.GetInsertPoint());
return RValue::get(Printf);
diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
index f02f7c841b91f..cde8fb13785ac 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -4626,6 +4626,16 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
Args.ClaimAllArgs(options::OPT_gen_cdb_fragment_path);
}
+ if (TC.getTriple().isAMDGPU() && types::isOpenCL(Input.getType())) {
+ if (Args.getLastArg(options::OPT_mprintf_kind_EQ)) {
+ CmdArgs.push_back(Args.MakeArgString(
+ "-mprintf-kind=" +
+ Args.getLastArgValue(options::OPT_mprintf_kind_EQ)));
+ // Force compiler error on invalid conversion specifiers
+ CmdArgs.push_back(Args.MakeArgString("-Werror=format-invalid-specifier"));
+ }
+ }
+
if (IsCuda || IsHIP) {
// We have to pass the triple of the host if compiling for a CUDA/HIP device
// and vice-versa.
>From 2fc06013a5a9ff4ad41d049ade8231aaa03e7f01 Mon Sep 17 00:00:00 2001
From: Vikram <Vikram.Hegde at amd.com>
Date: Mon, 20 Nov 2023 05:28:10 +0000
Subject: [PATCH 3/5] [AMDGPU] Add vector processing support to AMDGPU printf
---
clang/test/CodeGenOpenCL/amdgpu-printf.cl | 205 +++++++++++++++++-
.../lib/Transforms/Utils/AMDGPUEmitPrintf.cpp | 135 +++++++-----
2 files changed, 288 insertions(+), 52 deletions(-)
diff --git a/clang/test/CodeGenOpenCL/amdgpu-printf.cl b/clang/test/CodeGenOpenCL/amdgpu-printf.cl
index 6c84485b66b4a..5ada787f723ec 100644
--- a/clang/test/CodeGenOpenCL/amdgpu-printf.cl
+++ b/clang/test/CodeGenOpenCL/amdgpu-printf.cl
@@ -1,5 +1,6 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
-// RUN: %clang_cc1 -cl-std=CL1.2 -triple amdgcn-amd-amdhsa -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -cl-std=CL1.2 -triple amdgcn-amd-amdhsa -mprintf-kind=buffered -disable-llvm-passes -emit-llvm -o - %s | FileCheck --check-prefix=CHECK_BUFFERED %s
+// RUN: %clang_cc1 -cl-std=CL1.2 -triple amdgcn-amd-amdhsa -mprintf-kind=hostcall -disable-llvm-passes -emit-llvm -o - %s | FileCheck --check-prefix=CHECK_HOSTCALL %s
int printf(__constant const char* st, ...) __attribute__((format(printf, 1, 2)));
@@ -7,6 +8,61 @@ int printf(__constant const char* st, ...) __attribute__((format(printf, 1, 2)))
// CHECK-NEXT: entry:
// CHECK-NEXT: [[CALL:%.*]] = call i32 (ptr addrspace(4), ...) @printf(ptr addrspace(4) noundef @.str) #[[ATTR4:[0-9]+]]
// CHECK-NEXT: ret void
+// CHECK_BUFFERED-LABEL: @test_printf_noargs(
+// CHECK_BUFFERED-NEXT: entry:
+// CHECK_BUFFERED-NEXT: br i1 icmp eq (ptr addrspacecast (ptr addrspace(4) @.str to ptr), ptr null), label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]
+// CHECK_BUFFERED: strlen.while:
+// CHECK_BUFFERED-NEXT: [[TMP0:%.*]] = phi ptr [ addrspacecast (ptr addrspace(4) @.str to ptr), [[ENTRY:%.*]] ], [ [[TMP1:%.*]], [[STRLEN_WHILE]] ]
+// CHECK_BUFFERED-NEXT: [[TMP1]] = getelementptr i8, ptr [[TMP0]], i64 1
+// CHECK_BUFFERED-NEXT: [[TMP2:%.*]] = load i8, ptr [[TMP0]], align 1
+// CHECK_BUFFERED-NEXT: [[TMP3:%.*]] = icmp eq i8 [[TMP2]], 0
+// CHECK_BUFFERED-NEXT: br i1 [[TMP3]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]]
+// CHECK_BUFFERED: strlen.while.done:
+// CHECK_BUFFERED-NEXT: [[TMP4:%.*]] = ptrtoint ptr [[TMP0]] to i64
+// CHECK_BUFFERED-NEXT: [[TMP5:%.*]] = sub i64 [[TMP4]], ptrtoint (ptr addrspacecast (ptr addrspace(4) @.str to ptr) to i64)
+// CHECK_BUFFERED-NEXT: [[TMP6:%.*]] = add i64 [[TMP5]], 1
+// CHECK_BUFFERED-NEXT: br label [[STRLEN_JOIN]]
+// CHECK_BUFFERED: strlen.join:
+// CHECK_BUFFERED-NEXT: [[TMP7:%.*]] = phi i64 [ [[TMP6]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]
+// CHECK_BUFFERED-NEXT: [[TMP8:%.*]] = add i64 [[TMP7]], 7
+// CHECK_BUFFERED-NEXT: [[TMP9:%.*]] = and i64 [[TMP8]], 4294967288
+// CHECK_BUFFERED-NEXT: [[TMP10:%.*]] = add i64 [[TMP9]], 4
+// CHECK_BUFFERED-NEXT: [[TMP11:%.*]] = trunc i64 [[TMP10]] to i32
+// CHECK_BUFFERED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[TMP11]])
+// CHECK_BUFFERED-NEXT: [[TMP12:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
+// CHECK_BUFFERED-NEXT: br i1 [[TMP12]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]
+// CHECK_BUFFERED: end.block:
+// CHECK_BUFFERED-NEXT: [[TMP13:%.*]] = xor i1 [[TMP12]], true
+// CHECK_BUFFERED-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP13]] to i32
+// CHECK_BUFFERED-NEXT: ret void
+// CHECK_BUFFERED: argpush.block:
+// CHECK_BUFFERED-NEXT: [[TMP14:%.*]] = shl i32 [[TMP11]], 2
+// CHECK_BUFFERED-NEXT: store i32 [[TMP14]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
+// CHECK_BUFFERED-NEXT: [[TMP15:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
+// CHECK_BUFFERED-NEXT: call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 1 [[TMP15]], ptr align 1 addrspacecast (ptr addrspace(4) @.str to ptr), i64 [[TMP7]], i1 false)
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP15]], i64 [[TMP9]]
+// CHECK_BUFFERED-NEXT: br label [[END_BLOCK]]
+//
+// CHECK_HOSTCALL-LABEL: @test_printf_noargs(
+// CHECK_HOSTCALL-NEXT: entry:
+// CHECK_HOSTCALL-NEXT: [[TMP0:%.*]] = call i64 @__ockl_printf_begin(i64 0)
+// CHECK_HOSTCALL-NEXT: br i1 icmp eq (ptr addrspacecast (ptr addrspace(4) @.str to ptr), ptr null), label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]
+// CHECK_HOSTCALL: strlen.while:
+// CHECK_HOSTCALL-NEXT: [[TMP1:%.*]] = phi ptr [ addrspacecast (ptr addrspace(4) @.str to ptr), [[ENTRY:%.*]] ], [ [[TMP2:%.*]], [[STRLEN_WHILE]] ]
+// CHECK_HOSTCALL-NEXT: [[TMP2]] = getelementptr i8, ptr [[TMP1]], i64 1
+// CHECK_HOSTCALL-NEXT: [[TMP3:%.*]] = load i8, ptr [[TMP1]], align 1
+// CHECK_HOSTCALL-NEXT: [[TMP4:%.*]] = icmp eq i8 [[TMP3]], 0
+// CHECK_HOSTCALL-NEXT: br i1 [[TMP4]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]]
+// CHECK_HOSTCALL: strlen.while.done:
+// CHECK_HOSTCALL-NEXT: [[TMP5:%.*]] = ptrtoint ptr [[TMP1]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP6:%.*]] = sub i64 [[TMP5]], ptrtoint (ptr addrspacecast (ptr addrspace(4) @.str to ptr) to i64)
+// CHECK_HOSTCALL-NEXT: [[TMP7:%.*]] = add i64 [[TMP6]], 1
+// CHECK_HOSTCALL-NEXT: br label [[STRLEN_JOIN]]
+// CHECK_HOSTCALL: strlen.join:
+// CHECK_HOSTCALL-NEXT: [[TMP8:%.*]] = phi i64 [ [[TMP7]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]
+// CHECK_HOSTCALL-NEXT: [[TMP9:%.*]] = call i64 @__ockl_printf_append_string_n(i64 [[TMP0]], ptr addrspacecast (ptr addrspace(4) @.str to ptr), i64 [[TMP8]], i32 1)
+// CHECK_HOSTCALL-NEXT: [[TMP10:%.*]] = trunc i64 [[TMP9]] to i32
+// CHECK_HOSTCALL-NEXT: ret void
//
__kernel void test_printf_noargs() {
printf("");
@@ -19,6 +75,53 @@ __kernel void test_printf_noargs() {
// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA8]]
// CHECK-NEXT: [[CALL:%.*]] = call i32 (ptr addrspace(4), ...) @printf(ptr addrspace(4) noundef @.str.1, i32 noundef [[TMP0]]) #[[ATTR4]]
// CHECK-NEXT: ret void
+// CHECK_BUFFERED-LABEL: @test_printf_int(
+// CHECK_BUFFERED-NEXT: entry:
+// CHECK_BUFFERED-NEXT: [[I_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK_BUFFERED-NEXT: store i32 [[I:%.*]], ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA12:![0-9]+]]
+// CHECK_BUFFERED-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA12]]
+// CHECK_BUFFERED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 20)
+// CHECK_BUFFERED-NEXT: [[TMP1:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
+// CHECK_BUFFERED-NEXT: br i1 [[TMP1]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]
+// CHECK_BUFFERED: end.block:
+// CHECK_BUFFERED-NEXT: [[TMP2:%.*]] = xor i1 [[TMP1]], true
+// CHECK_BUFFERED-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP2]] to i32
+// CHECK_BUFFERED-NEXT: ret void
+// CHECK_BUFFERED: argpush.block:
+// CHECK_BUFFERED-NEXT: store i32 82, ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
+// CHECK_BUFFERED-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
+// CHECK_BUFFERED-NEXT: store i64 -2582314622382785113, ptr addrspace(1) [[TMP3]], align 8
+// CHECK_BUFFERED-NEXT: [[TMP4:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP3]], i32 8
+// CHECK_BUFFERED-NEXT: [[TMP5:%.*]] = zext i32 [[TMP0]] to i64
+// CHECK_BUFFERED-NEXT: store i64 [[TMP5]], ptr addrspace(1) [[TMP4]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP4]], i32 8
+// CHECK_BUFFERED-NEXT: br label [[END_BLOCK]]
+//
+// CHECK_HOSTCALL-LABEL: @test_printf_int(
+// CHECK_HOSTCALL-NEXT: entry:
+// CHECK_HOSTCALL-NEXT: [[I_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK_HOSTCALL-NEXT: store i32 [[I:%.*]], ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA9:![0-9]+]]
+// CHECK_HOSTCALL-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA9]]
+// CHECK_HOSTCALL-NEXT: [[TMP1:%.*]] = call i64 @__ockl_printf_begin(i64 0)
+// CHECK_HOSTCALL-NEXT: br i1 icmp eq (ptr addrspacecast (ptr addrspace(4) @.str.1 to ptr), ptr null), label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]
+// CHECK_HOSTCALL: strlen.while:
+// CHECK_HOSTCALL-NEXT: [[TMP2:%.*]] = phi ptr [ addrspacecast (ptr addrspace(4) @.str.1 to ptr), [[ENTRY:%.*]] ], [ [[TMP3:%.*]], [[STRLEN_WHILE]] ]
+// CHECK_HOSTCALL-NEXT: [[TMP3]] = getelementptr i8, ptr [[TMP2]], i64 1
+// CHECK_HOSTCALL-NEXT: [[TMP4:%.*]] = load i8, ptr [[TMP2]], align 1
+// CHECK_HOSTCALL-NEXT: [[TMP5:%.*]] = icmp eq i8 [[TMP4]], 0
+// CHECK_HOSTCALL-NEXT: br i1 [[TMP5]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]]
+// CHECK_HOSTCALL: strlen.while.done:
+// CHECK_HOSTCALL-NEXT: [[TMP6:%.*]] = ptrtoint ptr [[TMP2]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP7:%.*]] = sub i64 [[TMP6]], ptrtoint (ptr addrspacecast (ptr addrspace(4) @.str.1 to ptr) to i64)
+// CHECK_HOSTCALL-NEXT: [[TMP8:%.*]] = add i64 [[TMP7]], 1
+// CHECK_HOSTCALL-NEXT: br label [[STRLEN_JOIN]]
+// CHECK_HOSTCALL: strlen.join:
+// CHECK_HOSTCALL-NEXT: [[TMP9:%.*]] = phi i64 [ [[TMP8]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]
+// CHECK_HOSTCALL-NEXT: [[TMP10:%.*]] = call i64 @__ockl_printf_append_string_n(i64 [[TMP1]], ptr addrspacecast (ptr addrspace(4) @.str.1 to ptr), i64 [[TMP9]], i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP11:%.*]] = zext i32 [[TMP0]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP12:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP10]], i32 1, i64 [[TMP11]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 1)
+// CHECK_HOSTCALL-NEXT: [[TMP13:%.*]] = trunc i64 [[TMP12]] to i32
+// CHECK_HOSTCALL-NEXT: ret void
//
__kernel void test_printf_int(int i) {
printf("%d", i);
@@ -43,6 +146,106 @@ __kernel void test_printf_int(int i) {
// CHECK-NEXT: [[CALL:%.*]] = call i32 (ptr addrspace(4), ...) @printf(ptr addrspace(4) noundef @.str.2, ptr addrspace(5) noundef [[ARRAYDECAY]], i32 noundef [[TMP2]]) #[[ATTR4]]
// CHECK-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[S]]) #[[ATTR5]]
// CHECK-NEXT: ret void
+// CHECK_BUFFERED-LABEL: @test_printf_str_int(
+// CHECK_BUFFERED-NEXT: entry:
+// CHECK_BUFFERED-NEXT: [[I_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK_BUFFERED-NEXT: [[S:%.*]] = alloca [4 x i8], align 1, addrspace(5)
+// CHECK_BUFFERED-NEXT: store i32 [[I:%.*]], ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA12]]
+// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[S]]) #[[ATTR1:[0-9]+]]
+// CHECK_BUFFERED-NEXT: call void @llvm.memcpy.p5.p4.i64(ptr addrspace(5) align 1 [[S]], ptr addrspace(4) align 1 @__const.test_printf_str_int.s, i64 4, i1 false)
+// CHECK_BUFFERED-NEXT: [[ARRAYDECAY:%.*]] = getelementptr inbounds [4 x i8], ptr addrspace(5) [[S]], i64 0, i64 0
+// CHECK_BUFFERED-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA12]]
+// CHECK_BUFFERED-NEXT: [[TMP1:%.*]] = addrspacecast ptr addrspace(5) [[ARRAYDECAY]] to ptr
+// CHECK_BUFFERED-NEXT: [[TMP2:%.*]] = icmp eq ptr [[TMP1]], null
+// CHECK_BUFFERED-NEXT: br i1 [[TMP2]], label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]
+// CHECK_BUFFERED: strlen.while:
+// CHECK_BUFFERED-NEXT: [[TMP3:%.*]] = phi ptr [ [[TMP1]], [[ENTRY:%.*]] ], [ [[TMP4:%.*]], [[STRLEN_WHILE]] ]
+// CHECK_BUFFERED-NEXT: [[TMP4]] = getelementptr i8, ptr [[TMP3]], i64 1
+// CHECK_BUFFERED-NEXT: [[TMP5:%.*]] = load i8, ptr [[TMP3]], align 1
+// CHECK_BUFFERED-NEXT: [[TMP6:%.*]] = icmp eq i8 [[TMP5]], 0
+// CHECK_BUFFERED-NEXT: br i1 [[TMP6]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]]
+// CHECK_BUFFERED: strlen.while.done:
+// CHECK_BUFFERED-NEXT: [[TMP7:%.*]] = ptrtoint ptr [[TMP1]] to i64
+// CHECK_BUFFERED-NEXT: [[TMP8:%.*]] = ptrtoint ptr [[TMP3]] to i64
+// CHECK_BUFFERED-NEXT: [[TMP9:%.*]] = sub i64 [[TMP8]], [[TMP7]]
+// CHECK_BUFFERED-NEXT: [[TMP10:%.*]] = add i64 [[TMP9]], 1
+// CHECK_BUFFERED-NEXT: br label [[STRLEN_JOIN]]
+// CHECK_BUFFERED: strlen.join:
+// CHECK_BUFFERED-NEXT: [[TMP11:%.*]] = phi i64 [ [[TMP10]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]
+// CHECK_BUFFERED-NEXT: [[TMP12:%.*]] = add i64 [[TMP11]], 7
+// CHECK_BUFFERED-NEXT: [[TMP13:%.*]] = and i64 [[TMP12]], 4294967288
+// CHECK_BUFFERED-NEXT: [[TMP14:%.*]] = add i64 [[TMP13]], 20
+// CHECK_BUFFERED-NEXT: [[TMP15:%.*]] = trunc i64 [[TMP14]] to i32
+// CHECK_BUFFERED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[TMP15]])
+// CHECK_BUFFERED-NEXT: [[TMP16:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
+// CHECK_BUFFERED-NEXT: br i1 [[TMP16]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]
+// CHECK_BUFFERED: end.block:
+// CHECK_BUFFERED-NEXT: [[TMP17:%.*]] = xor i1 [[TMP16]], true
+// CHECK_BUFFERED-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP17]] to i32
+// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[S]]) #[[ATTR1]]
+// CHECK_BUFFERED-NEXT: ret void
+// CHECK_BUFFERED: argpush.block:
+// CHECK_BUFFERED-NEXT: [[TMP18:%.*]] = shl i32 [[TMP15]], 2
+// CHECK_BUFFERED-NEXT: [[TMP19:%.*]] = or i32 [[TMP18]], 2
+// CHECK_BUFFERED-NEXT: store i32 [[TMP19]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
+// CHECK_BUFFERED-NEXT: [[TMP20:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
+// CHECK_BUFFERED-NEXT: store i64 -2942283388077972797, ptr addrspace(1) [[TMP20]], align 8
+// CHECK_BUFFERED-NEXT: [[TMP21:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP20]], i32 8
+// CHECK_BUFFERED-NEXT: call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 1 [[TMP21]], ptr align 1 [[TMP1]], i64 [[TMP11]], i1 false)
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP21]], i64 [[TMP13]]
+// CHECK_BUFFERED-NEXT: [[TMP22:%.*]] = zext i32 [[TMP0]] to i64
+// CHECK_BUFFERED-NEXT: store i64 [[TMP22]], ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8
+// CHECK_BUFFERED-NEXT: br label [[END_BLOCK]]
+//
+// CHECK_HOSTCALL-LABEL: @test_printf_str_int(
+// CHECK_HOSTCALL-NEXT: entry:
+// CHECK_HOSTCALL-NEXT: [[I_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK_HOSTCALL-NEXT: [[S:%.*]] = alloca [4 x i8], align 1, addrspace(5)
+// CHECK_HOSTCALL-NEXT: store i32 [[I:%.*]], ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA9]]
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[S]]) #[[ATTR3:[0-9]+]]
+// CHECK_HOSTCALL-NEXT: call void @llvm.memcpy.p5.p4.i64(ptr addrspace(5) align 1 [[S]], ptr addrspace(4) align 1 @__const.test_printf_str_int.s, i64 4, i1 false)
+// CHECK_HOSTCALL-NEXT: [[ARRAYDECAY:%.*]] = getelementptr inbounds [4 x i8], ptr addrspace(5) [[S]], i64 0, i64 0
+// CHECK_HOSTCALL-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA9]]
+// CHECK_HOSTCALL-NEXT: [[TMP1:%.*]] = addrspacecast ptr addrspace(5) [[ARRAYDECAY]] to ptr
+// CHECK_HOSTCALL-NEXT: [[TMP2:%.*]] = call i64 @__ockl_printf_begin(i64 0)
+// CHECK_HOSTCALL-NEXT: br i1 icmp eq (ptr addrspacecast (ptr addrspace(4) @.str.2 to ptr), ptr null), label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]
+// CHECK_HOSTCALL: strlen.while:
+// CHECK_HOSTCALL-NEXT: [[TMP3:%.*]] = phi ptr [ addrspacecast (ptr addrspace(4) @.str.2 to ptr), [[ENTRY:%.*]] ], [ [[TMP4:%.*]], [[STRLEN_WHILE]] ]
+// CHECK_HOSTCALL-NEXT: [[TMP4]] = getelementptr i8, ptr [[TMP3]], i64 1
+// CHECK_HOSTCALL-NEXT: [[TMP5:%.*]] = load i8, ptr [[TMP3]], align 1
+// CHECK_HOSTCALL-NEXT: [[TMP6:%.*]] = icmp eq i8 [[TMP5]], 0
+// CHECK_HOSTCALL-NEXT: br i1 [[TMP6]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]]
+// CHECK_HOSTCALL: strlen.while.done:
+// CHECK_HOSTCALL-NEXT: [[TMP7:%.*]] = ptrtoint ptr [[TMP3]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP8:%.*]] = sub i64 [[TMP7]], ptrtoint (ptr addrspacecast (ptr addrspace(4) @.str.2 to ptr) to i64)
+// CHECK_HOSTCALL-NEXT: [[TMP9:%.*]] = add i64 [[TMP8]], 1
+// CHECK_HOSTCALL-NEXT: br label [[STRLEN_JOIN]]
+// CHECK_HOSTCALL: strlen.join:
+// CHECK_HOSTCALL-NEXT: [[TMP10:%.*]] = phi i64 [ [[TMP9]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]
+// CHECK_HOSTCALL-NEXT: [[TMP11:%.*]] = call i64 @__ockl_printf_append_string_n(i64 [[TMP2]], ptr addrspacecast (ptr addrspace(4) @.str.2 to ptr), i64 [[TMP10]], i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP12:%.*]] = icmp eq ptr [[TMP1]], null
+// CHECK_HOSTCALL-NEXT: br i1 [[TMP12]], label [[STRLEN_JOIN1:%.*]], label [[STRLEN_WHILE2:%.*]]
+// CHECK_HOSTCALL: strlen.while2:
+// CHECK_HOSTCALL-NEXT: [[TMP13:%.*]] = phi ptr [ [[TMP1]], [[STRLEN_JOIN]] ], [ [[TMP14:%.*]], [[STRLEN_WHILE2]] ]
+// CHECK_HOSTCALL-NEXT: [[TMP14]] = getelementptr i8, ptr [[TMP13]], i64 1
+// CHECK_HOSTCALL-NEXT: [[TMP15:%.*]] = load i8, ptr [[TMP13]], align 1
+// CHECK_HOSTCALL-NEXT: [[TMP16:%.*]] = icmp eq i8 [[TMP15]], 0
+// CHECK_HOSTCALL-NEXT: br i1 [[TMP16]], label [[STRLEN_WHILE_DONE3:%.*]], label [[STRLEN_WHILE2]]
+// CHECK_HOSTCALL: strlen.while.done3:
+// CHECK_HOSTCALL-NEXT: [[TMP17:%.*]] = ptrtoint ptr [[TMP1]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP18:%.*]] = ptrtoint ptr [[TMP13]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP19:%.*]] = sub i64 [[TMP18]], [[TMP17]]
+// CHECK_HOSTCALL-NEXT: [[TMP20:%.*]] = add i64 [[TMP19]], 1
+// CHECK_HOSTCALL-NEXT: br label [[STRLEN_JOIN1]]
+// CHECK_HOSTCALL: strlen.join1:
+// CHECK_HOSTCALL-NEXT: [[TMP21:%.*]] = phi i64 [ [[TMP20]], [[STRLEN_WHILE_DONE3]] ], [ 0, [[STRLEN_JOIN]] ]
+// CHECK_HOSTCALL-NEXT: [[TMP22:%.*]] = call i64 @__ockl_printf_append_string_n(i64 [[TMP11]], ptr [[TMP1]], i64 [[TMP21]], i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP23:%.*]] = zext i32 [[TMP0]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP24:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP22]], i32 1, i64 [[TMP23]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 1)
+// CHECK_HOSTCALL-NEXT: [[TMP25:%.*]] = trunc i64 [[TMP24]] to i32
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[S]]) #[[ATTR3]]
+// CHECK_HOSTCALL-NEXT: ret void
//
__kernel void test_printf_str_int(int i) {
char s[] = "foo";
diff --git a/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp b/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp
index 6ca737df49b95..bc3f20443504f 100644
--- a/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp
+++ b/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp
@@ -26,28 +26,31 @@ using namespace llvm;
#define DEBUG_TYPE "amdgpu-emit-printf"
-static Value *fitArgInto64Bits(IRBuilder<> &Builder, Value *Arg) {
+static Value *fitArgInto64Bits(IRBuilder<> &Builder, Value *Arg,
+ bool IsBuffered) {
+ const DataLayout &DL = Builder.GetInsertBlock()->getModule()->getDataLayout();
auto Int64Ty = Builder.getInt64Ty();
auto Ty = Arg->getType();
if (auto IntTy = dyn_cast<IntegerType>(Ty)) {
- switch (IntTy->getBitWidth()) {
- case 32:
- return Builder.CreateZExt(Arg, Int64Ty);
- case 64:
- return Arg;
+ if (IntTy->getBitWidth() < 64) {
+ return Builder.CreateZExt(Arg, Builder.getInt64Ty());
}
}
- if (Ty->getTypeID() == Type::DoubleTyID) {
+ if (Ty->isFloatingPointTy()) {
+ if (DL.getTypeAllocSize(Ty) < 8)
+ Arg = Builder.CreateFPExt(Arg, Builder.getDoubleTy());
+ if (IsBuffered)
+ return Arg;
return Builder.CreateBitCast(Arg, Int64Ty);
}
- if (isa<PointerType>(Ty)) {
+ if (!IsBuffered && isa<PointerType>(Ty)) {
return Builder.CreatePtrToInt(Arg, Int64Ty);
}
- llvm_unreachable("unexpected type");
+ return Arg;
}
static Value *callPrintfBegin(IRBuilder<> &Builder, Value *Version) {
@@ -74,8 +77,8 @@ static Value *callAppendArgs(IRBuilder<> &Builder, Value *Desc, int NumArgs,
}
static Value *appendArg(IRBuilder<> &Builder, Value *Desc, Value *Arg,
- bool IsLast) {
- auto Arg0 = fitArgInto64Bits(Builder, Arg);
+ bool IsLast, bool IsBuffered) {
+ auto Arg0 = fitArgInto64Bits(Builder, Arg, IsBuffered);
auto Zero = Builder.getInt64(0);
return callAppendArgs(Builder, Desc, 1, Arg0, Zero, Zero, Zero, Zero, Zero,
Zero, IsLast);
@@ -168,20 +171,49 @@ static Value *appendString(IRBuilder<> &Builder, Value *Desc, Value *Arg,
return callAppendStringN(Builder, Desc, Arg, Length, IsLast);
}
+static Value *appendVectorArg(IRBuilder<> &Builder, Value *Desc, Value *Arg,
+ bool IsLast, bool IsBuffered) {
+ assert(Arg->getType()->isVectorTy() && "incorrent append* function");
+ auto VectorTy = dyn_cast<FixedVectorType>(Arg->getType());
+ auto Zero = Builder.getInt64(0);
+ if (VectorTy) {
+ for (unsigned int i = 0; i < VectorTy->getNumElements() - 1; i++) {
+ auto Val = Builder.CreateExtractElement(Arg, i);
+ Desc = callAppendArgs(Builder, Desc, 1,
+ fitArgInto64Bits(Builder, Val, IsBuffered), Zero,
+ Zero, Zero, Zero, Zero, Zero, false);
+ }
+
+ Value* Val =
+ Builder.CreateExtractElement(Arg, VectorTy->getNumElements() - 1);
+ return callAppendArgs(Builder, Desc, 1,
+ fitArgInto64Bits(Builder, Val, IsBuffered), Zero,
+ Zero, Zero, Zero, Zero, Zero, IsLast);
+ }
+ return nullptr;
+}
+
static Value *processArg(IRBuilder<> &Builder, Value *Desc, Value *Arg,
- bool SpecIsCString, bool IsLast) {
+ bool SpecIsCString, bool IsVector, bool IsLast,
+ bool IsBuffered) {
if (SpecIsCString && isa<PointerType>(Arg->getType())) {
return appendString(Builder, Desc, Arg, IsLast);
}
- // 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);
+
+ if (IsVector) {
+ return appendVectorArg(Builder, Desc, Arg, IsLast, IsBuffered);
+ }
+
+ // 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, IsBuffered);
}
// 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, StringRef Str) {
+static void locateCStringsAndVectors(SparseBitVector<8> &BV,
+ SparseBitVector<8> &OV, StringRef Str) {
static const char ConvSpecifiers[] = "diouxXfFeEgGaAcspn";
size_t SpecPos = 0;
// Skip the first argument, the format string.
@@ -192,6 +224,8 @@ static void locateCStrings(SparseBitVector<8> &BV, StringRef Str) {
SpecPos += 2;
continue;
}
+ if (Str.find_first_of("v", SpecPos) != StringRef::npos)
+ OV.set(ArgIdx);
auto SpecEnd = Str.find_first_of(ConvSpecifiers, SpecPos);
if (SpecEnd == StringRef::npos)
return;
@@ -222,7 +256,8 @@ struct StringData {
static Value *callBufferedPrintfStart(
IRBuilder<> &Builder, ArrayRef<Value *> Args, Value *Fmt,
bool isConstFmtStr, SparseBitVector<8> &SpecIsCString,
- SmallVectorImpl<StringData> &StringContents, Value *&ArgSize) {
+ SparseBitVector<8> &OCLVectors, SmallVectorImpl<StringData> &StringContents,
+ Value *&ArgSize) {
Module *M = Builder.GetInsertBlock()->getModule();
Value *NonConstStrLen = nullptr;
Value *LenWithNull = nullptr;
@@ -276,7 +311,12 @@ static Value *callBufferedPrintfStart(
StringData(StringRef(), LenWithNull, LenWithNullAligned, false));
}
} else {
- int AllocSize = M->getDataLayout().getTypeAllocSize(Args[i]->getType());
+ int AllocSize = 0;
+ if (OCLVectors.test(i)) {
+ auto VecArg = cast<FixedVectorType>(Args[i]->getType());
+ AllocSize = VecArg->getNumElements() * 8;
+ } else
+ AllocSize = M->getDataLayout().getTypeAllocSize(Args[i]->getType());
// We end up expanding non string arguments to 8 bytes
// (args smaller than 8 bytes)
BufSize += std::max(AllocSize, 8);
@@ -350,30 +390,10 @@ static void processConstantStringArg(StringData *SD, IRBuilder<> &Builder,
WhatToStore.push_back(ConstantInt::get(Builder.getInt32Ty(), 0));
}
-static Value *processNonStringArg(Value *Arg, IRBuilder<> &Builder) {
- const DataLayout &DL = Builder.GetInsertBlock()->getModule()->getDataLayout();
- auto Ty = Arg->getType();
-
- if (auto IntTy = dyn_cast<IntegerType>(Ty)) {
- if (IntTy->getBitWidth() < 64) {
- return Builder.CreateZExt(Arg, Builder.getInt64Ty());
- }
- }
-
- if (Ty->isFloatingPointTy()) {
- if (DL.getTypeAllocSize(Ty) < 8) {
- return Builder.CreateFPExt(Arg, Builder.getDoubleTy());
- }
- }
-
- return Arg;
-}
-
-static void
-callBufferedPrintfArgPush(IRBuilder<> &Builder, ArrayRef<Value *> Args,
- Value *PtrToStore, SparseBitVector<8> &SpecIsCString,
- SmallVectorImpl<StringData> &StringContents,
- bool IsConstFmtStr) {
+static void callBufferedPrintfArgPush(
+ IRBuilder<> &Builder, ArrayRef<Value *> Args, Value *PtrToStore,
+ SparseBitVector<8> &SpecIsCString, SparseBitVector<8> &OCLVectors,
+ SmallVectorImpl<StringData> &StringContents, bool IsConstFmtStr) {
Module *M = Builder.GetInsertBlock()->getModule();
const DataLayout &DL = M->getDataLayout();
auto StrIt = StringContents.begin();
@@ -405,7 +425,17 @@ callBufferedPrintfArgPush(IRBuilder<> &Builder, ArrayRef<Value *> Args,
continue;
}
} else {
- WhatToStore.push_back(processNonStringArg(Args[i], Builder));
+ if (OCLVectors.test(i)) {
+ auto VectorTy = dyn_cast<FixedVectorType>(Args[i]->getType());
+ auto VecArg = Args[i];
+ for (unsigned int Num = 0; Num < VectorTy->getNumElements(); Num++) {
+ auto Val = Builder.CreateExtractElement(VecArg, Num);
+ WhatToStore.push_back(
+ fitArgInto64Bits(Builder, Val, /*IsBuffered*/ true));
+ }
+ } else
+ WhatToStore.push_back(
+ fitArgInto64Bits(Builder, Args[i], /*IsBuffered*/ true));
}
for (unsigned I = 0, E = WhatToStore.size(); I != E; ++I) {
@@ -432,10 +462,11 @@ Value *llvm::emitAMDGPUPrintfCall(IRBuilder<> &Builder, ArrayRef<Value *> Args,
auto Fmt = Args[0];
SparseBitVector<8> SpecIsCString;
+ SparseBitVector<8> OCLVectors;
StringRef FmtStr;
if (getConstantStringInfo(Fmt, FmtStr))
- locateCStrings(SpecIsCString, FmtStr);
+ locateCStringsAndVectors(SpecIsCString, OCLVectors, FmtStr);
if (IsBuffered) {
SmallVector<StringData, 8> StringContents;
@@ -446,9 +477,9 @@ Value *llvm::emitAMDGPUPrintfCall(IRBuilder<> &Builder, ArrayRef<Value *> Args,
bool IsConstFmtStr = !FmtStr.empty();
Value *ArgSize = nullptr;
- Value *Ptr =
- callBufferedPrintfStart(Builder, Args, Fmt, IsConstFmtStr,
- SpecIsCString, StringContents, ArgSize);
+ Value *Ptr = callBufferedPrintfStart(Builder, Args, Fmt, IsConstFmtStr,
+ SpecIsCString, OCLVectors,
+ StringContents, ArgSize);
// The buffered version still follows OpenCL printf standards for
// printf return value, i.e 0 on success, -1 on failure.
@@ -511,8 +542,8 @@ Value *llvm::emitAMDGPUPrintfCall(IRBuilder<> &Builder, ArrayRef<Value *> Args,
}
// Push The printf arguments onto buffer
- callBufferedPrintfArgPush(Builder, Args, Ptr, SpecIsCString, StringContents,
- IsConstFmtStr);
+ callBufferedPrintfArgPush(Builder, Args, Ptr, SpecIsCString, OCLVectors,
+ StringContents, IsConstFmtStr);
// End block, returns -1 on failure
BranchInst::Create(End, ArgPush);
@@ -529,7 +560,9 @@ Value *llvm::emitAMDGPUPrintfCall(IRBuilder<> &Builder, ArrayRef<Value *> Args,
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);
+ bool IsVector = OCLVectors.test(i);
+ Desc = processArg(Builder, Desc, Args[i], IsCString, IsVector, IsLast,
+ IsBuffered);
}
return Builder.CreateTrunc(Desc, Builder.getInt32Ty());
>From 73e271569174677fd6df1143f040d3bd22b9b025 Mon Sep 17 00:00:00 2001
From: Vikram <Vikram.Hegde at amd.com>
Date: Fri, 24 Nov 2023 18:02:36 +0000
Subject: [PATCH 4/5] Review commnents handled, code refactoring
---
clang/lib/CodeGen/CGGPUBuiltin.cpp | 20 +-
clang/test/CodeGenOpenCL/amdgpu-printf.cl | 616 +++++++++++++++++-
.../llvm/Transforms/Utils/AMDGPUEmitPrintf.h | 2 +-
.../lib/Transforms/Utils/AMDGPUEmitPrintf.cpp | 50 +-
4 files changed, 629 insertions(+), 59 deletions(-)
diff --git a/clang/lib/CodeGen/CGGPUBuiltin.cpp b/clang/lib/CodeGen/CGGPUBuiltin.cpp
index 32b457812af60..21bc39772d009 100644
--- a/clang/lib/CodeGen/CGGPUBuiltin.cpp
+++ b/clang/lib/CodeGen/CGGPUBuiltin.cpp
@@ -14,7 +14,9 @@
#include "CodeGenFunction.h"
#include "clang/Basic/Builtins.h"
#include "clang/Basic/TargetBuiltins.h"
+#include "llvm/Analysis/ValueTracking.h"
#include "llvm/IR/DataLayout.h"
+#include "llvm/IR/DiagnosticInfo.h"
#include "llvm/IR/Instruction.h"
#include "llvm/Support/MathExtras.h"
#include "llvm/Transforms/Utils/AMDGPUEmitPrintf.h"
@@ -219,7 +221,23 @@ RValue CodeGenFunction::EmitAMDGPUDevicePrintfCallExpr(const CallExpr *E) {
auto PFK = CGM.getTarget().getTargetOpts().AMDGPUPrintfKindVal;
bool isBuffered =
(PFK == clang::TargetOptions::AMDGPUPrintfKind::Buffered);
- auto Printf = llvm::emitAMDGPUPrintfCall(IRB, Args, isBuffered);
+
+ StringRef FmtStr;
+ if (llvm::getConstantStringInfo(Args[0], FmtStr)) {
+ if (FmtStr.empty())
+ FmtStr = StringRef("", 1);
+ } else {
+ if (CGM.getLangOpts().OpenCL) {
+ llvm::DiagnosticInfoUnsupported UnsupportedFormatStr(
+ *IRB.GetInsertBlock()->getParent(),
+ "printf format string must be a trivially resolved constant string "
+ "global variable",
+ IRB.getCurrentDebugLocation());
+ IRB.getContext().diagnose(UnsupportedFormatStr);
+ }
+ }
+
+ auto Printf = llvm::emitAMDGPUPrintfCall(IRB, Args, FmtStr, isBuffered);
Builder.SetInsertPoint(IRB.GetInsertBlock(), IRB.GetInsertPoint());
return RValue::get(Printf);
}
diff --git a/clang/test/CodeGenOpenCL/amdgpu-printf.cl b/clang/test/CodeGenOpenCL/amdgpu-printf.cl
index 5ada787f723ec..3342749e24722 100644
--- a/clang/test/CodeGenOpenCL/amdgpu-printf.cl
+++ b/clang/test/CodeGenOpenCL/amdgpu-printf.cl
@@ -2,6 +2,8 @@
// RUN: %clang_cc1 -cl-std=CL1.2 -triple amdgcn-amd-amdhsa -mprintf-kind=buffered -disable-llvm-passes -emit-llvm -o - %s | FileCheck --check-prefix=CHECK_BUFFERED %s
// RUN: %clang_cc1 -cl-std=CL1.2 -triple amdgcn-amd-amdhsa -mprintf-kind=hostcall -disable-llvm-passes -emit-llvm -o - %s | FileCheck --check-prefix=CHECK_HOSTCALL %s
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
int printf(__constant const char* st, ...) __attribute__((format(printf, 1, 2)));
// CHECK-LABEL: @test_printf_noargs(
@@ -10,37 +12,18 @@ int printf(__constant const char* st, ...) __attribute__((format(printf, 1, 2)))
// CHECK-NEXT: ret void
// CHECK_BUFFERED-LABEL: @test_printf_noargs(
// CHECK_BUFFERED-NEXT: entry:
-// CHECK_BUFFERED-NEXT: br i1 icmp eq (ptr addrspacecast (ptr addrspace(4) @.str to ptr), ptr null), label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]
-// CHECK_BUFFERED: strlen.while:
-// CHECK_BUFFERED-NEXT: [[TMP0:%.*]] = phi ptr [ addrspacecast (ptr addrspace(4) @.str to ptr), [[ENTRY:%.*]] ], [ [[TMP1:%.*]], [[STRLEN_WHILE]] ]
-// CHECK_BUFFERED-NEXT: [[TMP1]] = getelementptr i8, ptr [[TMP0]], i64 1
-// CHECK_BUFFERED-NEXT: [[TMP2:%.*]] = load i8, ptr [[TMP0]], align 1
-// CHECK_BUFFERED-NEXT: [[TMP3:%.*]] = icmp eq i8 [[TMP2]], 0
-// CHECK_BUFFERED-NEXT: br i1 [[TMP3]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]]
-// CHECK_BUFFERED: strlen.while.done:
-// CHECK_BUFFERED-NEXT: [[TMP4:%.*]] = ptrtoint ptr [[TMP0]] to i64
-// CHECK_BUFFERED-NEXT: [[TMP5:%.*]] = sub i64 [[TMP4]], ptrtoint (ptr addrspacecast (ptr addrspace(4) @.str to ptr) to i64)
-// CHECK_BUFFERED-NEXT: [[TMP6:%.*]] = add i64 [[TMP5]], 1
-// CHECK_BUFFERED-NEXT: br label [[STRLEN_JOIN]]
-// CHECK_BUFFERED: strlen.join:
-// CHECK_BUFFERED-NEXT: [[TMP7:%.*]] = phi i64 [ [[TMP6]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]
-// CHECK_BUFFERED-NEXT: [[TMP8:%.*]] = add i64 [[TMP7]], 7
-// CHECK_BUFFERED-NEXT: [[TMP9:%.*]] = and i64 [[TMP8]], 4294967288
-// CHECK_BUFFERED-NEXT: [[TMP10:%.*]] = add i64 [[TMP9]], 4
-// CHECK_BUFFERED-NEXT: [[TMP11:%.*]] = trunc i64 [[TMP10]] to i32
-// CHECK_BUFFERED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[TMP11]])
-// CHECK_BUFFERED-NEXT: [[TMP12:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
-// CHECK_BUFFERED-NEXT: br i1 [[TMP12]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]
+// CHECK_BUFFERED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 12)
+// CHECK_BUFFERED-NEXT: [[TMP0:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
+// CHECK_BUFFERED-NEXT: br i1 [[TMP0]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]
// CHECK_BUFFERED: end.block:
-// CHECK_BUFFERED-NEXT: [[TMP13:%.*]] = xor i1 [[TMP12]], true
-// CHECK_BUFFERED-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP13]] to i32
+// CHECK_BUFFERED-NEXT: [[TMP1:%.*]] = xor i1 [[TMP0]], true
+// CHECK_BUFFERED-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP1]] to i32
// CHECK_BUFFERED-NEXT: ret void
// CHECK_BUFFERED: argpush.block:
-// CHECK_BUFFERED-NEXT: [[TMP14:%.*]] = shl i32 [[TMP11]], 2
-// CHECK_BUFFERED-NEXT: store i32 [[TMP14]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
-// CHECK_BUFFERED-NEXT: [[TMP15:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
-// CHECK_BUFFERED-NEXT: call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 1 [[TMP15]], ptr align 1 addrspacecast (ptr addrspace(4) @.str to ptr), i64 [[TMP7]], i1 false)
-// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP15]], i64 [[TMP9]]
+// CHECK_BUFFERED-NEXT: store i32 50, ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
+// CHECK_BUFFERED-NEXT: [[TMP2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
+// CHECK_BUFFERED-NEXT: store i64 -8529802306755643245, ptr addrspace(1) [[TMP2]], align 8
+// CHECK_BUFFERED-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP2]], i32 8
// CHECK_BUFFERED-NEXT: br label [[END_BLOCK]]
//
// CHECK_HOSTCALL-LABEL: @test_printf_noargs(
@@ -78,8 +61,8 @@ __kernel void test_printf_noargs() {
// CHECK_BUFFERED-LABEL: @test_printf_int(
// CHECK_BUFFERED-NEXT: entry:
// CHECK_BUFFERED-NEXT: [[I_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
-// CHECK_BUFFERED-NEXT: store i32 [[I:%.*]], ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA12:![0-9]+]]
-// CHECK_BUFFERED-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA12]]
+// CHECK_BUFFERED-NEXT: store i32 [[I:%.*]], ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA16:![0-9]+]]
+// CHECK_BUFFERED-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA16]]
// CHECK_BUFFERED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 20)
// CHECK_BUFFERED-NEXT: [[TMP1:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
// CHECK_BUFFERED-NEXT: br i1 [[TMP1]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]
@@ -150,11 +133,11 @@ __kernel void test_printf_int(int i) {
// CHECK_BUFFERED-NEXT: entry:
// CHECK_BUFFERED-NEXT: [[I_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK_BUFFERED-NEXT: [[S:%.*]] = alloca [4 x i8], align 1, addrspace(5)
-// CHECK_BUFFERED-NEXT: store i32 [[I:%.*]], ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA12]]
+// CHECK_BUFFERED-NEXT: store i32 [[I:%.*]], ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA16]]
// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[S]]) #[[ATTR1:[0-9]+]]
// CHECK_BUFFERED-NEXT: call void @llvm.memcpy.p5.p4.i64(ptr addrspace(5) align 1 [[S]], ptr addrspace(4) align 1 @__const.test_printf_str_int.s, i64 4, i1 false)
// CHECK_BUFFERED-NEXT: [[ARRAYDECAY:%.*]] = getelementptr inbounds [4 x i8], ptr addrspace(5) [[S]], i64 0, i64 0
-// CHECK_BUFFERED-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA12]]
+// CHECK_BUFFERED-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA16]]
// CHECK_BUFFERED-NEXT: [[TMP1:%.*]] = addrspacecast ptr addrspace(5) [[ARRAYDECAY]] to ptr
// CHECK_BUFFERED-NEXT: [[TMP2:%.*]] = icmp eq ptr [[TMP1]], null
// CHECK_BUFFERED-NEXT: br i1 [[TMP2]], label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]
@@ -251,3 +234,572 @@ __kernel void test_printf_str_int(int i) {
char s[] = "foo";
printf("%s:%d", s, i);
}
+
+// CHECK_BUFFERED-LABEL: @test_half(
+// CHECK_BUFFERED-NEXT: entry:
+// CHECK_BUFFERED-NEXT: [[VAR:%.*]] = alloca half, align 2, addrspace(5)
+// CHECK_BUFFERED-NEXT: [[VAR2:%.*]] = alloca half, align 2, addrspace(5)
+// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.start.p5(i64 2, ptr addrspace(5) [[VAR]]) #[[ATTR1]]
+// CHECK_BUFFERED-NEXT: store half 0xH2E66, ptr addrspace(5) [[VAR]], align 2, !tbaa [[TBAA20:![0-9]+]]
+// CHECK_BUFFERED-NEXT: [[TMP0:%.*]] = load half, ptr addrspace(5) [[VAR]], align 2, !tbaa [[TBAA20]]
+// CHECK_BUFFERED-NEXT: [[CONV:%.*]] = fpext half [[TMP0]] to double
+// CHECK_BUFFERED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 20)
+// CHECK_BUFFERED-NEXT: [[TMP1:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
+// CHECK_BUFFERED-NEXT: br i1 [[TMP1]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]
+// CHECK_BUFFERED: end.block:
+// CHECK_BUFFERED-NEXT: [[TMP2:%.*]] = xor i1 [[TMP1]], true
+// CHECK_BUFFERED-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP2]] to i32
+// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.start.p5(i64 2, ptr addrspace(5) [[VAR2]]) #[[ATTR1]]
+// CHECK_BUFFERED-NEXT: store half 0xH3266, ptr addrspace(5) [[VAR2]], align 2, !tbaa [[TBAA22:![0-9]+]]
+// CHECK_BUFFERED-NEXT: [[TMP3:%.*]] = load half, ptr addrspace(5) [[VAR2]], align 2, !tbaa [[TBAA22]]
+// CHECK_BUFFERED-NEXT: [[PRINTF_ALLOC_FN1:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 20)
+// CHECK_BUFFERED-NEXT: [[TMP4:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN1]], null
+// CHECK_BUFFERED-NEXT: br i1 [[TMP4]], label [[ARGPUSH_BLOCK3:%.*]], label [[END_BLOCK2:%.*]]
+// CHECK_BUFFERED: argpush.block:
+// CHECK_BUFFERED-NEXT: store i32 82, ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
+// CHECK_BUFFERED-NEXT: [[TMP5:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
+// CHECK_BUFFERED-NEXT: store i64 5482945341418147321, ptr addrspace(1) [[TMP5]], align 8
+// CHECK_BUFFERED-NEXT: [[TMP6:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP5]], i32 8
+// CHECK_BUFFERED-NEXT: store double [[CONV]], ptr addrspace(1) [[TMP6]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP6]], i32 8
+// CHECK_BUFFERED-NEXT: br label [[END_BLOCK]]
+// CHECK_BUFFERED: end.block2:
+// CHECK_BUFFERED-NEXT: [[TMP7:%.*]] = xor i1 [[TMP4]], true
+// CHECK_BUFFERED-NEXT: [[PRINTF_RESULT5:%.*]] = sext i1 [[TMP7]] to i32
+// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.end.p5(i64 2, ptr addrspace(5) [[VAR2]]) #[[ATTR1]]
+// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.end.p5(i64 2, ptr addrspace(5) [[VAR]]) #[[ATTR1]]
+// CHECK_BUFFERED-NEXT: ret void
+// CHECK_BUFFERED: argpush.block3:
+// CHECK_BUFFERED-NEXT: store i32 82, ptr addrspace(1) [[PRINTF_ALLOC_FN1]], align 4
+// CHECK_BUFFERED-NEXT: [[TMP8:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN1]], i32 4
+// CHECK_BUFFERED-NEXT: store i64 5482945341418147321, ptr addrspace(1) [[TMP8]], align 8
+// CHECK_BUFFERED-NEXT: [[TMP9:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP8]], i32 8
+// CHECK_BUFFERED-NEXT: [[TMP10:%.*]] = fpext half [[TMP3]] to double
+// CHECK_BUFFERED-NEXT: store double [[TMP10]], ptr addrspace(1) [[TMP9]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR4:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP9]], i32 8
+// CHECK_BUFFERED-NEXT: br label [[END_BLOCK2]]
+//
+// CHECK_HOSTCALL-LABEL: @test_half(
+// CHECK_HOSTCALL-NEXT: entry:
+// CHECK_HOSTCALL-NEXT: [[VAR:%.*]] = alloca half, align 2, addrspace(5)
+// CHECK_HOSTCALL-NEXT: [[VAR2:%.*]] = alloca half, align 2, addrspace(5)
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.start.p5(i64 2, ptr addrspace(5) [[VAR]]) #[[ATTR3]]
+// CHECK_HOSTCALL-NEXT: store half 0xH2E66, ptr addrspace(5) [[VAR]], align 2, !tbaa [[TBAA13:![0-9]+]]
+// CHECK_HOSTCALL-NEXT: [[TMP0:%.*]] = load half, ptr addrspace(5) [[VAR]], align 2, !tbaa [[TBAA13]]
+// CHECK_HOSTCALL-NEXT: [[CONV:%.*]] = fpext half [[TMP0]] to double
+// CHECK_HOSTCALL-NEXT: [[TMP1:%.*]] = call i64 @__ockl_printf_begin(i64 0)
+// CHECK_HOSTCALL-NEXT: br i1 icmp eq (ptr addrspacecast (ptr addrspace(4) @.str.3 to ptr), ptr null), label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]
+// CHECK_HOSTCALL: strlen.while:
+// CHECK_HOSTCALL-NEXT: [[TMP2:%.*]] = phi ptr [ addrspacecast (ptr addrspace(4) @.str.3 to ptr), [[ENTRY:%.*]] ], [ [[TMP3:%.*]], [[STRLEN_WHILE]] ]
+// CHECK_HOSTCALL-NEXT: [[TMP3]] = getelementptr i8, ptr [[TMP2]], i64 1
+// CHECK_HOSTCALL-NEXT: [[TMP4:%.*]] = load i8, ptr [[TMP2]], align 1
+// CHECK_HOSTCALL-NEXT: [[TMP5:%.*]] = icmp eq i8 [[TMP4]], 0
+// CHECK_HOSTCALL-NEXT: br i1 [[TMP5]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]]
+// CHECK_HOSTCALL: strlen.while.done:
+// CHECK_HOSTCALL-NEXT: [[TMP6:%.*]] = ptrtoint ptr [[TMP2]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP7:%.*]] = sub i64 [[TMP6]], ptrtoint (ptr addrspacecast (ptr addrspace(4) @.str.3 to ptr) to i64)
+// CHECK_HOSTCALL-NEXT: [[TMP8:%.*]] = add i64 [[TMP7]], 1
+// CHECK_HOSTCALL-NEXT: br label [[STRLEN_JOIN]]
+// CHECK_HOSTCALL: strlen.join:
+// CHECK_HOSTCALL-NEXT: [[TMP9:%.*]] = phi i64 [ [[TMP8]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]
+// CHECK_HOSTCALL-NEXT: [[TMP10:%.*]] = call i64 @__ockl_printf_append_string_n(i64 [[TMP1]], ptr addrspacecast (ptr addrspace(4) @.str.3 to ptr), i64 [[TMP9]], i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP11:%.*]] = bitcast double [[CONV]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP12:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP10]], i32 1, i64 [[TMP11]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 1)
+// CHECK_HOSTCALL-NEXT: [[TMP13:%.*]] = trunc i64 [[TMP12]] to i32
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.start.p5(i64 2, ptr addrspace(5) [[VAR2]]) #[[ATTR3]]
+// CHECK_HOSTCALL-NEXT: store half 0xH3266, ptr addrspace(5) [[VAR2]], align 2, !tbaa [[TBAA15:![0-9]+]]
+// CHECK_HOSTCALL-NEXT: [[TMP14:%.*]] = load half, ptr addrspace(5) [[VAR2]], align 2, !tbaa [[TBAA15]]
+// CHECK_HOSTCALL-NEXT: [[TMP15:%.*]] = call i64 @__ockl_printf_begin(i64 0)
+// CHECK_HOSTCALL-NEXT: br i1 icmp eq (ptr addrspacecast (ptr addrspace(4) @.str.3 to ptr), ptr null), label [[STRLEN_JOIN1:%.*]], label [[STRLEN_WHILE2:%.*]]
+// CHECK_HOSTCALL: strlen.while2:
+// CHECK_HOSTCALL-NEXT: [[TMP16:%.*]] = phi ptr [ addrspacecast (ptr addrspace(4) @.str.3 to ptr), [[STRLEN_JOIN]] ], [ [[TMP17:%.*]], [[STRLEN_WHILE2]] ]
+// CHECK_HOSTCALL-NEXT: [[TMP17]] = getelementptr i8, ptr [[TMP16]], i64 1
+// CHECK_HOSTCALL-NEXT: [[TMP18:%.*]] = load i8, ptr [[TMP16]], align 1
+// CHECK_HOSTCALL-NEXT: [[TMP19:%.*]] = icmp eq i8 [[TMP18]], 0
+// CHECK_HOSTCALL-NEXT: br i1 [[TMP19]], label [[STRLEN_WHILE_DONE3:%.*]], label [[STRLEN_WHILE2]]
+// CHECK_HOSTCALL: strlen.while.done3:
+// CHECK_HOSTCALL-NEXT: [[TMP20:%.*]] = ptrtoint ptr [[TMP16]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP21:%.*]] = sub i64 [[TMP20]], ptrtoint (ptr addrspacecast (ptr addrspace(4) @.str.3 to ptr) to i64)
+// CHECK_HOSTCALL-NEXT: [[TMP22:%.*]] = add i64 [[TMP21]], 1
+// CHECK_HOSTCALL-NEXT: br label [[STRLEN_JOIN1]]
+// CHECK_HOSTCALL: strlen.join1:
+// CHECK_HOSTCALL-NEXT: [[TMP23:%.*]] = phi i64 [ [[TMP22]], [[STRLEN_WHILE_DONE3]] ], [ 0, [[STRLEN_JOIN]] ]
+// CHECK_HOSTCALL-NEXT: [[TMP24:%.*]] = call i64 @__ockl_printf_append_string_n(i64 [[TMP15]], ptr addrspacecast (ptr addrspace(4) @.str.3 to ptr), i64 [[TMP23]], i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP25:%.*]] = fpext half [[TMP14]] to double
+// CHECK_HOSTCALL-NEXT: [[TMP26:%.*]] = bitcast double [[TMP25]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP27:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP24]], i32 1, i64 [[TMP26]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 1)
+// CHECK_HOSTCALL-NEXT: [[TMP28:%.*]] = trunc i64 [[TMP27]] to i32
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.end.p5(i64 2, ptr addrspace(5) [[VAR2]]) #[[ATTR3]]
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.end.p5(i64 2, ptr addrspace(5) [[VAR]]) #[[ATTR3]]
+// CHECK_HOSTCALL-NEXT: ret void
+//
+__kernel void test_half() {
+ half var = 0.1f;
+ printf("%f", var);
+ _Float16 var2 = 0.2f;
+ printf("%f", var2);
+}
+
+typedef __attribute__((ext_vector_type(3))) int int3;
+typedef __attribute__((ext_vector_type(4))) int int4;
+typedef __attribute__((ext_vector_type(2))) float float2;
+typedef __attribute__((ext_vector_type(8))) float float8;
+typedef __attribute__((ext_vector_type(16))) float float16;
+
+// CHECK_BUFFERED-LABEL: @test_vector(
+// CHECK_BUFFERED-NEXT: entry:
+// CHECK_BUFFERED-NEXT: [[A:%.*]] = alloca <4 x i32>, align 16, addrspace(5)
+// CHECK_BUFFERED-NEXT: [[DOTCOMPOUNDLITERAL:%.*]] = alloca <4 x i32>, align 16, addrspace(5)
+// CHECK_BUFFERED-NEXT: [[B:%.*]] = alloca <3 x i32>, align 16, addrspace(5)
+// CHECK_BUFFERED-NEXT: [[DOTCOMPOUNDLITERAL4:%.*]] = alloca <3 x i32>, align 16, addrspace(5)
+// CHECK_BUFFERED-NEXT: [[VAR3:%.*]] = alloca <2 x float>, align 8, addrspace(5)
+// CHECK_BUFFERED-NEXT: [[DOTCOMPOUNDLITERAL17:%.*]] = alloca <2 x float>, align 8, addrspace(5)
+// CHECK_BUFFERED-NEXT: [[VAR4:%.*]] = alloca <8 x float>, align 32, addrspace(5)
+// CHECK_BUFFERED-NEXT: [[VAR5:%.*]] = alloca <16 x float>, align 64, addrspace(5)
+// CHECK_BUFFERED-NEXT: [[DOTCOMPOUNDLITERAL18:%.*]] = alloca <16 x float>, align 64, addrspace(5)
+// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.start.p5(i64 16, ptr addrspace(5) [[A]]) #[[ATTR1]]
+// CHECK_BUFFERED-NEXT: store <4 x i32> <i32 1, i32 2, i32 3, i32 4>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL]], align 16, !tbaa [[TBAA24:![0-9]+]]
+// CHECK_BUFFERED-NEXT: [[TMP0:%.*]] = load <4 x i32>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL]], align 16, !tbaa [[TBAA24]]
+// CHECK_BUFFERED-NEXT: store <4 x i32> [[TMP0]], ptr addrspace(5) [[A]], align 16, !tbaa [[TBAA24]]
+// CHECK_BUFFERED-NEXT: [[TMP1:%.*]] = load <4 x i32>, ptr addrspace(5) [[A]], align 16, !tbaa [[TBAA24]]
+// CHECK_BUFFERED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 44)
+// CHECK_BUFFERED-NEXT: [[TMP2:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
+// CHECK_BUFFERED-NEXT: br i1 [[TMP2]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]
+// CHECK_BUFFERED: end.block:
+// CHECK_BUFFERED-NEXT: [[TMP3:%.*]] = xor i1 [[TMP2]], true
+// CHECK_BUFFERED-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP3]] to i32
+// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.start.p5(i64 16, ptr addrspace(5) [[B]]) #[[ATTR1]]
+// CHECK_BUFFERED-NEXT: [[TMP4:%.*]] = load <4 x i32>, ptr addrspace(5) [[A]], align 16, !tbaa [[TBAA24]]
+// CHECK_BUFFERED-NEXT: [[VECEXT:%.*]] = extractelement <4 x i32> [[TMP4]], i32 2
+// CHECK_BUFFERED-NEXT: [[VECINIT:%.*]] = insertelement <3 x i32> <i32 1, i32 undef, i32 undef>, i32 [[VECEXT]], i32 1
+// CHECK_BUFFERED-NEXT: [[VECINIT5:%.*]] = insertelement <3 x i32> [[VECINIT]], i32 3, i32 2
+// CHECK_BUFFERED-NEXT: [[EXTRACTVEC:%.*]] = shufflevector <3 x i32> [[VECINIT5]], <3 x i32> poison, <4 x i32> <i32 0, i32 1, i32 2, i32 poison>
+// CHECK_BUFFERED-NEXT: store <4 x i32> [[EXTRACTVEC]], ptr addrspace(5) [[DOTCOMPOUNDLITERAL4]], align 16, !tbaa [[TBAA24]]
+// CHECK_BUFFERED-NEXT: [[LOADVEC4:%.*]] = load <4 x i32>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL4]], align 16
+// CHECK_BUFFERED-NEXT: [[EXTRACTVEC6:%.*]] = shufflevector <4 x i32> [[LOADVEC4]], <4 x i32> poison, <3 x i32> <i32 0, i32 1, i32 2>
+// CHECK_BUFFERED-NEXT: [[EXTRACTVEC7:%.*]] = shufflevector <3 x i32> [[EXTRACTVEC6]], <3 x i32> poison, <4 x i32> <i32 0, i32 1, i32 2, i32 poison>
+// CHECK_BUFFERED-NEXT: store <4 x i32> [[EXTRACTVEC7]], ptr addrspace(5) [[B]], align 16, !tbaa [[TBAA24]]
+// CHECK_BUFFERED-NEXT: [[LOADVEC48:%.*]] = load <4 x i32>, ptr addrspace(5) [[B]], align 16
+// CHECK_BUFFERED-NEXT: [[EXTRACTVEC9:%.*]] = shufflevector <4 x i32> [[LOADVEC48]], <4 x i32> poison, <3 x i32> <i32 0, i32 1, i32 2>
+// CHECK_BUFFERED-NEXT: [[PRINTF_ALLOC_FN10:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 36)
+// CHECK_BUFFERED-NEXT: [[TMP5:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN10]], null
+// CHECK_BUFFERED-NEXT: br i1 [[TMP5]], label [[ARGPUSH_BLOCK12:%.*]], label [[END_BLOCK11:%.*]]
+// CHECK_BUFFERED: argpush.block:
+// CHECK_BUFFERED-NEXT: store i32 178, ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
+// CHECK_BUFFERED-NEXT: [[TMP6:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
+// CHECK_BUFFERED-NEXT: store i64 157641626367533276, ptr addrspace(1) [[TMP6]], align 8
+// CHECK_BUFFERED-NEXT: [[TMP7:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP6]], i32 8
+// CHECK_BUFFERED-NEXT: [[TMP8:%.*]] = extractelement <4 x i32> [[TMP1]], i64 0
+// CHECK_BUFFERED-NEXT: [[TMP9:%.*]] = zext i32 [[TMP8]] to i64
+// CHECK_BUFFERED-NEXT: [[TMP10:%.*]] = extractelement <4 x i32> [[TMP1]], i64 1
+// CHECK_BUFFERED-NEXT: [[TMP11:%.*]] = zext i32 [[TMP10]] to i64
+// CHECK_BUFFERED-NEXT: [[TMP12:%.*]] = extractelement <4 x i32> [[TMP1]], i64 2
+// CHECK_BUFFERED-NEXT: [[TMP13:%.*]] = zext i32 [[TMP12]] to i64
+// CHECK_BUFFERED-NEXT: [[TMP14:%.*]] = extractelement <4 x i32> [[TMP1]], i64 3
+// CHECK_BUFFERED-NEXT: [[TMP15:%.*]] = zext i32 [[TMP14]] to i64
+// CHECK_BUFFERED-NEXT: store i64 [[TMP9]], ptr addrspace(1) [[TMP7]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP7]], i32 8
+// CHECK_BUFFERED-NEXT: store i64 [[TMP11]], ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8
+// CHECK_BUFFERED-NEXT: store i64 [[TMP13]], ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR1]], i32 8
+// CHECK_BUFFERED-NEXT: store i64 [[TMP15]], ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR2]], i32 8
+// CHECK_BUFFERED-NEXT: br label [[END_BLOCK]]
+// CHECK_BUFFERED: end.block11:
+// CHECK_BUFFERED-NEXT: [[TMP16:%.*]] = xor i1 [[TMP5]], true
+// CHECK_BUFFERED-NEXT: [[PRINTF_RESULT16:%.*]] = sext i1 [[TMP16]] to i32
+// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[VAR3]]) #[[ATTR1]]
+// CHECK_BUFFERED-NEXT: store <2 x float> <float 0x4008F5C280000000, float 0x4003D70A40000000>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL17]], align 8, !tbaa [[TBAA24]]
+// CHECK_BUFFERED-NEXT: [[TMP17:%.*]] = load <2 x float>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL17]], align 8, !tbaa [[TBAA24]]
+// CHECK_BUFFERED-NEXT: store <2 x float> [[TMP17]], ptr addrspace(5) [[VAR3]], align 8, !tbaa [[TBAA24]]
+// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.start.p5(i64 32, ptr addrspace(5) [[VAR4]]) #[[ATTR1]]
+// CHECK_BUFFERED-NEXT: store <8 x float> <float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000>, ptr addrspace(5) [[VAR4]], align 32, !tbaa [[TBAA24]]
+// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.start.p5(i64 64, ptr addrspace(5) [[VAR5]]) #[[ATTR1]]
+// CHECK_BUFFERED-NEXT: store <16 x float> <float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL18]], align 64, !tbaa [[TBAA24]]
+// CHECK_BUFFERED-NEXT: [[TMP18:%.*]] = load <16 x float>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL18]], align 64, !tbaa [[TBAA24]]
+// CHECK_BUFFERED-NEXT: store <16 x float> [[TMP18]], ptr addrspace(5) [[VAR5]], align 64, !tbaa [[TBAA24]]
+// CHECK_BUFFERED-NEXT: [[TMP19:%.*]] = load <2 x float>, ptr addrspace(5) [[VAR3]], align 8, !tbaa [[TBAA24]]
+// CHECK_BUFFERED-NEXT: [[TMP20:%.*]] = load <8 x float>, ptr addrspace(5) [[VAR4]], align 32, !tbaa [[TBAA24]]
+// CHECK_BUFFERED-NEXT: [[TMP21:%.*]] = load <16 x float>, ptr addrspace(5) [[VAR5]], align 64, !tbaa [[TBAA24]]
+// CHECK_BUFFERED-NEXT: [[PRINTF_ALLOC_FN19:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 220)
+// CHECK_BUFFERED-NEXT: [[TMP22:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN19]], null
+// CHECK_BUFFERED-NEXT: br i1 [[TMP22]], label [[ARGPUSH_BLOCK21:%.*]], label [[END_BLOCK20:%.*]]
+// CHECK_BUFFERED: argpush.block12:
+// CHECK_BUFFERED-NEXT: store i32 146, ptr addrspace(1) [[PRINTF_ALLOC_FN10]], align 4
+// CHECK_BUFFERED-NEXT: [[TMP23:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN10]], i32 4
+// CHECK_BUFFERED-NEXT: store i64 8349377361424904666, ptr addrspace(1) [[TMP23]], align 8
+// CHECK_BUFFERED-NEXT: [[TMP24:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP23]], i32 8
+// CHECK_BUFFERED-NEXT: [[TMP25:%.*]] = extractelement <3 x i32> [[EXTRACTVEC9]], i64 0
+// CHECK_BUFFERED-NEXT: [[TMP26:%.*]] = zext i32 [[TMP25]] to i64
+// CHECK_BUFFERED-NEXT: [[TMP27:%.*]] = extractelement <3 x i32> [[EXTRACTVEC9]], i64 1
+// CHECK_BUFFERED-NEXT: [[TMP28:%.*]] = zext i32 [[TMP27]] to i64
+// CHECK_BUFFERED-NEXT: [[TMP29:%.*]] = extractelement <3 x i32> [[EXTRACTVEC9]], i64 2
+// CHECK_BUFFERED-NEXT: [[TMP30:%.*]] = zext i32 [[TMP29]] to i64
+// CHECK_BUFFERED-NEXT: store i64 [[TMP26]], ptr addrspace(1) [[TMP24]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR13:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP24]], i32 8
+// CHECK_BUFFERED-NEXT: store i64 [[TMP28]], ptr addrspace(1) [[PRINTBUFFNEXTPTR13]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR14:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR13]], i32 8
+// CHECK_BUFFERED-NEXT: store i64 [[TMP30]], ptr addrspace(1) [[PRINTBUFFNEXTPTR14]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR15:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR14]], i32 8
+// CHECK_BUFFERED-NEXT: br label [[END_BLOCK11]]
+// CHECK_BUFFERED: end.block20:
+// CHECK_BUFFERED-NEXT: [[TMP31:%.*]] = xor i1 [[TMP22]], true
+// CHECK_BUFFERED-NEXT: [[PRINTF_RESULT48:%.*]] = sext i1 [[TMP31]] to i32
+// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.end.p5(i64 64, ptr addrspace(5) [[VAR5]]) #[[ATTR1]]
+// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.end.p5(i64 32, ptr addrspace(5) [[VAR4]]) #[[ATTR1]]
+// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[VAR3]]) #[[ATTR1]]
+// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.end.p5(i64 16, ptr addrspace(5) [[B]]) #[[ATTR1]]
+// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.end.p5(i64 16, ptr addrspace(5) [[A]]) #[[ATTR1]]
+// CHECK_BUFFERED-NEXT: ret void
+// CHECK_BUFFERED: argpush.block21:
+// CHECK_BUFFERED-NEXT: store i32 882, ptr addrspace(1) [[PRINTF_ALLOC_FN19]], align 4
+// CHECK_BUFFERED-NEXT: [[TMP32:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN19]], i32 4
+// CHECK_BUFFERED-NEXT: store i64 -7312483467425511358, ptr addrspace(1) [[TMP32]], align 8
+// CHECK_BUFFERED-NEXT: [[TMP33:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP32]], i32 8
+// CHECK_BUFFERED-NEXT: [[TMP34:%.*]] = extractelement <2 x float> [[TMP19]], i64 0
+// CHECK_BUFFERED-NEXT: [[TMP35:%.*]] = fpext float [[TMP34]] to double
+// CHECK_BUFFERED-NEXT: [[TMP36:%.*]] = extractelement <2 x float> [[TMP19]], i64 1
+// CHECK_BUFFERED-NEXT: [[TMP37:%.*]] = fpext float [[TMP36]] to double
+// CHECK_BUFFERED-NEXT: store double [[TMP35]], ptr addrspace(1) [[TMP33]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR22:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP33]], i32 8
+// CHECK_BUFFERED-NEXT: store double [[TMP37]], ptr addrspace(1) [[PRINTBUFFNEXTPTR22]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR23:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR22]], i32 8
+// CHECK_BUFFERED-NEXT: [[TMP38:%.*]] = extractelement <8 x float> [[TMP20]], i64 0
+// CHECK_BUFFERED-NEXT: [[TMP39:%.*]] = fpext float [[TMP38]] to double
+// CHECK_BUFFERED-NEXT: [[TMP40:%.*]] = extractelement <8 x float> [[TMP20]], i64 1
+// CHECK_BUFFERED-NEXT: [[TMP41:%.*]] = fpext float [[TMP40]] to double
+// CHECK_BUFFERED-NEXT: [[TMP42:%.*]] = extractelement <8 x float> [[TMP20]], i64 2
+// CHECK_BUFFERED-NEXT: [[TMP43:%.*]] = fpext float [[TMP42]] to double
+// CHECK_BUFFERED-NEXT: [[TMP44:%.*]] = extractelement <8 x float> [[TMP20]], i64 3
+// CHECK_BUFFERED-NEXT: [[TMP45:%.*]] = fpext float [[TMP44]] to double
+// CHECK_BUFFERED-NEXT: [[TMP46:%.*]] = extractelement <8 x float> [[TMP20]], i64 4
+// CHECK_BUFFERED-NEXT: [[TMP47:%.*]] = fpext float [[TMP46]] to double
+// CHECK_BUFFERED-NEXT: [[TMP48:%.*]] = extractelement <8 x float> [[TMP20]], i64 5
+// CHECK_BUFFERED-NEXT: [[TMP49:%.*]] = fpext float [[TMP48]] to double
+// CHECK_BUFFERED-NEXT: [[TMP50:%.*]] = extractelement <8 x float> [[TMP20]], i64 6
+// CHECK_BUFFERED-NEXT: [[TMP51:%.*]] = fpext float [[TMP50]] to double
+// CHECK_BUFFERED-NEXT: [[TMP52:%.*]] = extractelement <8 x float> [[TMP20]], i64 7
+// CHECK_BUFFERED-NEXT: [[TMP53:%.*]] = fpext float [[TMP52]] to double
+// CHECK_BUFFERED-NEXT: store double [[TMP39]], ptr addrspace(1) [[PRINTBUFFNEXTPTR23]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR24:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR23]], i32 8
+// CHECK_BUFFERED-NEXT: store double [[TMP41]], ptr addrspace(1) [[PRINTBUFFNEXTPTR24]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR25:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR24]], i32 8
+// CHECK_BUFFERED-NEXT: store double [[TMP43]], ptr addrspace(1) [[PRINTBUFFNEXTPTR25]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR26:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR25]], i32 8
+// CHECK_BUFFERED-NEXT: store double [[TMP45]], ptr addrspace(1) [[PRINTBUFFNEXTPTR26]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR27:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR26]], i32 8
+// CHECK_BUFFERED-NEXT: store double [[TMP47]], ptr addrspace(1) [[PRINTBUFFNEXTPTR27]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR28:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR27]], i32 8
+// CHECK_BUFFERED-NEXT: store double [[TMP49]], ptr addrspace(1) [[PRINTBUFFNEXTPTR28]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR29:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR28]], i32 8
+// CHECK_BUFFERED-NEXT: store double [[TMP51]], ptr addrspace(1) [[PRINTBUFFNEXTPTR29]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR30:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR29]], i32 8
+// CHECK_BUFFERED-NEXT: store double [[TMP53]], ptr addrspace(1) [[PRINTBUFFNEXTPTR30]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR31:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR30]], i32 8
+// CHECK_BUFFERED-NEXT: [[TMP54:%.*]] = extractelement <16 x float> [[TMP21]], i64 0
+// CHECK_BUFFERED-NEXT: [[TMP55:%.*]] = fpext float [[TMP54]] to double
+// CHECK_BUFFERED-NEXT: [[TMP56:%.*]] = extractelement <16 x float> [[TMP21]], i64 1
+// CHECK_BUFFERED-NEXT: [[TMP57:%.*]] = fpext float [[TMP56]] to double
+// CHECK_BUFFERED-NEXT: [[TMP58:%.*]] = extractelement <16 x float> [[TMP21]], i64 2
+// CHECK_BUFFERED-NEXT: [[TMP59:%.*]] = fpext float [[TMP58]] to double
+// CHECK_BUFFERED-NEXT: [[TMP60:%.*]] = extractelement <16 x float> [[TMP21]], i64 3
+// CHECK_BUFFERED-NEXT: [[TMP61:%.*]] = fpext float [[TMP60]] to double
+// CHECK_BUFFERED-NEXT: [[TMP62:%.*]] = extractelement <16 x float> [[TMP21]], i64 4
+// CHECK_BUFFERED-NEXT: [[TMP63:%.*]] = fpext float [[TMP62]] to double
+// CHECK_BUFFERED-NEXT: [[TMP64:%.*]] = extractelement <16 x float> [[TMP21]], i64 5
+// CHECK_BUFFERED-NEXT: [[TMP65:%.*]] = fpext float [[TMP64]] to double
+// CHECK_BUFFERED-NEXT: [[TMP66:%.*]] = extractelement <16 x float> [[TMP21]], i64 6
+// CHECK_BUFFERED-NEXT: [[TMP67:%.*]] = fpext float [[TMP66]] to double
+// CHECK_BUFFERED-NEXT: [[TMP68:%.*]] = extractelement <16 x float> [[TMP21]], i64 7
+// CHECK_BUFFERED-NEXT: [[TMP69:%.*]] = fpext float [[TMP68]] to double
+// CHECK_BUFFERED-NEXT: [[TMP70:%.*]] = extractelement <16 x float> [[TMP21]], i64 8
+// CHECK_BUFFERED-NEXT: [[TMP71:%.*]] = fpext float [[TMP70]] to double
+// CHECK_BUFFERED-NEXT: [[TMP72:%.*]] = extractelement <16 x float> [[TMP21]], i64 9
+// CHECK_BUFFERED-NEXT: [[TMP73:%.*]] = fpext float [[TMP72]] to double
+// CHECK_BUFFERED-NEXT: [[TMP74:%.*]] = extractelement <16 x float> [[TMP21]], i64 10
+// CHECK_BUFFERED-NEXT: [[TMP75:%.*]] = fpext float [[TMP74]] to double
+// CHECK_BUFFERED-NEXT: [[TMP76:%.*]] = extractelement <16 x float> [[TMP21]], i64 11
+// CHECK_BUFFERED-NEXT: [[TMP77:%.*]] = fpext float [[TMP76]] to double
+// CHECK_BUFFERED-NEXT: [[TMP78:%.*]] = extractelement <16 x float> [[TMP21]], i64 12
+// CHECK_BUFFERED-NEXT: [[TMP79:%.*]] = fpext float [[TMP78]] to double
+// CHECK_BUFFERED-NEXT: [[TMP80:%.*]] = extractelement <16 x float> [[TMP21]], i64 13
+// CHECK_BUFFERED-NEXT: [[TMP81:%.*]] = fpext float [[TMP80]] to double
+// CHECK_BUFFERED-NEXT: [[TMP82:%.*]] = extractelement <16 x float> [[TMP21]], i64 14
+// CHECK_BUFFERED-NEXT: [[TMP83:%.*]] = fpext float [[TMP82]] to double
+// CHECK_BUFFERED-NEXT: [[TMP84:%.*]] = extractelement <16 x float> [[TMP21]], i64 15
+// CHECK_BUFFERED-NEXT: [[TMP85:%.*]] = fpext float [[TMP84]] to double
+// CHECK_BUFFERED-NEXT: store double [[TMP55]], ptr addrspace(1) [[PRINTBUFFNEXTPTR31]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR32:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR31]], i32 8
+// CHECK_BUFFERED-NEXT: store double [[TMP57]], ptr addrspace(1) [[PRINTBUFFNEXTPTR32]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR33:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR32]], i32 8
+// CHECK_BUFFERED-NEXT: store double [[TMP59]], ptr addrspace(1) [[PRINTBUFFNEXTPTR33]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR34:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR33]], i32 8
+// CHECK_BUFFERED-NEXT: store double [[TMP61]], ptr addrspace(1) [[PRINTBUFFNEXTPTR34]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR35:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR34]], i32 8
+// CHECK_BUFFERED-NEXT: store double [[TMP63]], ptr addrspace(1) [[PRINTBUFFNEXTPTR35]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR36:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR35]], i32 8
+// CHECK_BUFFERED-NEXT: store double [[TMP65]], ptr addrspace(1) [[PRINTBUFFNEXTPTR36]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR37:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR36]], i32 8
+// CHECK_BUFFERED-NEXT: store double [[TMP67]], ptr addrspace(1) [[PRINTBUFFNEXTPTR37]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR38:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR37]], i32 8
+// CHECK_BUFFERED-NEXT: store double [[TMP69]], ptr addrspace(1) [[PRINTBUFFNEXTPTR38]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR39:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR38]], i32 8
+// CHECK_BUFFERED-NEXT: store double [[TMP71]], ptr addrspace(1) [[PRINTBUFFNEXTPTR39]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR40:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR39]], i32 8
+// CHECK_BUFFERED-NEXT: store double [[TMP73]], ptr addrspace(1) [[PRINTBUFFNEXTPTR40]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR41:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR40]], i32 8
+// CHECK_BUFFERED-NEXT: store double [[TMP75]], ptr addrspace(1) [[PRINTBUFFNEXTPTR41]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR42:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR41]], i32 8
+// CHECK_BUFFERED-NEXT: store double [[TMP77]], ptr addrspace(1) [[PRINTBUFFNEXTPTR42]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR43:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR42]], i32 8
+// CHECK_BUFFERED-NEXT: store double [[TMP79]], ptr addrspace(1) [[PRINTBUFFNEXTPTR43]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR44:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR43]], i32 8
+// CHECK_BUFFERED-NEXT: store double [[TMP81]], ptr addrspace(1) [[PRINTBUFFNEXTPTR44]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR45:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR44]], i32 8
+// CHECK_BUFFERED-NEXT: store double [[TMP83]], ptr addrspace(1) [[PRINTBUFFNEXTPTR45]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR46:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR45]], i32 8
+// CHECK_BUFFERED-NEXT: store double [[TMP85]], ptr addrspace(1) [[PRINTBUFFNEXTPTR46]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR47:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR46]], i32 8
+// CHECK_BUFFERED-NEXT: br label [[END_BLOCK20]]
+//
+// CHECK_HOSTCALL-LABEL: @test_vector(
+// CHECK_HOSTCALL-NEXT: entry:
+// CHECK_HOSTCALL-NEXT: [[A:%.*]] = alloca <4 x i32>, align 16, addrspace(5)
+// CHECK_HOSTCALL-NEXT: [[DOTCOMPOUNDLITERAL:%.*]] = alloca <4 x i32>, align 16, addrspace(5)
+// CHECK_HOSTCALL-NEXT: [[B:%.*]] = alloca <3 x i32>, align 16, addrspace(5)
+// CHECK_HOSTCALL-NEXT: [[DOTCOMPOUNDLITERAL1:%.*]] = alloca <3 x i32>, align 16, addrspace(5)
+// CHECK_HOSTCALL-NEXT: [[VAR3:%.*]] = alloca <2 x float>, align 8, addrspace(5)
+// CHECK_HOSTCALL-NEXT: [[DOTCOMPOUNDLITERAL10:%.*]] = alloca <2 x float>, align 8, addrspace(5)
+// CHECK_HOSTCALL-NEXT: [[VAR4:%.*]] = alloca <8 x float>, align 32, addrspace(5)
+// CHECK_HOSTCALL-NEXT: [[VAR5:%.*]] = alloca <16 x float>, align 64, addrspace(5)
+// CHECK_HOSTCALL-NEXT: [[DOTCOMPOUNDLITERAL11:%.*]] = alloca <16 x float>, align 64, addrspace(5)
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.start.p5(i64 16, ptr addrspace(5) [[A]]) #[[ATTR3]]
+// CHECK_HOSTCALL-NEXT: store <4 x i32> <i32 1, i32 2, i32 3, i32 4>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL]], align 16, !tbaa [[TBAA17:![0-9]+]]
+// CHECK_HOSTCALL-NEXT: [[TMP0:%.*]] = load <4 x i32>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL]], align 16, !tbaa [[TBAA17]]
+// CHECK_HOSTCALL-NEXT: store <4 x i32> [[TMP0]], ptr addrspace(5) [[A]], align 16, !tbaa [[TBAA17]]
+// CHECK_HOSTCALL-NEXT: [[TMP1:%.*]] = load <4 x i32>, ptr addrspace(5) [[A]], align 16, !tbaa [[TBAA17]]
+// CHECK_HOSTCALL-NEXT: [[TMP2:%.*]] = call i64 @__ockl_printf_begin(i64 0)
+// CHECK_HOSTCALL-NEXT: br i1 icmp eq (ptr addrspacecast (ptr addrspace(4) @.str.4 to ptr), ptr null), label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]
+// CHECK_HOSTCALL: strlen.while:
+// CHECK_HOSTCALL-NEXT: [[TMP3:%.*]] = phi ptr [ addrspacecast (ptr addrspace(4) @.str.4 to ptr), [[ENTRY:%.*]] ], [ [[TMP4:%.*]], [[STRLEN_WHILE]] ]
+// CHECK_HOSTCALL-NEXT: [[TMP4]] = getelementptr i8, ptr [[TMP3]], i64 1
+// CHECK_HOSTCALL-NEXT: [[TMP5:%.*]] = load i8, ptr [[TMP3]], align 1
+// CHECK_HOSTCALL-NEXT: [[TMP6:%.*]] = icmp eq i8 [[TMP5]], 0
+// CHECK_HOSTCALL-NEXT: br i1 [[TMP6]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]]
+// CHECK_HOSTCALL: strlen.while.done:
+// CHECK_HOSTCALL-NEXT: [[TMP7:%.*]] = ptrtoint ptr [[TMP3]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP8:%.*]] = sub i64 [[TMP7]], ptrtoint (ptr addrspacecast (ptr addrspace(4) @.str.4 to ptr) to i64)
+// CHECK_HOSTCALL-NEXT: [[TMP9:%.*]] = add i64 [[TMP8]], 1
+// CHECK_HOSTCALL-NEXT: br label [[STRLEN_JOIN]]
+// CHECK_HOSTCALL: strlen.join:
+// CHECK_HOSTCALL-NEXT: [[TMP10:%.*]] = phi i64 [ [[TMP9]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]
+// CHECK_HOSTCALL-NEXT: [[TMP11:%.*]] = call i64 @__ockl_printf_append_string_n(i64 [[TMP2]], ptr addrspacecast (ptr addrspace(4) @.str.4 to ptr), i64 [[TMP10]], i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP12:%.*]] = extractelement <4 x i32> [[TMP1]], i64 0
+// CHECK_HOSTCALL-NEXT: [[TMP13:%.*]] = zext i32 [[TMP12]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP14:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP11]], i32 1, i64 [[TMP13]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP15:%.*]] = extractelement <4 x i32> [[TMP1]], i64 1
+// CHECK_HOSTCALL-NEXT: [[TMP16:%.*]] = zext i32 [[TMP15]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP17:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP14]], i32 1, i64 [[TMP16]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP18:%.*]] = extractelement <4 x i32> [[TMP1]], i64 2
+// CHECK_HOSTCALL-NEXT: [[TMP19:%.*]] = zext i32 [[TMP18]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP20:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP17]], i32 1, i64 [[TMP19]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP21:%.*]] = extractelement <4 x i32> [[TMP1]], i64 3
+// CHECK_HOSTCALL-NEXT: [[TMP22:%.*]] = zext i32 [[TMP21]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP23:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP20]], i32 1, i64 [[TMP22]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 1)
+// CHECK_HOSTCALL-NEXT: [[TMP24:%.*]] = trunc i64 [[TMP23]] to i32
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.start.p5(i64 16, ptr addrspace(5) [[B]]) #[[ATTR3]]
+// CHECK_HOSTCALL-NEXT: [[TMP25:%.*]] = load <4 x i32>, ptr addrspace(5) [[A]], align 16, !tbaa [[TBAA17]]
+// CHECK_HOSTCALL-NEXT: [[VECEXT:%.*]] = extractelement <4 x i32> [[TMP25]], i32 2
+// CHECK_HOSTCALL-NEXT: [[VECINIT:%.*]] = insertelement <3 x i32> <i32 1, i32 undef, i32 undef>, i32 [[VECEXT]], i32 1
+// CHECK_HOSTCALL-NEXT: [[VECINIT2:%.*]] = insertelement <3 x i32> [[VECINIT]], i32 3, i32 2
+// CHECK_HOSTCALL-NEXT: [[EXTRACTVEC:%.*]] = shufflevector <3 x i32> [[VECINIT2]], <3 x i32> poison, <4 x i32> <i32 0, i32 1, i32 2, i32 poison>
+// CHECK_HOSTCALL-NEXT: store <4 x i32> [[EXTRACTVEC]], ptr addrspace(5) [[DOTCOMPOUNDLITERAL1]], align 16, !tbaa [[TBAA17]]
+// CHECK_HOSTCALL-NEXT: [[LOADVEC4:%.*]] = load <4 x i32>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL1]], align 16
+// CHECK_HOSTCALL-NEXT: [[EXTRACTVEC3:%.*]] = shufflevector <4 x i32> [[LOADVEC4]], <4 x i32> poison, <3 x i32> <i32 0, i32 1, i32 2>
+// CHECK_HOSTCALL-NEXT: [[EXTRACTVEC4:%.*]] = shufflevector <3 x i32> [[EXTRACTVEC3]], <3 x i32> poison, <4 x i32> <i32 0, i32 1, i32 2, i32 poison>
+// CHECK_HOSTCALL-NEXT: store <4 x i32> [[EXTRACTVEC4]], ptr addrspace(5) [[B]], align 16, !tbaa [[TBAA17]]
+// CHECK_HOSTCALL-NEXT: [[LOADVEC45:%.*]] = load <4 x i32>, ptr addrspace(5) [[B]], align 16
+// CHECK_HOSTCALL-NEXT: [[EXTRACTVEC6:%.*]] = shufflevector <4 x i32> [[LOADVEC45]], <4 x i32> poison, <3 x i32> <i32 0, i32 1, i32 2>
+// CHECK_HOSTCALL-NEXT: [[TMP26:%.*]] = call i64 @__ockl_printf_begin(i64 0)
+// CHECK_HOSTCALL-NEXT: br i1 icmp eq (ptr addrspacecast (ptr addrspace(4) @.str.5 to ptr), ptr null), label [[STRLEN_JOIN7:%.*]], label [[STRLEN_WHILE8:%.*]]
+// CHECK_HOSTCALL: strlen.while8:
+// CHECK_HOSTCALL-NEXT: [[TMP27:%.*]] = phi ptr [ addrspacecast (ptr addrspace(4) @.str.5 to ptr), [[STRLEN_JOIN]] ], [ [[TMP28:%.*]], [[STRLEN_WHILE8]] ]
+// CHECK_HOSTCALL-NEXT: [[TMP28]] = getelementptr i8, ptr [[TMP27]], i64 1
+// CHECK_HOSTCALL-NEXT: [[TMP29:%.*]] = load i8, ptr [[TMP27]], align 1
+// CHECK_HOSTCALL-NEXT: [[TMP30:%.*]] = icmp eq i8 [[TMP29]], 0
+// CHECK_HOSTCALL-NEXT: br i1 [[TMP30]], label [[STRLEN_WHILE_DONE9:%.*]], label [[STRLEN_WHILE8]]
+// CHECK_HOSTCALL: strlen.while.done9:
+// CHECK_HOSTCALL-NEXT: [[TMP31:%.*]] = ptrtoint ptr [[TMP27]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP32:%.*]] = sub i64 [[TMP31]], ptrtoint (ptr addrspacecast (ptr addrspace(4) @.str.5 to ptr) to i64)
+// CHECK_HOSTCALL-NEXT: [[TMP33:%.*]] = add i64 [[TMP32]], 1
+// CHECK_HOSTCALL-NEXT: br label [[STRLEN_JOIN7]]
+// CHECK_HOSTCALL: strlen.join7:
+// CHECK_HOSTCALL-NEXT: [[TMP34:%.*]] = phi i64 [ [[TMP33]], [[STRLEN_WHILE_DONE9]] ], [ 0, [[STRLEN_JOIN]] ]
+// CHECK_HOSTCALL-NEXT: [[TMP35:%.*]] = call i64 @__ockl_printf_append_string_n(i64 [[TMP26]], ptr addrspacecast (ptr addrspace(4) @.str.5 to ptr), i64 [[TMP34]], i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP36:%.*]] = extractelement <3 x i32> [[EXTRACTVEC6]], i64 0
+// CHECK_HOSTCALL-NEXT: [[TMP37:%.*]] = zext i32 [[TMP36]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP38:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP35]], i32 1, i64 [[TMP37]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP39:%.*]] = extractelement <3 x i32> [[EXTRACTVEC6]], i64 1
+// CHECK_HOSTCALL-NEXT: [[TMP40:%.*]] = zext i32 [[TMP39]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP41:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP38]], i32 1, i64 [[TMP40]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP42:%.*]] = extractelement <3 x i32> [[EXTRACTVEC6]], i64 2
+// CHECK_HOSTCALL-NEXT: [[TMP43:%.*]] = zext i32 [[TMP42]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP44:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP41]], i32 1, i64 [[TMP43]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 1)
+// CHECK_HOSTCALL-NEXT: [[TMP45:%.*]] = trunc i64 [[TMP44]] to i32
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[VAR3]]) #[[ATTR3]]
+// CHECK_HOSTCALL-NEXT: store <2 x float> <float 0x4008F5C280000000, float 0x4003D70A40000000>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL10]], align 8, !tbaa [[TBAA17]]
+// CHECK_HOSTCALL-NEXT: [[TMP46:%.*]] = load <2 x float>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL10]], align 8, !tbaa [[TBAA17]]
+// CHECK_HOSTCALL-NEXT: store <2 x float> [[TMP46]], ptr addrspace(5) [[VAR3]], align 8, !tbaa [[TBAA17]]
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.start.p5(i64 32, ptr addrspace(5) [[VAR4]]) #[[ATTR3]]
+// CHECK_HOSTCALL-NEXT: store <8 x float> <float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000>, ptr addrspace(5) [[VAR4]], align 32, !tbaa [[TBAA17]]
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.start.p5(i64 64, ptr addrspace(5) [[VAR5]]) #[[ATTR3]]
+// CHECK_HOSTCALL-NEXT: store <16 x float> <float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL11]], align 64, !tbaa [[TBAA17]]
+// CHECK_HOSTCALL-NEXT: [[TMP47:%.*]] = load <16 x float>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL11]], align 64, !tbaa [[TBAA17]]
+// CHECK_HOSTCALL-NEXT: store <16 x float> [[TMP47]], ptr addrspace(5) [[VAR5]], align 64, !tbaa [[TBAA17]]
+// CHECK_HOSTCALL-NEXT: [[TMP48:%.*]] = load <2 x float>, ptr addrspace(5) [[VAR3]], align 8, !tbaa [[TBAA17]]
+// CHECK_HOSTCALL-NEXT: [[TMP49:%.*]] = load <8 x float>, ptr addrspace(5) [[VAR4]], align 32, !tbaa [[TBAA17]]
+// CHECK_HOSTCALL-NEXT: [[TMP50:%.*]] = load <16 x float>, ptr addrspace(5) [[VAR5]], align 64, !tbaa [[TBAA17]]
+// CHECK_HOSTCALL-NEXT: [[TMP51:%.*]] = call i64 @__ockl_printf_begin(i64 0)
+// CHECK_HOSTCALL-NEXT: br i1 icmp eq (ptr addrspacecast (ptr addrspace(4) @.str.6 to ptr), ptr null), label [[STRLEN_JOIN12:%.*]], label [[STRLEN_WHILE13:%.*]]
+// CHECK_HOSTCALL: strlen.while13:
+// CHECK_HOSTCALL-NEXT: [[TMP52:%.*]] = phi ptr [ addrspacecast (ptr addrspace(4) @.str.6 to ptr), [[STRLEN_JOIN7]] ], [ [[TMP53:%.*]], [[STRLEN_WHILE13]] ]
+// CHECK_HOSTCALL-NEXT: [[TMP53]] = getelementptr i8, ptr [[TMP52]], i64 1
+// CHECK_HOSTCALL-NEXT: [[TMP54:%.*]] = load i8, ptr [[TMP52]], align 1
+// CHECK_HOSTCALL-NEXT: [[TMP55:%.*]] = icmp eq i8 [[TMP54]], 0
+// CHECK_HOSTCALL-NEXT: br i1 [[TMP55]], label [[STRLEN_WHILE_DONE14:%.*]], label [[STRLEN_WHILE13]]
+// CHECK_HOSTCALL: strlen.while.done14:
+// CHECK_HOSTCALL-NEXT: [[TMP56:%.*]] = ptrtoint ptr [[TMP52]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP57:%.*]] = sub i64 [[TMP56]], ptrtoint (ptr addrspacecast (ptr addrspace(4) @.str.6 to ptr) to i64)
+// CHECK_HOSTCALL-NEXT: [[TMP58:%.*]] = add i64 [[TMP57]], 1
+// CHECK_HOSTCALL-NEXT: br label [[STRLEN_JOIN12]]
+// CHECK_HOSTCALL: strlen.join12:
+// CHECK_HOSTCALL-NEXT: [[TMP59:%.*]] = phi i64 [ [[TMP58]], [[STRLEN_WHILE_DONE14]] ], [ 0, [[STRLEN_JOIN7]] ]
+// CHECK_HOSTCALL-NEXT: [[TMP60:%.*]] = call i64 @__ockl_printf_append_string_n(i64 [[TMP51]], ptr addrspacecast (ptr addrspace(4) @.str.6 to ptr), i64 [[TMP59]], i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP61:%.*]] = extractelement <2 x float> [[TMP48]], i64 0
+// CHECK_HOSTCALL-NEXT: [[TMP62:%.*]] = fpext float [[TMP61]] to double
+// CHECK_HOSTCALL-NEXT: [[TMP63:%.*]] = bitcast double [[TMP62]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP64:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP60]], i32 1, i64 [[TMP63]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP65:%.*]] = extractelement <2 x float> [[TMP48]], i64 1
+// CHECK_HOSTCALL-NEXT: [[TMP66:%.*]] = fpext float [[TMP65]] to double
+// CHECK_HOSTCALL-NEXT: [[TMP67:%.*]] = bitcast double [[TMP66]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP68:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP64]], i32 1, i64 [[TMP67]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP69:%.*]] = extractelement <8 x float> [[TMP49]], i64 0
+// CHECK_HOSTCALL-NEXT: [[TMP70:%.*]] = fpext float [[TMP69]] to double
+// CHECK_HOSTCALL-NEXT: [[TMP71:%.*]] = bitcast double [[TMP70]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP72:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP68]], i32 1, i64 [[TMP71]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP73:%.*]] = extractelement <8 x float> [[TMP49]], i64 1
+// CHECK_HOSTCALL-NEXT: [[TMP74:%.*]] = fpext float [[TMP73]] to double
+// CHECK_HOSTCALL-NEXT: [[TMP75:%.*]] = bitcast double [[TMP74]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP76:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP72]], i32 1, i64 [[TMP75]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP77:%.*]] = extractelement <8 x float> [[TMP49]], i64 2
+// CHECK_HOSTCALL-NEXT: [[TMP78:%.*]] = fpext float [[TMP77]] to double
+// CHECK_HOSTCALL-NEXT: [[TMP79:%.*]] = bitcast double [[TMP78]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP80:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP76]], i32 1, i64 [[TMP79]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP81:%.*]] = extractelement <8 x float> [[TMP49]], i64 3
+// CHECK_HOSTCALL-NEXT: [[TMP82:%.*]] = fpext float [[TMP81]] to double
+// CHECK_HOSTCALL-NEXT: [[TMP83:%.*]] = bitcast double [[TMP82]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP84:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP80]], i32 1, i64 [[TMP83]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP85:%.*]] = extractelement <8 x float> [[TMP49]], i64 4
+// CHECK_HOSTCALL-NEXT: [[TMP86:%.*]] = fpext float [[TMP85]] to double
+// CHECK_HOSTCALL-NEXT: [[TMP87:%.*]] = bitcast double [[TMP86]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP88:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP84]], i32 1, i64 [[TMP87]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP89:%.*]] = extractelement <8 x float> [[TMP49]], i64 5
+// CHECK_HOSTCALL-NEXT: [[TMP90:%.*]] = fpext float [[TMP89]] to double
+// CHECK_HOSTCALL-NEXT: [[TMP91:%.*]] = bitcast double [[TMP90]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP92:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP88]], i32 1, i64 [[TMP91]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP93:%.*]] = extractelement <8 x float> [[TMP49]], i64 6
+// CHECK_HOSTCALL-NEXT: [[TMP94:%.*]] = fpext float [[TMP93]] to double
+// CHECK_HOSTCALL-NEXT: [[TMP95:%.*]] = bitcast double [[TMP94]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP96:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP92]], i32 1, i64 [[TMP95]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP97:%.*]] = extractelement <8 x float> [[TMP49]], i64 7
+// CHECK_HOSTCALL-NEXT: [[TMP98:%.*]] = fpext float [[TMP97]] to double
+// CHECK_HOSTCALL-NEXT: [[TMP99:%.*]] = bitcast double [[TMP98]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP100:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP96]], i32 1, i64 [[TMP99]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP101:%.*]] = extractelement <16 x float> [[TMP50]], i64 0
+// CHECK_HOSTCALL-NEXT: [[TMP102:%.*]] = fpext float [[TMP101]] to double
+// CHECK_HOSTCALL-NEXT: [[TMP103:%.*]] = bitcast double [[TMP102]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP104:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP100]], i32 1, i64 [[TMP103]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP105:%.*]] = extractelement <16 x float> [[TMP50]], i64 1
+// CHECK_HOSTCALL-NEXT: [[TMP106:%.*]] = fpext float [[TMP105]] to double
+// CHECK_HOSTCALL-NEXT: [[TMP107:%.*]] = bitcast double [[TMP106]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP108:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP104]], i32 1, i64 [[TMP107]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP109:%.*]] = extractelement <16 x float> [[TMP50]], i64 2
+// CHECK_HOSTCALL-NEXT: [[TMP110:%.*]] = fpext float [[TMP109]] to double
+// CHECK_HOSTCALL-NEXT: [[TMP111:%.*]] = bitcast double [[TMP110]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP112:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP108]], i32 1, i64 [[TMP111]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP113:%.*]] = extractelement <16 x float> [[TMP50]], i64 3
+// CHECK_HOSTCALL-NEXT: [[TMP114:%.*]] = fpext float [[TMP113]] to double
+// CHECK_HOSTCALL-NEXT: [[TMP115:%.*]] = bitcast double [[TMP114]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP116:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP112]], i32 1, i64 [[TMP115]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP117:%.*]] = extractelement <16 x float> [[TMP50]], i64 4
+// CHECK_HOSTCALL-NEXT: [[TMP118:%.*]] = fpext float [[TMP117]] to double
+// CHECK_HOSTCALL-NEXT: [[TMP119:%.*]] = bitcast double [[TMP118]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP120:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP116]], i32 1, i64 [[TMP119]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP121:%.*]] = extractelement <16 x float> [[TMP50]], i64 5
+// CHECK_HOSTCALL-NEXT: [[TMP122:%.*]] = fpext float [[TMP121]] to double
+// CHECK_HOSTCALL-NEXT: [[TMP123:%.*]] = bitcast double [[TMP122]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP124:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP120]], i32 1, i64 [[TMP123]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP125:%.*]] = extractelement <16 x float> [[TMP50]], i64 6
+// CHECK_HOSTCALL-NEXT: [[TMP126:%.*]] = fpext float [[TMP125]] to double
+// CHECK_HOSTCALL-NEXT: [[TMP127:%.*]] = bitcast double [[TMP126]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP128:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP124]], i32 1, i64 [[TMP127]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP129:%.*]] = extractelement <16 x float> [[TMP50]], i64 7
+// CHECK_HOSTCALL-NEXT: [[TMP130:%.*]] = fpext float [[TMP129]] to double
+// CHECK_HOSTCALL-NEXT: [[TMP131:%.*]] = bitcast double [[TMP130]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP132:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP128]], i32 1, i64 [[TMP131]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP133:%.*]] = extractelement <16 x float> [[TMP50]], i64 8
+// CHECK_HOSTCALL-NEXT: [[TMP134:%.*]] = fpext float [[TMP133]] to double
+// CHECK_HOSTCALL-NEXT: [[TMP135:%.*]] = bitcast double [[TMP134]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP136:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP132]], i32 1, i64 [[TMP135]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP137:%.*]] = extractelement <16 x float> [[TMP50]], i64 9
+// CHECK_HOSTCALL-NEXT: [[TMP138:%.*]] = fpext float [[TMP137]] to double
+// CHECK_HOSTCALL-NEXT: [[TMP139:%.*]] = bitcast double [[TMP138]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP140:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP136]], i32 1, i64 [[TMP139]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP141:%.*]] = extractelement <16 x float> [[TMP50]], i64 10
+// CHECK_HOSTCALL-NEXT: [[TMP142:%.*]] = fpext float [[TMP141]] to double
+// CHECK_HOSTCALL-NEXT: [[TMP143:%.*]] = bitcast double [[TMP142]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP144:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP140]], i32 1, i64 [[TMP143]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP145:%.*]] = extractelement <16 x float> [[TMP50]], i64 11
+// CHECK_HOSTCALL-NEXT: [[TMP146:%.*]] = fpext float [[TMP145]] to double
+// CHECK_HOSTCALL-NEXT: [[TMP147:%.*]] = bitcast double [[TMP146]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP148:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP144]], i32 1, i64 [[TMP147]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP149:%.*]] = extractelement <16 x float> [[TMP50]], i64 12
+// CHECK_HOSTCALL-NEXT: [[TMP150:%.*]] = fpext float [[TMP149]] to double
+// CHECK_HOSTCALL-NEXT: [[TMP151:%.*]] = bitcast double [[TMP150]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP152:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP148]], i32 1, i64 [[TMP151]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP153:%.*]] = extractelement <16 x float> [[TMP50]], i64 13
+// CHECK_HOSTCALL-NEXT: [[TMP154:%.*]] = fpext float [[TMP153]] to double
+// CHECK_HOSTCALL-NEXT: [[TMP155:%.*]] = bitcast double [[TMP154]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP156:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP152]], i32 1, i64 [[TMP155]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP157:%.*]] = extractelement <16 x float> [[TMP50]], i64 14
+// CHECK_HOSTCALL-NEXT: [[TMP158:%.*]] = fpext float [[TMP157]] to double
+// CHECK_HOSTCALL-NEXT: [[TMP159:%.*]] = bitcast double [[TMP158]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP160:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP156]], i32 1, i64 [[TMP159]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP161:%.*]] = extractelement <16 x float> [[TMP50]], i64 15
+// CHECK_HOSTCALL-NEXT: [[TMP162:%.*]] = fpext float [[TMP161]] to double
+// CHECK_HOSTCALL-NEXT: [[TMP163:%.*]] = bitcast double [[TMP162]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP164:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP160]], i32 1, i64 [[TMP163]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 1)
+// CHECK_HOSTCALL-NEXT: [[TMP165:%.*]] = trunc i64 [[TMP164]] to i32
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.end.p5(i64 64, ptr addrspace(5) [[VAR5]]) #[[ATTR3]]
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.end.p5(i64 32, ptr addrspace(5) [[VAR4]]) #[[ATTR3]]
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[VAR3]]) #[[ATTR3]]
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.end.p5(i64 16, ptr addrspace(5) [[B]]) #[[ATTR3]]
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.end.p5(i64 16, ptr addrspace(5) [[A]]) #[[ATTR3]]
+// CHECK_HOSTCALL-NEXT: ret void
+//
+__kernel void test_vector() {
+
+int4 a = (int4)(1, 2, 3, 4);
+printf("%v4d", a);
+
+int3 b = (int3)(1, a[2], 3);
+printf("%v3d", b);
+
+float2 var3 = (float2)(3.12, 2.48);
+float8 var4 = (float8)(3.14);
+float16 var5 = (float16)((float8)(3.12), (float8)(2.32));
+
+printf("%v2hlf, %v8hlf, %v16hlf", var3, var4, var5);
+
+}
diff --git a/llvm/include/llvm/Transforms/Utils/AMDGPUEmitPrintf.h b/llvm/include/llvm/Transforms/Utils/AMDGPUEmitPrintf.h
index 55e92c37a1676..087d34ad7d905 100644
--- a/llvm/include/llvm/Transforms/Utils/AMDGPUEmitPrintf.h
+++ b/llvm/include/llvm/Transforms/Utils/AMDGPUEmitPrintf.h
@@ -19,7 +19,7 @@
namespace llvm {
Value *emitAMDGPUPrintfCall(IRBuilder<> &Builder, ArrayRef<Value *> Args,
- bool isBuffered);
+ StringRef FmtStr, bool isBuffered);
} // end namespace llvm
diff --git a/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp b/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp
index bc3f20443504f..86af790f5fd76 100644
--- a/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp
+++ b/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp
@@ -173,24 +173,21 @@ static Value *appendString(IRBuilder<> &Builder, Value *Desc, Value *Arg,
static Value *appendVectorArg(IRBuilder<> &Builder, Value *Desc, Value *Arg,
bool IsLast, bool IsBuffered) {
- assert(Arg->getType()->isVectorTy() && "incorrent append* function");
- auto VectorTy = dyn_cast<FixedVectorType>(Arg->getType());
+ assert(Arg->getType()->isVectorTy() && "incorrect append* function");
+ auto VectorTy = cast<FixedVectorType>(Arg->getType());
auto Zero = Builder.getInt64(0);
- if (VectorTy) {
- for (unsigned int i = 0; i < VectorTy->getNumElements() - 1; i++) {
- auto Val = Builder.CreateExtractElement(Arg, i);
- Desc = callAppendArgs(Builder, Desc, 1,
- fitArgInto64Bits(Builder, Val, IsBuffered), Zero,
- Zero, Zero, Zero, Zero, Zero, false);
- }
-
- Value* Val =
- Builder.CreateExtractElement(Arg, VectorTy->getNumElements() - 1);
- return callAppendArgs(Builder, Desc, 1,
+ for (unsigned int i = 0; i < VectorTy->getNumElements() - 1; i++) {
+ auto Val = Builder.CreateExtractElement(Arg, i);
+ Desc = callAppendArgs(Builder, Desc, 1,
fitArgInto64Bits(Builder, Val, IsBuffered), Zero,
- Zero, Zero, Zero, Zero, Zero, IsLast);
+ Zero, Zero, Zero, Zero, Zero, false);
}
- return nullptr;
+
+ Value *Val =
+ Builder.CreateExtractElement(Arg, VectorTy->getNumElements() - 1);
+ return callAppendArgs(Builder, Desc, 1,
+ fitArgInto64Bits(Builder, Val, IsBuffered), Zero, Zero,
+ Zero, Zero, Zero, Zero, IsLast);
}
static Value *processArg(IRBuilder<> &Builder, Value *Desc, Value *Arg,
@@ -210,8 +207,10 @@ static Value *processArg(IRBuilder<> &Builder, Value *Desc, Value *Arg,
return appendArg(Builder, Desc, Arg, IsLast, IsBuffered);
}
-// Scan the format string to locate all specifiers, and mark the ones that
-// specify a string, i.e, the "%s" specifier with optional '*' characters.
+// Scan the format string to locate all specifiers and OCL vectors,
+// and mark the ones that specify a string/vector,
+// i.e, the "%s" specifier with optional '*' characters
+// or "%v" specifier.
static void locateCStringsAndVectors(SparseBitVector<8> &BV,
SparseBitVector<8> &OV, StringRef Str) {
static const char ConvSpecifiers[] = "diouxXfFeEgGaAcspn";
@@ -224,12 +223,14 @@ static void locateCStringsAndVectors(SparseBitVector<8> &BV,
SpecPos += 2;
continue;
}
- if (Str.find_first_of("v", SpecPos) != StringRef::npos)
- OV.set(ArgIdx);
auto SpecEnd = Str.find_first_of(ConvSpecifiers, SpecPos);
if (SpecEnd == StringRef::npos)
return;
auto Spec = Str.slice(SpecPos, SpecEnd + 1);
+
+ if ((Spec.find_first_of("v")) != StringRef::npos)
+ OV.set(ArgIdx);
+
ArgIdx += Spec.count('*');
if (Str[SpecEnd] == 's') {
BV.set(ArgIdx);
@@ -426,8 +427,8 @@ static void callBufferedPrintfArgPush(
}
} else {
if (OCLVectors.test(i)) {
- auto VectorTy = dyn_cast<FixedVectorType>(Args[i]->getType());
- auto VecArg = Args[i];
+ auto VectorTy = cast<FixedVectorType>(Args[i]->getType());
+ Value *VecArg = Args[i];
for (unsigned int Num = 0; Num < VectorTy->getNumElements(); Num++) {
auto Val = Builder.CreateExtractElement(VecArg, Num);
WhatToStore.push_back(
@@ -456,16 +457,16 @@ static void callBufferedPrintfArgPush(
}
Value *llvm::emitAMDGPUPrintfCall(IRBuilder<> &Builder, ArrayRef<Value *> Args,
- bool IsBuffered) {
+ StringRef FmtStr, bool IsBuffered) {
auto NumOps = Args.size();
assert(NumOps >= 1);
auto Fmt = Args[0];
SparseBitVector<8> SpecIsCString;
SparseBitVector<8> OCLVectors;
- StringRef FmtStr;
+ bool IsConstFmtStr = !FmtStr.empty();
- if (getConstantStringInfo(Fmt, FmtStr))
+ if (IsConstFmtStr)
locateCStringsAndVectors(SpecIsCString, OCLVectors, FmtStr);
if (IsBuffered) {
@@ -474,7 +475,6 @@ Value *llvm::emitAMDGPUPrintfCall(IRBuilder<> &Builder, ArrayRef<Value *> Args,
LLVMContext &Ctx = Builder.getContext();
auto Int8Ty = Builder.getInt8Ty();
auto Int32Ty = Builder.getInt32Ty();
- bool IsConstFmtStr = !FmtStr.empty();
Value *ArgSize = nullptr;
Value *Ptr = callBufferedPrintfStart(Builder, Args, Fmt, IsConstFmtStr,
>From aaff6ddb2082d11e5aa0c5b60baff9188e771518 Mon Sep 17 00:00:00 2001
From: Vikram <Vikram.Hegde at amd.com>
Date: Mon, 4 Dec 2023 06:03:24 +0000
Subject: [PATCH 5/5] Handled further review comments
---
clang/include/clang/Basic/BuiltinsAMDGPU.def | 5 +-
clang/test/CodeGenOpenCL/amdgpu-printf.cl | 440 ++++++++++--------
.../lib/Transforms/Utils/AMDGPUEmitPrintf.cpp | 5 +-
3 files changed, 256 insertions(+), 194 deletions(-)
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index 1799c72806bfd..240841fd283f9 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -410,7 +410,10 @@ TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_fp8_f32, "iffiIb", "nc", "fp8-insts")
TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_bf8_f32, "ifiiIi", "nc", "fp8-insts")
TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_fp8_f32, "ifiiIi", "nc", "fp8-insts")
-// OpenCL
+// OpenCL printf has the following signature
+// int printf(__constant const char* st, ...) __attribute__((format(printf, 1, 2)));
+// The "__constant" address space corresponds to number 4 in LLVM IR for AMDGPU.
+// Following entry makes sure printf is recognized as builtin for OCL inputs.
LANGBUILTIN(printf, "icC*4.", "fp:0:", ALL_OCL_LANGUAGES)
#undef BUILTIN
diff --git a/clang/test/CodeGenOpenCL/amdgpu-printf.cl b/clang/test/CodeGenOpenCL/amdgpu-printf.cl
index 3342749e24722..c42fb1086e3ab 100644
--- a/clang/test/CodeGenOpenCL/amdgpu-printf.cl
+++ b/clang/test/CodeGenOpenCL/amdgpu-printf.cl
@@ -61,8 +61,8 @@ __kernel void test_printf_noargs() {
// CHECK_BUFFERED-LABEL: @test_printf_int(
// CHECK_BUFFERED-NEXT: entry:
// CHECK_BUFFERED-NEXT: [[I_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
-// CHECK_BUFFERED-NEXT: store i32 [[I:%.*]], ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA16:![0-9]+]]
-// CHECK_BUFFERED-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA16]]
+// CHECK_BUFFERED-NEXT: store i32 [[I:%.*]], ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA17:![0-9]+]]
+// CHECK_BUFFERED-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA17]]
// CHECK_BUFFERED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 20)
// CHECK_BUFFERED-NEXT: [[TMP1:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
// CHECK_BUFFERED-NEXT: br i1 [[TMP1]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]
@@ -133,51 +133,58 @@ __kernel void test_printf_int(int i) {
// CHECK_BUFFERED-NEXT: entry:
// CHECK_BUFFERED-NEXT: [[I_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK_BUFFERED-NEXT: [[S:%.*]] = alloca [4 x i8], align 1, addrspace(5)
-// CHECK_BUFFERED-NEXT: store i32 [[I:%.*]], ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA16]]
+// CHECK_BUFFERED-NEXT: store i32 [[I:%.*]], ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA17]]
// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[S]]) #[[ATTR1:[0-9]+]]
-// CHECK_BUFFERED-NEXT: call void @llvm.memcpy.p5.p4.i64(ptr addrspace(5) align 1 [[S]], ptr addrspace(4) align 1 @__const.test_printf_str_int.s, i64 4, i1 false)
+// CHECK_BUFFERED-NEXT: [[TMP0:%.*]] = getelementptr i8, ptr addrspace(5) [[S]], i64 0
+// CHECK_BUFFERED-NEXT: store i8 102, ptr addrspace(5) [[TMP0]], align 1
+// CHECK_BUFFERED-NEXT: [[TMP1:%.*]] = getelementptr i8, ptr addrspace(5) [[S]], i64 1
+// CHECK_BUFFERED-NEXT: store i8 111, ptr addrspace(5) [[TMP1]], align 1
+// CHECK_BUFFERED-NEXT: [[TMP2:%.*]] = getelementptr i8, ptr addrspace(5) [[S]], i64 2
+// CHECK_BUFFERED-NEXT: store i8 111, ptr addrspace(5) [[TMP2]], align 1
+// CHECK_BUFFERED-NEXT: [[TMP3:%.*]] = getelementptr i8, ptr addrspace(5) [[S]], i64 3
+// CHECK_BUFFERED-NEXT: store i8 0, ptr addrspace(5) [[TMP3]], align 1
// CHECK_BUFFERED-NEXT: [[ARRAYDECAY:%.*]] = getelementptr inbounds [4 x i8], ptr addrspace(5) [[S]], i64 0, i64 0
-// CHECK_BUFFERED-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA16]]
-// CHECK_BUFFERED-NEXT: [[TMP1:%.*]] = addrspacecast ptr addrspace(5) [[ARRAYDECAY]] to ptr
-// CHECK_BUFFERED-NEXT: [[TMP2:%.*]] = icmp eq ptr [[TMP1]], null
-// CHECK_BUFFERED-NEXT: br i1 [[TMP2]], label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]
+// CHECK_BUFFERED-NEXT: [[TMP4:%.*]] = load i32, ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA17]]
+// CHECK_BUFFERED-NEXT: [[TMP5:%.*]] = addrspacecast ptr addrspace(5) [[ARRAYDECAY]] to ptr
+// CHECK_BUFFERED-NEXT: [[TMP6:%.*]] = icmp eq ptr [[TMP5]], null
+// CHECK_BUFFERED-NEXT: br i1 [[TMP6]], label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]
// CHECK_BUFFERED: strlen.while:
-// CHECK_BUFFERED-NEXT: [[TMP3:%.*]] = phi ptr [ [[TMP1]], [[ENTRY:%.*]] ], [ [[TMP4:%.*]], [[STRLEN_WHILE]] ]
-// CHECK_BUFFERED-NEXT: [[TMP4]] = getelementptr i8, ptr [[TMP3]], i64 1
-// CHECK_BUFFERED-NEXT: [[TMP5:%.*]] = load i8, ptr [[TMP3]], align 1
-// CHECK_BUFFERED-NEXT: [[TMP6:%.*]] = icmp eq i8 [[TMP5]], 0
-// CHECK_BUFFERED-NEXT: br i1 [[TMP6]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]]
+// CHECK_BUFFERED-NEXT: [[TMP7:%.*]] = phi ptr [ [[TMP5]], [[ENTRY:%.*]] ], [ [[TMP8:%.*]], [[STRLEN_WHILE]] ]
+// CHECK_BUFFERED-NEXT: [[TMP8]] = getelementptr i8, ptr [[TMP7]], i64 1
+// CHECK_BUFFERED-NEXT: [[TMP9:%.*]] = load i8, ptr [[TMP7]], align 1
+// CHECK_BUFFERED-NEXT: [[TMP10:%.*]] = icmp eq i8 [[TMP9]], 0
+// CHECK_BUFFERED-NEXT: br i1 [[TMP10]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]]
// CHECK_BUFFERED: strlen.while.done:
-// CHECK_BUFFERED-NEXT: [[TMP7:%.*]] = ptrtoint ptr [[TMP1]] to i64
-// CHECK_BUFFERED-NEXT: [[TMP8:%.*]] = ptrtoint ptr [[TMP3]] to i64
-// CHECK_BUFFERED-NEXT: [[TMP9:%.*]] = sub i64 [[TMP8]], [[TMP7]]
-// CHECK_BUFFERED-NEXT: [[TMP10:%.*]] = add i64 [[TMP9]], 1
+// CHECK_BUFFERED-NEXT: [[TMP11:%.*]] = ptrtoint ptr [[TMP5]] to i64
+// CHECK_BUFFERED-NEXT: [[TMP12:%.*]] = ptrtoint ptr [[TMP7]] to i64
+// CHECK_BUFFERED-NEXT: [[TMP13:%.*]] = sub i64 [[TMP12]], [[TMP11]]
+// CHECK_BUFFERED-NEXT: [[TMP14:%.*]] = add i64 [[TMP13]], 1
// CHECK_BUFFERED-NEXT: br label [[STRLEN_JOIN]]
// CHECK_BUFFERED: strlen.join:
-// CHECK_BUFFERED-NEXT: [[TMP11:%.*]] = phi i64 [ [[TMP10]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]
-// CHECK_BUFFERED-NEXT: [[TMP12:%.*]] = add i64 [[TMP11]], 7
-// CHECK_BUFFERED-NEXT: [[TMP13:%.*]] = and i64 [[TMP12]], 4294967288
-// CHECK_BUFFERED-NEXT: [[TMP14:%.*]] = add i64 [[TMP13]], 20
-// CHECK_BUFFERED-NEXT: [[TMP15:%.*]] = trunc i64 [[TMP14]] to i32
-// CHECK_BUFFERED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[TMP15]])
-// CHECK_BUFFERED-NEXT: [[TMP16:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
-// CHECK_BUFFERED-NEXT: br i1 [[TMP16]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]
+// CHECK_BUFFERED-NEXT: [[TMP15:%.*]] = phi i64 [ [[TMP14]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]
+// CHECK_BUFFERED-NEXT: [[TMP16:%.*]] = add i64 [[TMP15]], 7
+// CHECK_BUFFERED-NEXT: [[TMP17:%.*]] = and i64 [[TMP16]], 4294967288
+// CHECK_BUFFERED-NEXT: [[TMP18:%.*]] = add i64 [[TMP17]], 20
+// CHECK_BUFFERED-NEXT: [[TMP19:%.*]] = trunc i64 [[TMP18]] to i32
+// CHECK_BUFFERED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[TMP19]])
+// CHECK_BUFFERED-NEXT: [[TMP20:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
+// CHECK_BUFFERED-NEXT: br i1 [[TMP20]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]
// CHECK_BUFFERED: end.block:
-// CHECK_BUFFERED-NEXT: [[TMP17:%.*]] = xor i1 [[TMP16]], true
-// CHECK_BUFFERED-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP17]] to i32
+// CHECK_BUFFERED-NEXT: [[TMP21:%.*]] = xor i1 [[TMP20]], true
+// CHECK_BUFFERED-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP21]] to i32
// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[S]]) #[[ATTR1]]
// CHECK_BUFFERED-NEXT: ret void
// CHECK_BUFFERED: argpush.block:
-// CHECK_BUFFERED-NEXT: [[TMP18:%.*]] = shl i32 [[TMP15]], 2
-// CHECK_BUFFERED-NEXT: [[TMP19:%.*]] = or i32 [[TMP18]], 2
-// CHECK_BUFFERED-NEXT: store i32 [[TMP19]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
-// CHECK_BUFFERED-NEXT: [[TMP20:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
-// CHECK_BUFFERED-NEXT: store i64 -2942283388077972797, ptr addrspace(1) [[TMP20]], align 8
-// CHECK_BUFFERED-NEXT: [[TMP21:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP20]], i32 8
-// CHECK_BUFFERED-NEXT: call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 1 [[TMP21]], ptr align 1 [[TMP1]], i64 [[TMP11]], i1 false)
-// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP21]], i64 [[TMP13]]
-// CHECK_BUFFERED-NEXT: [[TMP22:%.*]] = zext i32 [[TMP0]] to i64
-// CHECK_BUFFERED-NEXT: store i64 [[TMP22]], ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8
+// CHECK_BUFFERED-NEXT: [[TMP22:%.*]] = shl i32 [[TMP19]], 2
+// CHECK_BUFFERED-NEXT: [[TMP23:%.*]] = or i32 [[TMP22]], 2
+// CHECK_BUFFERED-NEXT: store i32 [[TMP23]], ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
+// CHECK_BUFFERED-NEXT: [[TMP24:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
+// CHECK_BUFFERED-NEXT: store i64 -2942283388077972797, ptr addrspace(1) [[TMP24]], align 8
+// CHECK_BUFFERED-NEXT: [[TMP25:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP24]], i32 8
+// CHECK_BUFFERED-NEXT: call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) align 1 [[TMP25]], ptr align 1 [[TMP5]], i64 [[TMP15]], i1 false)
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP25]], i64 [[TMP17]]
+// CHECK_BUFFERED-NEXT: [[TMP26:%.*]] = zext i32 [[TMP4]] to i64
+// CHECK_BUFFERED-NEXT: store i64 [[TMP26]], ptr addrspace(1) [[PRINTBUFFNEXTPTR]], align 8
// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR]], i32 8
// CHECK_BUFFERED-NEXT: br label [[END_BLOCK]]
//
@@ -186,48 +193,55 @@ __kernel void test_printf_int(int i) {
// CHECK_HOSTCALL-NEXT: [[I_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK_HOSTCALL-NEXT: [[S:%.*]] = alloca [4 x i8], align 1, addrspace(5)
// CHECK_HOSTCALL-NEXT: store i32 [[I:%.*]], ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA9]]
-// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[S]]) #[[ATTR3:[0-9]+]]
-// CHECK_HOSTCALL-NEXT: call void @llvm.memcpy.p5.p4.i64(ptr addrspace(5) align 1 [[S]], ptr addrspace(4) align 1 @__const.test_printf_str_int.s, i64 4, i1 false)
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[S]]) #[[ATTR2:[0-9]+]]
+// CHECK_HOSTCALL-NEXT: [[TMP0:%.*]] = getelementptr i8, ptr addrspace(5) [[S]], i64 0
+// CHECK_HOSTCALL-NEXT: store i8 102, ptr addrspace(5) [[TMP0]], align 1
+// CHECK_HOSTCALL-NEXT: [[TMP1:%.*]] = getelementptr i8, ptr addrspace(5) [[S]], i64 1
+// CHECK_HOSTCALL-NEXT: store i8 111, ptr addrspace(5) [[TMP1]], align 1
+// CHECK_HOSTCALL-NEXT: [[TMP2:%.*]] = getelementptr i8, ptr addrspace(5) [[S]], i64 2
+// CHECK_HOSTCALL-NEXT: store i8 111, ptr addrspace(5) [[TMP2]], align 1
+// CHECK_HOSTCALL-NEXT: [[TMP3:%.*]] = getelementptr i8, ptr addrspace(5) [[S]], i64 3
+// CHECK_HOSTCALL-NEXT: store i8 0, ptr addrspace(5) [[TMP3]], align 1
// CHECK_HOSTCALL-NEXT: [[ARRAYDECAY:%.*]] = getelementptr inbounds [4 x i8], ptr addrspace(5) [[S]], i64 0, i64 0
-// CHECK_HOSTCALL-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA9]]
-// CHECK_HOSTCALL-NEXT: [[TMP1:%.*]] = addrspacecast ptr addrspace(5) [[ARRAYDECAY]] to ptr
-// CHECK_HOSTCALL-NEXT: [[TMP2:%.*]] = call i64 @__ockl_printf_begin(i64 0)
+// CHECK_HOSTCALL-NEXT: [[TMP4:%.*]] = load i32, ptr addrspace(5) [[I_ADDR]], align 4, !tbaa [[TBAA9]]
+// CHECK_HOSTCALL-NEXT: [[TMP5:%.*]] = addrspacecast ptr addrspace(5) [[ARRAYDECAY]] to ptr
+// CHECK_HOSTCALL-NEXT: [[TMP6:%.*]] = call i64 @__ockl_printf_begin(i64 0)
// CHECK_HOSTCALL-NEXT: br i1 icmp eq (ptr addrspacecast (ptr addrspace(4) @.str.2 to ptr), ptr null), label [[STRLEN_JOIN:%.*]], label [[STRLEN_WHILE:%.*]]
// CHECK_HOSTCALL: strlen.while:
-// CHECK_HOSTCALL-NEXT: [[TMP3:%.*]] = phi ptr [ addrspacecast (ptr addrspace(4) @.str.2 to ptr), [[ENTRY:%.*]] ], [ [[TMP4:%.*]], [[STRLEN_WHILE]] ]
-// CHECK_HOSTCALL-NEXT: [[TMP4]] = getelementptr i8, ptr [[TMP3]], i64 1
-// CHECK_HOSTCALL-NEXT: [[TMP5:%.*]] = load i8, ptr [[TMP3]], align 1
-// CHECK_HOSTCALL-NEXT: [[TMP6:%.*]] = icmp eq i8 [[TMP5]], 0
-// CHECK_HOSTCALL-NEXT: br i1 [[TMP6]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]]
+// CHECK_HOSTCALL-NEXT: [[TMP7:%.*]] = phi ptr [ addrspacecast (ptr addrspace(4) @.str.2 to ptr), [[ENTRY:%.*]] ], [ [[TMP8:%.*]], [[STRLEN_WHILE]] ]
+// CHECK_HOSTCALL-NEXT: [[TMP8]] = getelementptr i8, ptr [[TMP7]], i64 1
+// CHECK_HOSTCALL-NEXT: [[TMP9:%.*]] = load i8, ptr [[TMP7]], align 1
+// CHECK_HOSTCALL-NEXT: [[TMP10:%.*]] = icmp eq i8 [[TMP9]], 0
+// CHECK_HOSTCALL-NEXT: br i1 [[TMP10]], label [[STRLEN_WHILE_DONE:%.*]], label [[STRLEN_WHILE]]
// CHECK_HOSTCALL: strlen.while.done:
-// CHECK_HOSTCALL-NEXT: [[TMP7:%.*]] = ptrtoint ptr [[TMP3]] to i64
-// CHECK_HOSTCALL-NEXT: [[TMP8:%.*]] = sub i64 [[TMP7]], ptrtoint (ptr addrspacecast (ptr addrspace(4) @.str.2 to ptr) to i64)
-// CHECK_HOSTCALL-NEXT: [[TMP9:%.*]] = add i64 [[TMP8]], 1
+// CHECK_HOSTCALL-NEXT: [[TMP11:%.*]] = ptrtoint ptr [[TMP7]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP12:%.*]] = sub i64 [[TMP11]], ptrtoint (ptr addrspacecast (ptr addrspace(4) @.str.2 to ptr) to i64)
+// CHECK_HOSTCALL-NEXT: [[TMP13:%.*]] = add i64 [[TMP12]], 1
// CHECK_HOSTCALL-NEXT: br label [[STRLEN_JOIN]]
// CHECK_HOSTCALL: strlen.join:
-// CHECK_HOSTCALL-NEXT: [[TMP10:%.*]] = phi i64 [ [[TMP9]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]
-// CHECK_HOSTCALL-NEXT: [[TMP11:%.*]] = call i64 @__ockl_printf_append_string_n(i64 [[TMP2]], ptr addrspacecast (ptr addrspace(4) @.str.2 to ptr), i64 [[TMP10]], i32 0)
-// CHECK_HOSTCALL-NEXT: [[TMP12:%.*]] = icmp eq ptr [[TMP1]], null
-// CHECK_HOSTCALL-NEXT: br i1 [[TMP12]], label [[STRLEN_JOIN1:%.*]], label [[STRLEN_WHILE2:%.*]]
+// CHECK_HOSTCALL-NEXT: [[TMP14:%.*]] = phi i64 [ [[TMP13]], [[STRLEN_WHILE_DONE]] ], [ 0, [[ENTRY]] ]
+// CHECK_HOSTCALL-NEXT: [[TMP15:%.*]] = call i64 @__ockl_printf_append_string_n(i64 [[TMP6]], ptr addrspacecast (ptr addrspace(4) @.str.2 to ptr), i64 [[TMP14]], i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP16:%.*]] = icmp eq ptr [[TMP5]], null
+// CHECK_HOSTCALL-NEXT: br i1 [[TMP16]], label [[STRLEN_JOIN1:%.*]], label [[STRLEN_WHILE2:%.*]]
// CHECK_HOSTCALL: strlen.while2:
-// CHECK_HOSTCALL-NEXT: [[TMP13:%.*]] = phi ptr [ [[TMP1]], [[STRLEN_JOIN]] ], [ [[TMP14:%.*]], [[STRLEN_WHILE2]] ]
-// CHECK_HOSTCALL-NEXT: [[TMP14]] = getelementptr i8, ptr [[TMP13]], i64 1
-// CHECK_HOSTCALL-NEXT: [[TMP15:%.*]] = load i8, ptr [[TMP13]], align 1
-// CHECK_HOSTCALL-NEXT: [[TMP16:%.*]] = icmp eq i8 [[TMP15]], 0
-// CHECK_HOSTCALL-NEXT: br i1 [[TMP16]], label [[STRLEN_WHILE_DONE3:%.*]], label [[STRLEN_WHILE2]]
+// CHECK_HOSTCALL-NEXT: [[TMP17:%.*]] = phi ptr [ [[TMP5]], [[STRLEN_JOIN]] ], [ [[TMP18:%.*]], [[STRLEN_WHILE2]] ]
+// CHECK_HOSTCALL-NEXT: [[TMP18]] = getelementptr i8, ptr [[TMP17]], i64 1
+// CHECK_HOSTCALL-NEXT: [[TMP19:%.*]] = load i8, ptr [[TMP17]], align 1
+// CHECK_HOSTCALL-NEXT: [[TMP20:%.*]] = icmp eq i8 [[TMP19]], 0
+// CHECK_HOSTCALL-NEXT: br i1 [[TMP20]], label [[STRLEN_WHILE_DONE3:%.*]], label [[STRLEN_WHILE2]]
// CHECK_HOSTCALL: strlen.while.done3:
-// CHECK_HOSTCALL-NEXT: [[TMP17:%.*]] = ptrtoint ptr [[TMP1]] to i64
-// CHECK_HOSTCALL-NEXT: [[TMP18:%.*]] = ptrtoint ptr [[TMP13]] to i64
-// CHECK_HOSTCALL-NEXT: [[TMP19:%.*]] = sub i64 [[TMP18]], [[TMP17]]
-// CHECK_HOSTCALL-NEXT: [[TMP20:%.*]] = add i64 [[TMP19]], 1
+// CHECK_HOSTCALL-NEXT: [[TMP21:%.*]] = ptrtoint ptr [[TMP5]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP22:%.*]] = ptrtoint ptr [[TMP17]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP23:%.*]] = sub i64 [[TMP22]], [[TMP21]]
+// CHECK_HOSTCALL-NEXT: [[TMP24:%.*]] = add i64 [[TMP23]], 1
// CHECK_HOSTCALL-NEXT: br label [[STRLEN_JOIN1]]
// CHECK_HOSTCALL: strlen.join1:
-// CHECK_HOSTCALL-NEXT: [[TMP21:%.*]] = phi i64 [ [[TMP20]], [[STRLEN_WHILE_DONE3]] ], [ 0, [[STRLEN_JOIN]] ]
-// CHECK_HOSTCALL-NEXT: [[TMP22:%.*]] = call i64 @__ockl_printf_append_string_n(i64 [[TMP11]], ptr [[TMP1]], i64 [[TMP21]], i32 0)
-// CHECK_HOSTCALL-NEXT: [[TMP23:%.*]] = zext i32 [[TMP0]] to i64
-// CHECK_HOSTCALL-NEXT: [[TMP24:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP22]], i32 1, i64 [[TMP23]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 1)
-// CHECK_HOSTCALL-NEXT: [[TMP25:%.*]] = trunc i64 [[TMP24]] to i32
-// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[S]]) #[[ATTR3]]
+// CHECK_HOSTCALL-NEXT: [[TMP25:%.*]] = phi i64 [ [[TMP24]], [[STRLEN_WHILE_DONE3]] ], [ 0, [[STRLEN_JOIN]] ]
+// CHECK_HOSTCALL-NEXT: [[TMP26:%.*]] = call i64 @__ockl_printf_append_string_n(i64 [[TMP15]], ptr [[TMP5]], i64 [[TMP25]], i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP27:%.*]] = zext i32 [[TMP4]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP28:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP26]], i32 1, i64 [[TMP27]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 1)
+// CHECK_HOSTCALL-NEXT: [[TMP29:%.*]] = trunc i64 [[TMP28]] to i32
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[S]]) #[[ATTR2]]
// CHECK_HOSTCALL-NEXT: ret void
//
__kernel void test_printf_str_int(int i) {
@@ -240,8 +254,8 @@ __kernel void test_printf_str_int(int i) {
// CHECK_BUFFERED-NEXT: [[VAR:%.*]] = alloca half, align 2, addrspace(5)
// CHECK_BUFFERED-NEXT: [[VAR2:%.*]] = alloca half, align 2, addrspace(5)
// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.start.p5(i64 2, ptr addrspace(5) [[VAR]]) #[[ATTR1]]
-// CHECK_BUFFERED-NEXT: store half 0xH2E66, ptr addrspace(5) [[VAR]], align 2, !tbaa [[TBAA20:![0-9]+]]
-// CHECK_BUFFERED-NEXT: [[TMP0:%.*]] = load half, ptr addrspace(5) [[VAR]], align 2, !tbaa [[TBAA20]]
+// CHECK_BUFFERED-NEXT: store half 0xH2E66, ptr addrspace(5) [[VAR]], align 2, !tbaa [[TBAA21:![0-9]+]]
+// CHECK_BUFFERED-NEXT: [[TMP0:%.*]] = load half, ptr addrspace(5) [[VAR]], align 2, !tbaa [[TBAA21]]
// CHECK_BUFFERED-NEXT: [[CONV:%.*]] = fpext half [[TMP0]] to double
// CHECK_BUFFERED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 20)
// CHECK_BUFFERED-NEXT: [[TMP1:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
@@ -250,8 +264,8 @@ __kernel void test_printf_str_int(int i) {
// CHECK_BUFFERED-NEXT: [[TMP2:%.*]] = xor i1 [[TMP1]], true
// CHECK_BUFFERED-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP2]] to i32
// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.start.p5(i64 2, ptr addrspace(5) [[VAR2]]) #[[ATTR1]]
-// CHECK_BUFFERED-NEXT: store half 0xH3266, ptr addrspace(5) [[VAR2]], align 2, !tbaa [[TBAA22:![0-9]+]]
-// CHECK_BUFFERED-NEXT: [[TMP3:%.*]] = load half, ptr addrspace(5) [[VAR2]], align 2, !tbaa [[TBAA22]]
+// CHECK_BUFFERED-NEXT: store half 0xH3266, ptr addrspace(5) [[VAR2]], align 2, !tbaa [[TBAA23:![0-9]+]]
+// CHECK_BUFFERED-NEXT: [[TMP3:%.*]] = load half, ptr addrspace(5) [[VAR2]], align 2, !tbaa [[TBAA23]]
// CHECK_BUFFERED-NEXT: [[PRINTF_ALLOC_FN1:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 20)
// CHECK_BUFFERED-NEXT: [[TMP4:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN1]], null
// CHECK_BUFFERED-NEXT: br i1 [[TMP4]], label [[ARGPUSH_BLOCK3:%.*]], label [[END_BLOCK2:%.*]]
@@ -283,7 +297,7 @@ __kernel void test_printf_str_int(int i) {
// CHECK_HOSTCALL-NEXT: entry:
// CHECK_HOSTCALL-NEXT: [[VAR:%.*]] = alloca half, align 2, addrspace(5)
// CHECK_HOSTCALL-NEXT: [[VAR2:%.*]] = alloca half, align 2, addrspace(5)
-// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.start.p5(i64 2, ptr addrspace(5) [[VAR]]) #[[ATTR3]]
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.start.p5(i64 2, ptr addrspace(5) [[VAR]]) #[[ATTR2]]
// CHECK_HOSTCALL-NEXT: store half 0xH2E66, ptr addrspace(5) [[VAR]], align 2, !tbaa [[TBAA13:![0-9]+]]
// CHECK_HOSTCALL-NEXT: [[TMP0:%.*]] = load half, ptr addrspace(5) [[VAR]], align 2, !tbaa [[TBAA13]]
// CHECK_HOSTCALL-NEXT: [[CONV:%.*]] = fpext half [[TMP0]] to double
@@ -306,7 +320,7 @@ __kernel void test_printf_str_int(int i) {
// CHECK_HOSTCALL-NEXT: [[TMP11:%.*]] = bitcast double [[CONV]] to i64
// CHECK_HOSTCALL-NEXT: [[TMP12:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP10]], i32 1, i64 [[TMP11]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 1)
// CHECK_HOSTCALL-NEXT: [[TMP13:%.*]] = trunc i64 [[TMP12]] to i32
-// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.start.p5(i64 2, ptr addrspace(5) [[VAR2]]) #[[ATTR3]]
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.start.p5(i64 2, ptr addrspace(5) [[VAR2]]) #[[ATTR2]]
// CHECK_HOSTCALL-NEXT: store half 0xH3266, ptr addrspace(5) [[VAR2]], align 2, !tbaa [[TBAA15:![0-9]+]]
// CHECK_HOSTCALL-NEXT: [[TMP14:%.*]] = load half, ptr addrspace(5) [[VAR2]], align 2, !tbaa [[TBAA15]]
// CHECK_HOSTCALL-NEXT: [[TMP15:%.*]] = call i64 @__ockl_printf_begin(i64 0)
@@ -329,8 +343,8 @@ __kernel void test_printf_str_int(int i) {
// CHECK_HOSTCALL-NEXT: [[TMP26:%.*]] = bitcast double [[TMP25]] to i64
// CHECK_HOSTCALL-NEXT: [[TMP27:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP24]], i32 1, i64 [[TMP26]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 1)
// CHECK_HOSTCALL-NEXT: [[TMP28:%.*]] = trunc i64 [[TMP27]] to i32
-// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.end.p5(i64 2, ptr addrspace(5) [[VAR2]]) #[[ATTR3]]
-// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.end.p5(i64 2, ptr addrspace(5) [[VAR]]) #[[ATTR3]]
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.end.p5(i64 2, ptr addrspace(5) [[VAR2]]) #[[ATTR2]]
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.end.p5(i64 2, ptr addrspace(5) [[VAR]]) #[[ATTR2]]
// CHECK_HOSTCALL-NEXT: ret void
//
__kernel void test_half() {
@@ -358,10 +372,10 @@ typedef __attribute__((ext_vector_type(16))) float float16;
// CHECK_BUFFERED-NEXT: [[VAR5:%.*]] = alloca <16 x float>, align 64, addrspace(5)
// CHECK_BUFFERED-NEXT: [[DOTCOMPOUNDLITERAL18:%.*]] = alloca <16 x float>, align 64, addrspace(5)
// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.start.p5(i64 16, ptr addrspace(5) [[A]]) #[[ATTR1]]
-// CHECK_BUFFERED-NEXT: store <4 x i32> <i32 1, i32 2, i32 3, i32 4>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL]], align 16, !tbaa [[TBAA24:![0-9]+]]
-// CHECK_BUFFERED-NEXT: [[TMP0:%.*]] = load <4 x i32>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL]], align 16, !tbaa [[TBAA24]]
-// CHECK_BUFFERED-NEXT: store <4 x i32> [[TMP0]], ptr addrspace(5) [[A]], align 16, !tbaa [[TBAA24]]
-// CHECK_BUFFERED-NEXT: [[TMP1:%.*]] = load <4 x i32>, ptr addrspace(5) [[A]], align 16, !tbaa [[TBAA24]]
+// CHECK_BUFFERED-NEXT: store <4 x i32> <i32 1, i32 2, i32 3, i32 4>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL]], align 16, !tbaa [[TBAA25:![0-9]+]]
+// CHECK_BUFFERED-NEXT: [[TMP0:%.*]] = load <4 x i32>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL]], align 16, !tbaa [[TBAA25]]
+// CHECK_BUFFERED-NEXT: store <4 x i32> [[TMP0]], ptr addrspace(5) [[A]], align 16, !tbaa [[TBAA25]]
+// CHECK_BUFFERED-NEXT: [[TMP1:%.*]] = load <4 x i32>, ptr addrspace(5) [[A]], align 16, !tbaa [[TBAA25]]
// CHECK_BUFFERED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 44)
// CHECK_BUFFERED-NEXT: [[TMP2:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
// CHECK_BUFFERED-NEXT: br i1 [[TMP2]], label [[ARGPUSH_BLOCK:%.*]], label [[END_BLOCK:%.*]]
@@ -369,16 +383,16 @@ typedef __attribute__((ext_vector_type(16))) float float16;
// CHECK_BUFFERED-NEXT: [[TMP3:%.*]] = xor i1 [[TMP2]], true
// CHECK_BUFFERED-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP3]] to i32
// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.start.p5(i64 16, ptr addrspace(5) [[B]]) #[[ATTR1]]
-// CHECK_BUFFERED-NEXT: [[TMP4:%.*]] = load <4 x i32>, ptr addrspace(5) [[A]], align 16, !tbaa [[TBAA24]]
+// CHECK_BUFFERED-NEXT: [[TMP4:%.*]] = load <4 x i32>, ptr addrspace(5) [[A]], align 16, !tbaa [[TBAA25]]
// CHECK_BUFFERED-NEXT: [[VECEXT:%.*]] = extractelement <4 x i32> [[TMP4]], i32 2
// CHECK_BUFFERED-NEXT: [[VECINIT:%.*]] = insertelement <3 x i32> <i32 1, i32 undef, i32 undef>, i32 [[VECEXT]], i32 1
// CHECK_BUFFERED-NEXT: [[VECINIT5:%.*]] = insertelement <3 x i32> [[VECINIT]], i32 3, i32 2
// CHECK_BUFFERED-NEXT: [[EXTRACTVEC:%.*]] = shufflevector <3 x i32> [[VECINIT5]], <3 x i32> poison, <4 x i32> <i32 0, i32 1, i32 2, i32 poison>
-// CHECK_BUFFERED-NEXT: store <4 x i32> [[EXTRACTVEC]], ptr addrspace(5) [[DOTCOMPOUNDLITERAL4]], align 16, !tbaa [[TBAA24]]
+// CHECK_BUFFERED-NEXT: store <4 x i32> [[EXTRACTVEC]], ptr addrspace(5) [[DOTCOMPOUNDLITERAL4]], align 16, !tbaa [[TBAA25]]
// CHECK_BUFFERED-NEXT: [[LOADVEC4:%.*]] = load <4 x i32>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL4]], align 16
// CHECK_BUFFERED-NEXT: [[EXTRACTVEC6:%.*]] = shufflevector <4 x i32> [[LOADVEC4]], <4 x i32> poison, <3 x i32> <i32 0, i32 1, i32 2>
// CHECK_BUFFERED-NEXT: [[EXTRACTVEC7:%.*]] = shufflevector <3 x i32> [[EXTRACTVEC6]], <3 x i32> poison, <4 x i32> <i32 0, i32 1, i32 2, i32 poison>
-// CHECK_BUFFERED-NEXT: store <4 x i32> [[EXTRACTVEC7]], ptr addrspace(5) [[B]], align 16, !tbaa [[TBAA24]]
+// CHECK_BUFFERED-NEXT: store <4 x i32> [[EXTRACTVEC7]], ptr addrspace(5) [[B]], align 16, !tbaa [[TBAA25]]
// CHECK_BUFFERED-NEXT: [[LOADVEC48:%.*]] = load <4 x i32>, ptr addrspace(5) [[B]], align 16
// CHECK_BUFFERED-NEXT: [[EXTRACTVEC9:%.*]] = shufflevector <4 x i32> [[LOADVEC48]], <4 x i32> poison, <3 x i32> <i32 0, i32 1, i32 2>
// CHECK_BUFFERED-NEXT: [[PRINTF_ALLOC_FN10:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 36)
@@ -410,18 +424,18 @@ typedef __attribute__((ext_vector_type(16))) float float16;
// CHECK_BUFFERED-NEXT: [[TMP16:%.*]] = xor i1 [[TMP5]], true
// CHECK_BUFFERED-NEXT: [[PRINTF_RESULT16:%.*]] = sext i1 [[TMP16]] to i32
// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[VAR3]]) #[[ATTR1]]
-// CHECK_BUFFERED-NEXT: store <2 x float> <float 0x4008F5C280000000, float 0x4003D70A40000000>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL17]], align 8, !tbaa [[TBAA24]]
-// CHECK_BUFFERED-NEXT: [[TMP17:%.*]] = load <2 x float>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL17]], align 8, !tbaa [[TBAA24]]
-// CHECK_BUFFERED-NEXT: store <2 x float> [[TMP17]], ptr addrspace(5) [[VAR3]], align 8, !tbaa [[TBAA24]]
+// CHECK_BUFFERED-NEXT: store <2 x float> <float 0x4008F5C280000000, float 0x4003D70A40000000>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL17]], align 8, !tbaa [[TBAA25]]
+// CHECK_BUFFERED-NEXT: [[TMP17:%.*]] = load <2 x float>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL17]], align 8, !tbaa [[TBAA25]]
+// CHECK_BUFFERED-NEXT: store <2 x float> [[TMP17]], ptr addrspace(5) [[VAR3]], align 8, !tbaa [[TBAA25]]
// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.start.p5(i64 32, ptr addrspace(5) [[VAR4]]) #[[ATTR1]]
-// CHECK_BUFFERED-NEXT: store <8 x float> <float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000>, ptr addrspace(5) [[VAR4]], align 32, !tbaa [[TBAA24]]
+// CHECK_BUFFERED-NEXT: store <8 x float> <float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000>, ptr addrspace(5) [[VAR4]], align 32, !tbaa [[TBAA25]]
// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.start.p5(i64 64, ptr addrspace(5) [[VAR5]]) #[[ATTR1]]
-// CHECK_BUFFERED-NEXT: store <16 x float> <float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL18]], align 64, !tbaa [[TBAA24]]
-// CHECK_BUFFERED-NEXT: [[TMP18:%.*]] = load <16 x float>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL18]], align 64, !tbaa [[TBAA24]]
-// CHECK_BUFFERED-NEXT: store <16 x float> [[TMP18]], ptr addrspace(5) [[VAR5]], align 64, !tbaa [[TBAA24]]
-// CHECK_BUFFERED-NEXT: [[TMP19:%.*]] = load <2 x float>, ptr addrspace(5) [[VAR3]], align 8, !tbaa [[TBAA24]]
-// CHECK_BUFFERED-NEXT: [[TMP20:%.*]] = load <8 x float>, ptr addrspace(5) [[VAR4]], align 32, !tbaa [[TBAA24]]
-// CHECK_BUFFERED-NEXT: [[TMP21:%.*]] = load <16 x float>, ptr addrspace(5) [[VAR5]], align 64, !tbaa [[TBAA24]]
+// CHECK_BUFFERED-NEXT: store <16 x float> <float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL18]], align 64, !tbaa [[TBAA25]]
+// CHECK_BUFFERED-NEXT: [[TMP18:%.*]] = load <16 x float>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL18]], align 64, !tbaa [[TBAA25]]
+// CHECK_BUFFERED-NEXT: store <16 x float> [[TMP18]], ptr addrspace(5) [[VAR5]], align 64, !tbaa [[TBAA25]]
+// CHECK_BUFFERED-NEXT: [[TMP19:%.*]] = load <2 x float>, ptr addrspace(5) [[VAR3]], align 8, !tbaa [[TBAA25]]
+// CHECK_BUFFERED-NEXT: [[TMP20:%.*]] = load <8 x float>, ptr addrspace(5) [[VAR4]], align 32, !tbaa [[TBAA25]]
+// CHECK_BUFFERED-NEXT: [[TMP21:%.*]] = load <16 x float>, ptr addrspace(5) [[VAR5]], align 64, !tbaa [[TBAA25]]
// CHECK_BUFFERED-NEXT: [[PRINTF_ALLOC_FN19:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 220)
// CHECK_BUFFERED-NEXT: [[TMP22:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN19]], null
// CHECK_BUFFERED-NEXT: br i1 [[TMP22]], label [[ARGPUSH_BLOCK21:%.*]], label [[END_BLOCK20:%.*]]
@@ -446,122 +460,140 @@ typedef __attribute__((ext_vector_type(16))) float float16;
// CHECK_BUFFERED: end.block20:
// CHECK_BUFFERED-NEXT: [[TMP31:%.*]] = xor i1 [[TMP22]], true
// CHECK_BUFFERED-NEXT: [[PRINTF_RESULT48:%.*]] = sext i1 [[TMP31]] to i32
-// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.end.p5(i64 64, ptr addrspace(5) [[VAR5]]) #[[ATTR1]]
-// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.end.p5(i64 32, ptr addrspace(5) [[VAR4]]) #[[ATTR1]]
-// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[VAR3]]) #[[ATTR1]]
-// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.end.p5(i64 16, ptr addrspace(5) [[B]]) #[[ATTR1]]
-// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.end.p5(i64 16, ptr addrspace(5) [[A]]) #[[ATTR1]]
-// CHECK_BUFFERED-NEXT: ret void
+// CHECK_BUFFERED-NEXT: [[PRINTF_ALLOC_FN49:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 36)
+// CHECK_BUFFERED-NEXT: [[TMP32:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN49]], null
+// CHECK_BUFFERED-NEXT: br i1 [[TMP32]], label [[ARGPUSH_BLOCK51:%.*]], label [[END_BLOCK50:%.*]]
// CHECK_BUFFERED: argpush.block21:
// CHECK_BUFFERED-NEXT: store i32 882, ptr addrspace(1) [[PRINTF_ALLOC_FN19]], align 4
-// CHECK_BUFFERED-NEXT: [[TMP32:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN19]], i32 4
-// CHECK_BUFFERED-NEXT: store i64 -7312483467425511358, ptr addrspace(1) [[TMP32]], align 8
-// CHECK_BUFFERED-NEXT: [[TMP33:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP32]], i32 8
-// CHECK_BUFFERED-NEXT: [[TMP34:%.*]] = extractelement <2 x float> [[TMP19]], i64 0
-// CHECK_BUFFERED-NEXT: [[TMP35:%.*]] = fpext float [[TMP34]] to double
-// CHECK_BUFFERED-NEXT: [[TMP36:%.*]] = extractelement <2 x float> [[TMP19]], i64 1
-// CHECK_BUFFERED-NEXT: [[TMP37:%.*]] = fpext float [[TMP36]] to double
-// CHECK_BUFFERED-NEXT: store double [[TMP35]], ptr addrspace(1) [[TMP33]], align 8
-// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR22:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP33]], i32 8
-// CHECK_BUFFERED-NEXT: store double [[TMP37]], ptr addrspace(1) [[PRINTBUFFNEXTPTR22]], align 8
+// CHECK_BUFFERED-NEXT: [[TMP33:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN19]], i32 4
+// CHECK_BUFFERED-NEXT: store i64 -7312483467425511358, ptr addrspace(1) [[TMP33]], align 8
+// CHECK_BUFFERED-NEXT: [[TMP34:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP33]], i32 8
+// CHECK_BUFFERED-NEXT: [[TMP35:%.*]] = extractelement <2 x float> [[TMP19]], i64 0
+// CHECK_BUFFERED-NEXT: [[TMP36:%.*]] = fpext float [[TMP35]] to double
+// CHECK_BUFFERED-NEXT: [[TMP37:%.*]] = extractelement <2 x float> [[TMP19]], i64 1
+// CHECK_BUFFERED-NEXT: [[TMP38:%.*]] = fpext float [[TMP37]] to double
+// CHECK_BUFFERED-NEXT: store double [[TMP36]], ptr addrspace(1) [[TMP34]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR22:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP34]], i32 8
+// CHECK_BUFFERED-NEXT: store double [[TMP38]], ptr addrspace(1) [[PRINTBUFFNEXTPTR22]], align 8
// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR23:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR22]], i32 8
-// CHECK_BUFFERED-NEXT: [[TMP38:%.*]] = extractelement <8 x float> [[TMP20]], i64 0
-// CHECK_BUFFERED-NEXT: [[TMP39:%.*]] = fpext float [[TMP38]] to double
-// CHECK_BUFFERED-NEXT: [[TMP40:%.*]] = extractelement <8 x float> [[TMP20]], i64 1
-// CHECK_BUFFERED-NEXT: [[TMP41:%.*]] = fpext float [[TMP40]] to double
-// CHECK_BUFFERED-NEXT: [[TMP42:%.*]] = extractelement <8 x float> [[TMP20]], i64 2
-// CHECK_BUFFERED-NEXT: [[TMP43:%.*]] = fpext float [[TMP42]] to double
-// CHECK_BUFFERED-NEXT: [[TMP44:%.*]] = extractelement <8 x float> [[TMP20]], i64 3
-// CHECK_BUFFERED-NEXT: [[TMP45:%.*]] = fpext float [[TMP44]] to double
-// CHECK_BUFFERED-NEXT: [[TMP46:%.*]] = extractelement <8 x float> [[TMP20]], i64 4
-// CHECK_BUFFERED-NEXT: [[TMP47:%.*]] = fpext float [[TMP46]] to double
-// CHECK_BUFFERED-NEXT: [[TMP48:%.*]] = extractelement <8 x float> [[TMP20]], i64 5
-// CHECK_BUFFERED-NEXT: [[TMP49:%.*]] = fpext float [[TMP48]] to double
-// CHECK_BUFFERED-NEXT: [[TMP50:%.*]] = extractelement <8 x float> [[TMP20]], i64 6
-// CHECK_BUFFERED-NEXT: [[TMP51:%.*]] = fpext float [[TMP50]] to double
-// CHECK_BUFFERED-NEXT: [[TMP52:%.*]] = extractelement <8 x float> [[TMP20]], i64 7
-// CHECK_BUFFERED-NEXT: [[TMP53:%.*]] = fpext float [[TMP52]] to double
-// CHECK_BUFFERED-NEXT: store double [[TMP39]], ptr addrspace(1) [[PRINTBUFFNEXTPTR23]], align 8
+// CHECK_BUFFERED-NEXT: [[TMP39:%.*]] = extractelement <8 x float> [[TMP20]], i64 0
+// CHECK_BUFFERED-NEXT: [[TMP40:%.*]] = fpext float [[TMP39]] to double
+// CHECK_BUFFERED-NEXT: [[TMP41:%.*]] = extractelement <8 x float> [[TMP20]], i64 1
+// CHECK_BUFFERED-NEXT: [[TMP42:%.*]] = fpext float [[TMP41]] to double
+// CHECK_BUFFERED-NEXT: [[TMP43:%.*]] = extractelement <8 x float> [[TMP20]], i64 2
+// CHECK_BUFFERED-NEXT: [[TMP44:%.*]] = fpext float [[TMP43]] to double
+// CHECK_BUFFERED-NEXT: [[TMP45:%.*]] = extractelement <8 x float> [[TMP20]], i64 3
+// CHECK_BUFFERED-NEXT: [[TMP46:%.*]] = fpext float [[TMP45]] to double
+// CHECK_BUFFERED-NEXT: [[TMP47:%.*]] = extractelement <8 x float> [[TMP20]], i64 4
+// CHECK_BUFFERED-NEXT: [[TMP48:%.*]] = fpext float [[TMP47]] to double
+// CHECK_BUFFERED-NEXT: [[TMP49:%.*]] = extractelement <8 x float> [[TMP20]], i64 5
+// CHECK_BUFFERED-NEXT: [[TMP50:%.*]] = fpext float [[TMP49]] to double
+// CHECK_BUFFERED-NEXT: [[TMP51:%.*]] = extractelement <8 x float> [[TMP20]], i64 6
+// CHECK_BUFFERED-NEXT: [[TMP52:%.*]] = fpext float [[TMP51]] to double
+// CHECK_BUFFERED-NEXT: [[TMP53:%.*]] = extractelement <8 x float> [[TMP20]], i64 7
+// CHECK_BUFFERED-NEXT: [[TMP54:%.*]] = fpext float [[TMP53]] to double
+// CHECK_BUFFERED-NEXT: store double [[TMP40]], ptr addrspace(1) [[PRINTBUFFNEXTPTR23]], align 8
// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR24:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR23]], i32 8
-// CHECK_BUFFERED-NEXT: store double [[TMP41]], ptr addrspace(1) [[PRINTBUFFNEXTPTR24]], align 8
+// CHECK_BUFFERED-NEXT: store double [[TMP42]], ptr addrspace(1) [[PRINTBUFFNEXTPTR24]], align 8
// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR25:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR24]], i32 8
-// CHECK_BUFFERED-NEXT: store double [[TMP43]], ptr addrspace(1) [[PRINTBUFFNEXTPTR25]], align 8
+// CHECK_BUFFERED-NEXT: store double [[TMP44]], ptr addrspace(1) [[PRINTBUFFNEXTPTR25]], align 8
// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR26:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR25]], i32 8
-// CHECK_BUFFERED-NEXT: store double [[TMP45]], ptr addrspace(1) [[PRINTBUFFNEXTPTR26]], align 8
+// CHECK_BUFFERED-NEXT: store double [[TMP46]], ptr addrspace(1) [[PRINTBUFFNEXTPTR26]], align 8
// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR27:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR26]], i32 8
-// CHECK_BUFFERED-NEXT: store double [[TMP47]], ptr addrspace(1) [[PRINTBUFFNEXTPTR27]], align 8
+// CHECK_BUFFERED-NEXT: store double [[TMP48]], ptr addrspace(1) [[PRINTBUFFNEXTPTR27]], align 8
// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR28:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR27]], i32 8
-// CHECK_BUFFERED-NEXT: store double [[TMP49]], ptr addrspace(1) [[PRINTBUFFNEXTPTR28]], align 8
+// CHECK_BUFFERED-NEXT: store double [[TMP50]], ptr addrspace(1) [[PRINTBUFFNEXTPTR28]], align 8
// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR29:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR28]], i32 8
-// CHECK_BUFFERED-NEXT: store double [[TMP51]], ptr addrspace(1) [[PRINTBUFFNEXTPTR29]], align 8
+// CHECK_BUFFERED-NEXT: store double [[TMP52]], ptr addrspace(1) [[PRINTBUFFNEXTPTR29]], align 8
// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR30:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR29]], i32 8
-// CHECK_BUFFERED-NEXT: store double [[TMP53]], ptr addrspace(1) [[PRINTBUFFNEXTPTR30]], align 8
+// CHECK_BUFFERED-NEXT: store double [[TMP54]], ptr addrspace(1) [[PRINTBUFFNEXTPTR30]], align 8
// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR31:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR30]], i32 8
-// CHECK_BUFFERED-NEXT: [[TMP54:%.*]] = extractelement <16 x float> [[TMP21]], i64 0
-// CHECK_BUFFERED-NEXT: [[TMP55:%.*]] = fpext float [[TMP54]] to double
-// CHECK_BUFFERED-NEXT: [[TMP56:%.*]] = extractelement <16 x float> [[TMP21]], i64 1
-// CHECK_BUFFERED-NEXT: [[TMP57:%.*]] = fpext float [[TMP56]] to double
-// CHECK_BUFFERED-NEXT: [[TMP58:%.*]] = extractelement <16 x float> [[TMP21]], i64 2
-// CHECK_BUFFERED-NEXT: [[TMP59:%.*]] = fpext float [[TMP58]] to double
-// CHECK_BUFFERED-NEXT: [[TMP60:%.*]] = extractelement <16 x float> [[TMP21]], i64 3
-// CHECK_BUFFERED-NEXT: [[TMP61:%.*]] = fpext float [[TMP60]] to double
-// CHECK_BUFFERED-NEXT: [[TMP62:%.*]] = extractelement <16 x float> [[TMP21]], i64 4
-// CHECK_BUFFERED-NEXT: [[TMP63:%.*]] = fpext float [[TMP62]] to double
-// CHECK_BUFFERED-NEXT: [[TMP64:%.*]] = extractelement <16 x float> [[TMP21]], i64 5
-// CHECK_BUFFERED-NEXT: [[TMP65:%.*]] = fpext float [[TMP64]] to double
-// CHECK_BUFFERED-NEXT: [[TMP66:%.*]] = extractelement <16 x float> [[TMP21]], i64 6
-// CHECK_BUFFERED-NEXT: [[TMP67:%.*]] = fpext float [[TMP66]] to double
-// CHECK_BUFFERED-NEXT: [[TMP68:%.*]] = extractelement <16 x float> [[TMP21]], i64 7
-// CHECK_BUFFERED-NEXT: [[TMP69:%.*]] = fpext float [[TMP68]] to double
-// CHECK_BUFFERED-NEXT: [[TMP70:%.*]] = extractelement <16 x float> [[TMP21]], i64 8
-// CHECK_BUFFERED-NEXT: [[TMP71:%.*]] = fpext float [[TMP70]] to double
-// CHECK_BUFFERED-NEXT: [[TMP72:%.*]] = extractelement <16 x float> [[TMP21]], i64 9
-// CHECK_BUFFERED-NEXT: [[TMP73:%.*]] = fpext float [[TMP72]] to double
-// CHECK_BUFFERED-NEXT: [[TMP74:%.*]] = extractelement <16 x float> [[TMP21]], i64 10
-// CHECK_BUFFERED-NEXT: [[TMP75:%.*]] = fpext float [[TMP74]] to double
-// CHECK_BUFFERED-NEXT: [[TMP76:%.*]] = extractelement <16 x float> [[TMP21]], i64 11
-// CHECK_BUFFERED-NEXT: [[TMP77:%.*]] = fpext float [[TMP76]] to double
-// CHECK_BUFFERED-NEXT: [[TMP78:%.*]] = extractelement <16 x float> [[TMP21]], i64 12
-// CHECK_BUFFERED-NEXT: [[TMP79:%.*]] = fpext float [[TMP78]] to double
-// CHECK_BUFFERED-NEXT: [[TMP80:%.*]] = extractelement <16 x float> [[TMP21]], i64 13
-// CHECK_BUFFERED-NEXT: [[TMP81:%.*]] = fpext float [[TMP80]] to double
-// CHECK_BUFFERED-NEXT: [[TMP82:%.*]] = extractelement <16 x float> [[TMP21]], i64 14
-// CHECK_BUFFERED-NEXT: [[TMP83:%.*]] = fpext float [[TMP82]] to double
-// CHECK_BUFFERED-NEXT: [[TMP84:%.*]] = extractelement <16 x float> [[TMP21]], i64 15
-// CHECK_BUFFERED-NEXT: [[TMP85:%.*]] = fpext float [[TMP84]] to double
-// CHECK_BUFFERED-NEXT: store double [[TMP55]], ptr addrspace(1) [[PRINTBUFFNEXTPTR31]], align 8
+// CHECK_BUFFERED-NEXT: [[TMP55:%.*]] = extractelement <16 x float> [[TMP21]], i64 0
+// CHECK_BUFFERED-NEXT: [[TMP56:%.*]] = fpext float [[TMP55]] to double
+// CHECK_BUFFERED-NEXT: [[TMP57:%.*]] = extractelement <16 x float> [[TMP21]], i64 1
+// CHECK_BUFFERED-NEXT: [[TMP58:%.*]] = fpext float [[TMP57]] to double
+// CHECK_BUFFERED-NEXT: [[TMP59:%.*]] = extractelement <16 x float> [[TMP21]], i64 2
+// CHECK_BUFFERED-NEXT: [[TMP60:%.*]] = fpext float [[TMP59]] to double
+// CHECK_BUFFERED-NEXT: [[TMP61:%.*]] = extractelement <16 x float> [[TMP21]], i64 3
+// CHECK_BUFFERED-NEXT: [[TMP62:%.*]] = fpext float [[TMP61]] to double
+// CHECK_BUFFERED-NEXT: [[TMP63:%.*]] = extractelement <16 x float> [[TMP21]], i64 4
+// CHECK_BUFFERED-NEXT: [[TMP64:%.*]] = fpext float [[TMP63]] to double
+// CHECK_BUFFERED-NEXT: [[TMP65:%.*]] = extractelement <16 x float> [[TMP21]], i64 5
+// CHECK_BUFFERED-NEXT: [[TMP66:%.*]] = fpext float [[TMP65]] to double
+// CHECK_BUFFERED-NEXT: [[TMP67:%.*]] = extractelement <16 x float> [[TMP21]], i64 6
+// CHECK_BUFFERED-NEXT: [[TMP68:%.*]] = fpext float [[TMP67]] to double
+// CHECK_BUFFERED-NEXT: [[TMP69:%.*]] = extractelement <16 x float> [[TMP21]], i64 7
+// CHECK_BUFFERED-NEXT: [[TMP70:%.*]] = fpext float [[TMP69]] to double
+// CHECK_BUFFERED-NEXT: [[TMP71:%.*]] = extractelement <16 x float> [[TMP21]], i64 8
+// CHECK_BUFFERED-NEXT: [[TMP72:%.*]] = fpext float [[TMP71]] to double
+// CHECK_BUFFERED-NEXT: [[TMP73:%.*]] = extractelement <16 x float> [[TMP21]], i64 9
+// CHECK_BUFFERED-NEXT: [[TMP74:%.*]] = fpext float [[TMP73]] to double
+// CHECK_BUFFERED-NEXT: [[TMP75:%.*]] = extractelement <16 x float> [[TMP21]], i64 10
+// CHECK_BUFFERED-NEXT: [[TMP76:%.*]] = fpext float [[TMP75]] to double
+// CHECK_BUFFERED-NEXT: [[TMP77:%.*]] = extractelement <16 x float> [[TMP21]], i64 11
+// CHECK_BUFFERED-NEXT: [[TMP78:%.*]] = fpext float [[TMP77]] to double
+// CHECK_BUFFERED-NEXT: [[TMP79:%.*]] = extractelement <16 x float> [[TMP21]], i64 12
+// CHECK_BUFFERED-NEXT: [[TMP80:%.*]] = fpext float [[TMP79]] to double
+// CHECK_BUFFERED-NEXT: [[TMP81:%.*]] = extractelement <16 x float> [[TMP21]], i64 13
+// CHECK_BUFFERED-NEXT: [[TMP82:%.*]] = fpext float [[TMP81]] to double
+// CHECK_BUFFERED-NEXT: [[TMP83:%.*]] = extractelement <16 x float> [[TMP21]], i64 14
+// CHECK_BUFFERED-NEXT: [[TMP84:%.*]] = fpext float [[TMP83]] to double
+// CHECK_BUFFERED-NEXT: [[TMP85:%.*]] = extractelement <16 x float> [[TMP21]], i64 15
+// CHECK_BUFFERED-NEXT: [[TMP86:%.*]] = fpext float [[TMP85]] to double
+// CHECK_BUFFERED-NEXT: store double [[TMP56]], ptr addrspace(1) [[PRINTBUFFNEXTPTR31]], align 8
// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR32:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR31]], i32 8
-// CHECK_BUFFERED-NEXT: store double [[TMP57]], ptr addrspace(1) [[PRINTBUFFNEXTPTR32]], align 8
+// CHECK_BUFFERED-NEXT: store double [[TMP58]], ptr addrspace(1) [[PRINTBUFFNEXTPTR32]], align 8
// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR33:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR32]], i32 8
-// CHECK_BUFFERED-NEXT: store double [[TMP59]], ptr addrspace(1) [[PRINTBUFFNEXTPTR33]], align 8
+// CHECK_BUFFERED-NEXT: store double [[TMP60]], ptr addrspace(1) [[PRINTBUFFNEXTPTR33]], align 8
// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR34:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR33]], i32 8
-// CHECK_BUFFERED-NEXT: store double [[TMP61]], ptr addrspace(1) [[PRINTBUFFNEXTPTR34]], align 8
+// CHECK_BUFFERED-NEXT: store double [[TMP62]], ptr addrspace(1) [[PRINTBUFFNEXTPTR34]], align 8
// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR35:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR34]], i32 8
-// CHECK_BUFFERED-NEXT: store double [[TMP63]], ptr addrspace(1) [[PRINTBUFFNEXTPTR35]], align 8
+// CHECK_BUFFERED-NEXT: store double [[TMP64]], ptr addrspace(1) [[PRINTBUFFNEXTPTR35]], align 8
// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR36:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR35]], i32 8
-// CHECK_BUFFERED-NEXT: store double [[TMP65]], ptr addrspace(1) [[PRINTBUFFNEXTPTR36]], align 8
+// CHECK_BUFFERED-NEXT: store double [[TMP66]], ptr addrspace(1) [[PRINTBUFFNEXTPTR36]], align 8
// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR37:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR36]], i32 8
-// CHECK_BUFFERED-NEXT: store double [[TMP67]], ptr addrspace(1) [[PRINTBUFFNEXTPTR37]], align 8
+// CHECK_BUFFERED-NEXT: store double [[TMP68]], ptr addrspace(1) [[PRINTBUFFNEXTPTR37]], align 8
// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR38:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR37]], i32 8
-// CHECK_BUFFERED-NEXT: store double [[TMP69]], ptr addrspace(1) [[PRINTBUFFNEXTPTR38]], align 8
+// CHECK_BUFFERED-NEXT: store double [[TMP70]], ptr addrspace(1) [[PRINTBUFFNEXTPTR38]], align 8
// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR39:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR38]], i32 8
-// CHECK_BUFFERED-NEXT: store double [[TMP71]], ptr addrspace(1) [[PRINTBUFFNEXTPTR39]], align 8
+// CHECK_BUFFERED-NEXT: store double [[TMP72]], ptr addrspace(1) [[PRINTBUFFNEXTPTR39]], align 8
// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR40:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR39]], i32 8
-// CHECK_BUFFERED-NEXT: store double [[TMP73]], ptr addrspace(1) [[PRINTBUFFNEXTPTR40]], align 8
+// CHECK_BUFFERED-NEXT: store double [[TMP74]], ptr addrspace(1) [[PRINTBUFFNEXTPTR40]], align 8
// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR41:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR40]], i32 8
-// CHECK_BUFFERED-NEXT: store double [[TMP75]], ptr addrspace(1) [[PRINTBUFFNEXTPTR41]], align 8
+// CHECK_BUFFERED-NEXT: store double [[TMP76]], ptr addrspace(1) [[PRINTBUFFNEXTPTR41]], align 8
// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR42:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR41]], i32 8
-// CHECK_BUFFERED-NEXT: store double [[TMP77]], ptr addrspace(1) [[PRINTBUFFNEXTPTR42]], align 8
+// CHECK_BUFFERED-NEXT: store double [[TMP78]], ptr addrspace(1) [[PRINTBUFFNEXTPTR42]], align 8
// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR43:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR42]], i32 8
-// CHECK_BUFFERED-NEXT: store double [[TMP79]], ptr addrspace(1) [[PRINTBUFFNEXTPTR43]], align 8
+// CHECK_BUFFERED-NEXT: store double [[TMP80]], ptr addrspace(1) [[PRINTBUFFNEXTPTR43]], align 8
// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR44:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR43]], i32 8
-// CHECK_BUFFERED-NEXT: store double [[TMP81]], ptr addrspace(1) [[PRINTBUFFNEXTPTR44]], align 8
+// CHECK_BUFFERED-NEXT: store double [[TMP82]], ptr addrspace(1) [[PRINTBUFFNEXTPTR44]], align 8
// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR45:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR44]], i32 8
-// CHECK_BUFFERED-NEXT: store double [[TMP83]], ptr addrspace(1) [[PRINTBUFFNEXTPTR45]], align 8
+// CHECK_BUFFERED-NEXT: store double [[TMP84]], ptr addrspace(1) [[PRINTBUFFNEXTPTR45]], align 8
// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR46:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR45]], i32 8
-// CHECK_BUFFERED-NEXT: store double [[TMP85]], ptr addrspace(1) [[PRINTBUFFNEXTPTR46]], align 8
+// CHECK_BUFFERED-NEXT: store double [[TMP86]], ptr addrspace(1) [[PRINTBUFFNEXTPTR46]], align 8
// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR47:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR46]], i32 8
// CHECK_BUFFERED-NEXT: br label [[END_BLOCK20]]
+// CHECK_BUFFERED: end.block50:
+// CHECK_BUFFERED-NEXT: [[TMP87:%.*]] = xor i1 [[TMP32]], true
+// CHECK_BUFFERED-NEXT: [[PRINTF_RESULT55:%.*]] = sext i1 [[TMP87]] to i32
+// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.end.p5(i64 64, ptr addrspace(5) [[VAR5]]) #[[ATTR1]]
+// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.end.p5(i64 32, ptr addrspace(5) [[VAR4]]) #[[ATTR1]]
+// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[VAR3]]) #[[ATTR1]]
+// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.end.p5(i64 16, ptr addrspace(5) [[B]]) #[[ATTR1]]
+// CHECK_BUFFERED-NEXT: call void @llvm.lifetime.end.p5(i64 16, ptr addrspace(5) [[A]]) #[[ATTR1]]
+// CHECK_BUFFERED-NEXT: ret void
+// CHECK_BUFFERED: argpush.block51:
+// CHECK_BUFFERED-NEXT: store i32 146, ptr addrspace(1) [[PRINTF_ALLOC_FN49]], align 4
+// CHECK_BUFFERED-NEXT: [[TMP88:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN49]], i32 4
+// CHECK_BUFFERED-NEXT: store i64 7812276053070244211, ptr addrspace(1) [[TMP88]], align 8
+// CHECK_BUFFERED-NEXT: [[TMP89:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP88]], i32 8
+// CHECK_BUFFERED-NEXT: store i64 1, ptr addrspace(1) [[TMP89]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR52:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP89]], i32 8
+// CHECK_BUFFERED-NEXT: store double 0x4008F5C280000000, ptr addrspace(1) [[PRINTBUFFNEXTPTR52]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR53:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR52]], i32 8
+// CHECK_BUFFERED-NEXT: store i64 2, ptr addrspace(1) [[PRINTBUFFNEXTPTR53]], align 8
+// CHECK_BUFFERED-NEXT: [[PRINTBUFFNEXTPTR54:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTBUFFNEXTPTR53]], i32 8
+// CHECK_BUFFERED-NEXT: br label [[END_BLOCK50]]
//
// CHECK_HOSTCALL-LABEL: @test_vector(
// CHECK_HOSTCALL-NEXT: entry:
@@ -574,7 +606,7 @@ typedef __attribute__((ext_vector_type(16))) float float16;
// CHECK_HOSTCALL-NEXT: [[VAR4:%.*]] = alloca <8 x float>, align 32, addrspace(5)
// CHECK_HOSTCALL-NEXT: [[VAR5:%.*]] = alloca <16 x float>, align 64, addrspace(5)
// CHECK_HOSTCALL-NEXT: [[DOTCOMPOUNDLITERAL11:%.*]] = alloca <16 x float>, align 64, addrspace(5)
-// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.start.p5(i64 16, ptr addrspace(5) [[A]]) #[[ATTR3]]
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.start.p5(i64 16, ptr addrspace(5) [[A]]) #[[ATTR2]]
// CHECK_HOSTCALL-NEXT: store <4 x i32> <i32 1, i32 2, i32 3, i32 4>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL]], align 16, !tbaa [[TBAA17:![0-9]+]]
// CHECK_HOSTCALL-NEXT: [[TMP0:%.*]] = load <4 x i32>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL]], align 16, !tbaa [[TBAA17]]
// CHECK_HOSTCALL-NEXT: store <4 x i32> [[TMP0]], ptr addrspace(5) [[A]], align 16, !tbaa [[TBAA17]]
@@ -608,7 +640,7 @@ typedef __attribute__((ext_vector_type(16))) float float16;
// CHECK_HOSTCALL-NEXT: [[TMP22:%.*]] = zext i32 [[TMP21]] to i64
// CHECK_HOSTCALL-NEXT: [[TMP23:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP20]], i32 1, i64 [[TMP22]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 1)
// CHECK_HOSTCALL-NEXT: [[TMP24:%.*]] = trunc i64 [[TMP23]] to i32
-// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.start.p5(i64 16, ptr addrspace(5) [[B]]) #[[ATTR3]]
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.start.p5(i64 16, ptr addrspace(5) [[B]]) #[[ATTR2]]
// CHECK_HOSTCALL-NEXT: [[TMP25:%.*]] = load <4 x i32>, ptr addrspace(5) [[A]], align 16, !tbaa [[TBAA17]]
// CHECK_HOSTCALL-NEXT: [[VECEXT:%.*]] = extractelement <4 x i32> [[TMP25]], i32 2
// CHECK_HOSTCALL-NEXT: [[VECINIT:%.*]] = insertelement <3 x i32> <i32 1, i32 undef, i32 undef>, i32 [[VECEXT]], i32 1
@@ -647,13 +679,13 @@ typedef __attribute__((ext_vector_type(16))) float float16;
// CHECK_HOSTCALL-NEXT: [[TMP43:%.*]] = zext i32 [[TMP42]] to i64
// CHECK_HOSTCALL-NEXT: [[TMP44:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP41]], i32 1, i64 [[TMP43]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 1)
// CHECK_HOSTCALL-NEXT: [[TMP45:%.*]] = trunc i64 [[TMP44]] to i32
-// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[VAR3]]) #[[ATTR3]]
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[VAR3]]) #[[ATTR2]]
// CHECK_HOSTCALL-NEXT: store <2 x float> <float 0x4008F5C280000000, float 0x4003D70A40000000>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL10]], align 8, !tbaa [[TBAA17]]
// CHECK_HOSTCALL-NEXT: [[TMP46:%.*]] = load <2 x float>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL10]], align 8, !tbaa [[TBAA17]]
// CHECK_HOSTCALL-NEXT: store <2 x float> [[TMP46]], ptr addrspace(5) [[VAR3]], align 8, !tbaa [[TBAA17]]
-// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.start.p5(i64 32, ptr addrspace(5) [[VAR4]]) #[[ATTR3]]
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.start.p5(i64 32, ptr addrspace(5) [[VAR4]]) #[[ATTR2]]
// CHECK_HOSTCALL-NEXT: store <8 x float> <float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000, float 0x40091EB860000000>, ptr addrspace(5) [[VAR4]], align 32, !tbaa [[TBAA17]]
-// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.start.p5(i64 64, ptr addrspace(5) [[VAR5]]) #[[ATTR3]]
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.start.p5(i64 64, ptr addrspace(5) [[VAR5]]) #[[ATTR2]]
// CHECK_HOSTCALL-NEXT: store <16 x float> <float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x4008F5C280000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000, float 0x40028F5C20000000>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL11]], align 64, !tbaa [[TBAA17]]
// CHECK_HOSTCALL-NEXT: [[TMP47:%.*]] = load <16 x float>, ptr addrspace(5) [[DOTCOMPOUNDLITERAL11]], align 64, !tbaa [[TBAA17]]
// CHECK_HOSTCALL-NEXT: store <16 x float> [[TMP47]], ptr addrspace(5) [[VAR5]], align 64, !tbaa [[TBAA17]]
@@ -781,11 +813,31 @@ typedef __attribute__((ext_vector_type(16))) float float16;
// CHECK_HOSTCALL-NEXT: [[TMP163:%.*]] = bitcast double [[TMP162]] to i64
// CHECK_HOSTCALL-NEXT: [[TMP164:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP160]], i32 1, i64 [[TMP163]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 1)
// CHECK_HOSTCALL-NEXT: [[TMP165:%.*]] = trunc i64 [[TMP164]] to i32
-// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.end.p5(i64 64, ptr addrspace(5) [[VAR5]]) #[[ATTR3]]
-// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.end.p5(i64 32, ptr addrspace(5) [[VAR4]]) #[[ATTR3]]
-// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[VAR3]]) #[[ATTR3]]
-// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.end.p5(i64 16, ptr addrspace(5) [[B]]) #[[ATTR3]]
-// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.end.p5(i64 16, ptr addrspace(5) [[A]]) #[[ATTR3]]
+// CHECK_HOSTCALL-NEXT: [[TMP166:%.*]] = call i64 @__ockl_printf_begin(i64 0)
+// CHECK_HOSTCALL-NEXT: br i1 icmp eq (ptr addrspacecast (ptr addrspace(4) @.str.7 to ptr), ptr null), label [[STRLEN_JOIN15:%.*]], label [[STRLEN_WHILE16:%.*]]
+// CHECK_HOSTCALL: strlen.while16:
+// CHECK_HOSTCALL-NEXT: [[TMP167:%.*]] = phi ptr [ addrspacecast (ptr addrspace(4) @.str.7 to ptr), [[STRLEN_JOIN12]] ], [ [[TMP168:%.*]], [[STRLEN_WHILE16]] ]
+// CHECK_HOSTCALL-NEXT: [[TMP168]] = getelementptr i8, ptr [[TMP167]], i64 1
+// CHECK_HOSTCALL-NEXT: [[TMP169:%.*]] = load i8, ptr [[TMP167]], align 1
+// CHECK_HOSTCALL-NEXT: [[TMP170:%.*]] = icmp eq i8 [[TMP169]], 0
+// CHECK_HOSTCALL-NEXT: br i1 [[TMP170]], label [[STRLEN_WHILE_DONE17:%.*]], label [[STRLEN_WHILE16]]
+// CHECK_HOSTCALL: strlen.while.done17:
+// CHECK_HOSTCALL-NEXT: [[TMP171:%.*]] = ptrtoint ptr [[TMP167]] to i64
+// CHECK_HOSTCALL-NEXT: [[TMP172:%.*]] = sub i64 [[TMP171]], ptrtoint (ptr addrspacecast (ptr addrspace(4) @.str.7 to ptr) to i64)
+// CHECK_HOSTCALL-NEXT: [[TMP173:%.*]] = add i64 [[TMP172]], 1
+// CHECK_HOSTCALL-NEXT: br label [[STRLEN_JOIN15]]
+// CHECK_HOSTCALL: strlen.join15:
+// CHECK_HOSTCALL-NEXT: [[TMP174:%.*]] = phi i64 [ [[TMP173]], [[STRLEN_WHILE_DONE17]] ], [ 0, [[STRLEN_JOIN12]] ]
+// CHECK_HOSTCALL-NEXT: [[TMP175:%.*]] = call i64 @__ockl_printf_append_string_n(i64 [[TMP166]], ptr addrspacecast (ptr addrspace(4) @.str.7 to ptr), i64 [[TMP174]], i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP176:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP175]], i32 1, i64 1, i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP177:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP176]], i32 1, i64 4614208033961017344, i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// CHECK_HOSTCALL-NEXT: [[TMP178:%.*]] = call i64 @__ockl_printf_append_args(i64 [[TMP177]], i32 1, i64 2, i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 1)
+// CHECK_HOSTCALL-NEXT: [[TMP179:%.*]] = trunc i64 [[TMP178]] to i32
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.end.p5(i64 64, ptr addrspace(5) [[VAR5]]) #[[ATTR2]]
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.end.p5(i64 32, ptr addrspace(5) [[VAR4]]) #[[ATTR2]]
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[VAR3]]) #[[ATTR2]]
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.end.p5(i64 16, ptr addrspace(5) [[B]]) #[[ATTR2]]
+// CHECK_HOSTCALL-NEXT: call void @llvm.lifetime.end.p5(i64 16, ptr addrspace(5) [[A]]) #[[ATTR2]]
// CHECK_HOSTCALL-NEXT: ret void
//
__kernel void test_vector() {
@@ -802,4 +854,8 @@ float16 var5 = (float16)((float8)(3.12), (float8)(2.32));
printf("%v2hlf, %v8hlf, %v16hlf", var3, var4, var5);
+// cases where vector specifier is not part of conversion
+// specifier
+printf("%dv, v8%f, %dv16", 1, 3.12f, 2);
+
}
diff --git a/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp b/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp
index 86af790f5fd76..215dfb6f07464 100644
--- a/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp
+++ b/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp
@@ -46,6 +46,9 @@ static Value *fitArgInto64Bits(IRBuilder<> &Builder, Value *Arg,
return Builder.CreateBitCast(Arg, Int64Ty);
}
+ // The cast is necessary for the hostcall case
+ // for the argument to be compatible with device lib
+ // functions.
if (!IsBuffered && isa<PointerType>(Ty)) {
return Builder.CreatePtrToInt(Arg, Int64Ty);
}
@@ -228,7 +231,7 @@ static void locateCStringsAndVectors(SparseBitVector<8> &BV,
return;
auto Spec = Str.slice(SpecPos, SpecEnd + 1);
- if ((Spec.find_first_of("v")) != StringRef::npos)
+ if ((Spec.find("v")) != StringRef::npos)
OV.set(ArgIdx);
ArgIdx += Spec.count('*');
More information about the cfe-commits
mailing list