[clang] [llvm] [AMDGPU] Fix out-of-bounds read of a trailing '%' in printf formats (PR #229088)
Steffen Larsen via cfe-commits
cfe-commits at lists.llvm.org
Mon Oct 5 23:52:03 PDT 2026
https://github.com/steffenlarsen updated https://github.com/llvm/llvm-project/pull/229088
>From 1eef8cbf52f6ab04b272ed5ea0a44e231cc4a31b Mon Sep 17 00:00:00 2001
From: Steffen Holst Larsen <sholstla at amd.com>
Date: Mon, 5 Oct 2026 08:47:02 -0500
Subject: [PATCH 1/2] [AMDGPU] Fix out-of-bounds read of a trailing '%' in
printf formats
locateCStrings checked whether the character after a '%' was another '%'
without checking that there was one. A device printf whose format string
ends in '%' hit the StringRef index assertion in both hostcall and
buffered mode.
Assisted-by: Claude Opus 5.5
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
---
.../CodeGenHIP/printf-trailing-percent.hip | 30 +++++++++++++++++++
.../lib/Transforms/Utils/AMDGPUEmitPrintf.cpp | 2 +-
2 files changed, 31 insertions(+), 1 deletion(-)
create mode 100644 clang/test/CodeGenHIP/printf-trailing-percent.hip
diff --git a/clang/test/CodeGenHIP/printf-trailing-percent.hip b/clang/test/CodeGenHIP/printf-trailing-percent.hip
new file mode 100644
index 0000000000000..1794b0f48ecaf
--- /dev/null
+++ b/clang/test/CodeGenHIP/printf-trailing-percent.hip
@@ -0,0 +1,30 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -emit-llvm -fcuda-is-device \
+// RUN: -Wno-format -mprintf-kind=hostcall -o - %s \
+// RUN: | FileCheck --check-prefix=HOSTCALL %s
+// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -emit-llvm -fcuda-is-device \
+// RUN: -Wno-format -mprintf-kind=buffered -o - %s \
+// RUN: | FileCheck --check-prefix=BUFFERED %s
+
+// A format string ending in an incomplete specifier must not be read past its
+// end, and the specifiers before it must still be found.
+
+#define __device__ __attribute__((device))
+
+extern "C" __device__ int printf(const char *format, ...);
+
+__device__ int trailing_percent(const char *s) { return printf("%s%", s); }
+
+// The %s argument is printed as a string, not as a pointer.
+
+// HOSTCALL-LABEL: define {{.*}} @_Z16trailing_percentPKc(
+// HOSTCALL: call i64 @__ockl_printf_begin(i64 0)
+// HOSTCALL: call i64 @__ockl_printf_append_string_n(i64 %{{.*}}, ptr {{.*}}@.str{{.*}}, i64 %{{.*}}, i32 0)
+// HOSTCALL: call i64 @__ockl_printf_append_string_n(i64 %{{.*}}, ptr %{{.*}}, i64 %{{.*}}, i32 1)
+// HOSTCALL-NOT: call i64 @__ockl_printf_append_args
+
+// BUFFERED-LABEL: define {{.*}} @_Z16trailing_percentPKc(
+// BUFFERED: call ptr addrspace(1) @__printf_alloc(
+// BUFFERED: call void @llvm.memcpy.p1.p0.i64(
+// BUFFERED-NOT: ptrtoint
+// BUFFERED: !{!"0:0:{{[0-9a-f]+}},%s%"}
diff --git a/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp b/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp
index 8fa120f94e1df..3afb7af38fce2 100644
--- a/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp
+++ b/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp
@@ -187,7 +187,7 @@ static void locateCStrings(SparseBitVector<8> &BV, StringRef Str) {
unsigned ArgIdx = 1;
while ((SpecPos = Str.find_first_of('%', SpecPos)) != StringRef::npos) {
- if (Str[SpecPos + 1] == '%') {
+ if (SpecPos + 1 < Str.size() && Str[SpecPos + 1] == '%') {
SpecPos += 2;
continue;
}
>From ecb1c029ce808dd7b93c4593a2e3d4714c4c55ca Mon Sep 17 00:00:00 2001
From: Steffen Holst Larsen <sholstla at amd.com>
Date: Tue, 6 Oct 2026 01:49:25 -0500
Subject: [PATCH 2/2] Move test to existing file and auto-gen
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
---
clang/test/CodeGenHIP/printf-builtin.hip | 169 ++++++++++++++++--
.../CodeGenHIP/printf-trailing-percent.hip | 30 ----
2 files changed, 153 insertions(+), 46 deletions(-)
delete mode 100644 clang/test/CodeGenHIP/printf-trailing-percent.hip
diff --git a/clang/test/CodeGenHIP/printf-builtin.hip b/clang/test/CodeGenHIP/printf-builtin.hip
index 5a95eb4862feb..0bd72ef1f5b57 100644
--- a/clang/test/CodeGenHIP/printf-builtin.hip
+++ b/clang/test/CodeGenHIP/printf-builtin.hip
@@ -1,31 +1,168 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
// REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -emit-llvm -disable-llvm-optzns -mprintf-kind=hostcall -fno-builtin-printf -fcuda-is-device \
-// RUN: -o - %s | FileCheck --check-prefixes=CHECK,HOSTCALL %s
+// RUN: %clang_cc1 -triple amdgpu9-amd-amdhsa -emit-llvm -disable-llvm-optzns -mprintf-kind=hostcall -fno-builtin-printf -fcuda-is-device \
+// RUN: -o - %s | FileCheck --check-prefixes=HOSTCALL %s
// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -emit-llvm -disable-llvm-optzns -mprintf-kind=hostcall -fno-builtin-printf -fcuda-is-device \
-// RUN: -o - %s | FileCheck --check-prefixes=CHECK-AMDGCNSPIRV,HOSTCALL-AMDGCNSPIRV %s
-// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -emit-llvm -disable-llvm-optzns -mprintf-kind=buffered -fno-builtin-printf -fcuda-is-device \
-// RUN: -o - %s | FileCheck --check-prefixes=CHECK,BUFFERED %s
+// RUN: -o - %s | FileCheck --check-prefixes=HOSTCALL-AMDGCNSPIRV %s
+// RUN: %clang_cc1 -triple amdgpu9-amd-amdhsa -emit-llvm -disable-llvm-optzns -mprintf-kind=buffered -fno-builtin-printf -fcuda-is-device \
+// RUN: -o - %s | FileCheck --check-prefixes=BUFFERED %s
// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -emit-llvm -disable-llvm-optzns -mprintf-kind=buffered -fno-builtin-printf -fcuda-is-device \
-// RUN: -o - %s | FileCheck --check-prefixes=CHECK-AMDGCNSPIRV,BUFFERED-AMDGCNSPIRV %s
+// RUN: -o - %s | FileCheck --check-prefixes=BUFFERED-AMDGCNSPIRV %s
#define __device__ __attribute__((device))
extern "C" __device__ int printf(const char *format, ...);
-// CHECK-LABEL: @_Z4foo1v()
+// HOSTCALL-LABEL: define dso_local noundef i32 @_Z4foo1v(
+// HOSTCALL-SAME: ) #[[ATTR0:[0-9]+]] {
+// HOSTCALL-NEXT: [[ENTRY:.*]]:
+// HOSTCALL-NEXT: [[TMP0:%.*]] = call i64 @__ockl_printf_begin(i64 0)
+// HOSTCALL-NEXT: [[TMP1:%.*]] = icmp eq ptr addrspacecast (ptr addrspace(4) @.str to ptr), null
+// HOSTCALL-NEXT: br i1 [[TMP1]], label %[[STRLEN_JOIN:.*]], label %[[STRLEN_WHILE:.*]]
+// HOSTCALL: [[STRLEN_WHILE]]:
+// HOSTCALL-NEXT: [[TMP2:%.*]] = phi ptr [ addrspacecast (ptr addrspace(4) @.str to ptr), %[[ENTRY]] ], [ [[TMP3:%.*]], %[[STRLEN_WHILE]] ]
+// HOSTCALL-NEXT: [[TMP3]] = getelementptr i8, ptr [[TMP2]], i64 1
+// HOSTCALL-NEXT: [[TMP4:%.*]] = load i8, ptr [[TMP2]], align 1
+// HOSTCALL-NEXT: [[TMP5:%.*]] = icmp eq i8 [[TMP4]], 0
+// HOSTCALL-NEXT: br i1 [[TMP5]], label %[[STRLEN_WHILE_DONE:.*]], label %[[STRLEN_WHILE]]
+// HOSTCALL: [[STRLEN_WHILE_DONE]]:
+// HOSTCALL-NEXT: [[TMP6:%.*]] = ptrtoaddr ptr [[TMP2]] to i64
+// HOSTCALL-NEXT: [[TMP7:%.*]] = sub i64 [[TMP6]], ptrtoaddr (ptr addrspacecast (ptr addrspace(4) @.str to ptr) to i64)
+// HOSTCALL-NEXT: [[TMP8:%.*]] = add i64 [[TMP7]], 1
+// HOSTCALL-NEXT: br label %[[STRLEN_JOIN]]
+// HOSTCALL: [[STRLEN_JOIN]]:
+// HOSTCALL-NEXT: [[TMP9:%.*]] = phi i64 [ [[TMP8]], %[[STRLEN_WHILE_DONE]] ], [ 0, %[[ENTRY]] ]
+// HOSTCALL-NEXT: [[TMP10:%.*]] = call i64 @__ockl_printf_append_string_n(i64 [[TMP0]], ptr addrspacecast (ptr addrspace(4) @.str to ptr), i64 [[TMP9]], i32 1)
+// HOSTCALL-NEXT: [[TMP11:%.*]] = trunc i64 [[TMP10]] to i32
+// HOSTCALL-NEXT: ret i32 [[TMP11]]
+//
+// HOSTCALL-AMDGCNSPIRV-LABEL: define spir_func noundef i32 @_Z4foo1v(
+// HOSTCALL-AMDGCNSPIRV-SAME: ) addrspace(4) #[[ATTR0:[0-9]+]] {
+// HOSTCALL-AMDGCNSPIRV-NEXT: [[ENTRY:.*]]:
+// HOSTCALL-AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = call addrspace(4) i64 @__ockl_printf_begin(i64 0)
+// HOSTCALL-AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = icmp eq ptr addrspace(4) addrspacecast (ptr addrspace(1) @.str to ptr addrspace(4)), null
+// HOSTCALL-AMDGCNSPIRV-NEXT: br i1 [[TMP1]], label %[[STRLEN_JOIN:.*]], label %[[STRLEN_WHILE:.*]]
+// HOSTCALL-AMDGCNSPIRV: [[STRLEN_WHILE]]:
+// HOSTCALL-AMDGCNSPIRV-NEXT: [[TMP2:%.*]] = phi ptr addrspace(4) [ addrspacecast (ptr addrspace(1) @.str to ptr addrspace(4)), %[[ENTRY]] ], [ [[TMP3:%.*]], %[[STRLEN_WHILE]] ]
+// HOSTCALL-AMDGCNSPIRV-NEXT: [[TMP3]] = getelementptr i8, ptr addrspace(4) [[TMP2]], i64 1
+// HOSTCALL-AMDGCNSPIRV-NEXT: [[TMP4:%.*]] = load i8, ptr addrspace(4) [[TMP2]], align 1
+// HOSTCALL-AMDGCNSPIRV-NEXT: [[TMP5:%.*]] = icmp eq i8 [[TMP4]], 0
+// HOSTCALL-AMDGCNSPIRV-NEXT: br i1 [[TMP5]], label %[[STRLEN_WHILE_DONE:.*]], label %[[STRLEN_WHILE]]
+// HOSTCALL-AMDGCNSPIRV: [[STRLEN_WHILE_DONE]]:
+// HOSTCALL-AMDGCNSPIRV-NEXT: [[TMP6:%.*]] = ptrtoaddr ptr addrspace(4) [[TMP2]] to i64
+// HOSTCALL-AMDGCNSPIRV-NEXT: [[TMP7:%.*]] = sub i64 [[TMP6]], ptrtoaddr (ptr addrspace(4) addrspacecast (ptr addrspace(1) @.str to ptr addrspace(4)) to i64)
+// HOSTCALL-AMDGCNSPIRV-NEXT: [[TMP8:%.*]] = add i64 [[TMP7]], 1
+// HOSTCALL-AMDGCNSPIRV-NEXT: br label %[[STRLEN_JOIN]]
+// HOSTCALL-AMDGCNSPIRV: [[STRLEN_JOIN]]:
+// HOSTCALL-AMDGCNSPIRV-NEXT: [[TMP9:%.*]] = phi i64 [ [[TMP8]], %[[STRLEN_WHILE_DONE]] ], [ 0, %[[ENTRY]] ]
+// HOSTCALL-AMDGCNSPIRV-NEXT: [[TMP10:%.*]] = call addrspace(4) i64 @__ockl_printf_append_string_n(i64 [[TMP0]], ptr addrspace(4) addrspacecast (ptr addrspace(1) @.str to ptr addrspace(4)), i64 [[TMP9]], i32 1)
+// HOSTCALL-AMDGCNSPIRV-NEXT: [[TMP11:%.*]] = trunc i64 [[TMP10]] to i32
+// HOSTCALL-AMDGCNSPIRV-NEXT: ret i32 [[TMP11]]
+//
+// BUFFERED-LABEL: define dso_local noundef i32 @_Z4foo1v(
+// BUFFERED-SAME: ) #[[ATTR0:[0-9]+]] {
+// BUFFERED-NEXT: [[ENTRY:.*:]]
+// BUFFERED-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 12)
+// BUFFERED-NEXT: [[TMP0:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
+// BUFFERED-NEXT: br i1 [[TMP0]], label %[[ARGPUSH_BLOCK:.*]], label %[[END_BLOCK:.*]]
+// BUFFERED: [[END_BLOCK]]:
+// BUFFERED-NEXT: [[TMP1:%.*]] = xor i1 [[TMP0]], true
+// BUFFERED-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP1]] to i32
+// BUFFERED-NEXT: ret i32 [[PRINTF_RESULT]]
+// BUFFERED: [[ARGPUSH_BLOCK]]:
+// BUFFERED-NEXT: store i32 50, ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
+// BUFFERED-NEXT: [[TMP2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
+// BUFFERED-NEXT: store i64 -8840842864239206427, ptr addrspace(1) [[TMP2]], align 8
+// BUFFERED-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP2]], i32 8
+// BUFFERED-NEXT: br label %[[END_BLOCK]]
+//
+// BUFFERED-AMDGCNSPIRV-LABEL: define spir_func noundef i32 @_Z4foo1v(
+// BUFFERED-AMDGCNSPIRV-SAME: ) addrspace(4) #[[ATTR0:[0-9]+]] {
+// BUFFERED-AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
+// BUFFERED-AMDGCNSPIRV-NEXT: [[PRINTF_ALLOC_FN:%.*]] = call addrspace(4) ptr addrspace(1) @__printf_alloc(i32 12)
+// BUFFERED-AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = icmp ne ptr addrspace(1) [[PRINTF_ALLOC_FN]], null
+// BUFFERED-AMDGCNSPIRV-NEXT: br i1 [[TMP0]], label %[[ARGPUSH_BLOCK:.*]], label %[[END_BLOCK:.*]]
+// BUFFERED-AMDGCNSPIRV: [[END_BLOCK]]:
+// BUFFERED-AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = xor i1 [[TMP0]], true
+// BUFFERED-AMDGCNSPIRV-NEXT: [[PRINTF_RESULT:%.*]] = sext i1 [[TMP1]] to i32
+// BUFFERED-AMDGCNSPIRV-NEXT: ret i32 [[PRINTF_RESULT]]
+// BUFFERED-AMDGCNSPIRV: [[ARGPUSH_BLOCK]]:
+// BUFFERED-AMDGCNSPIRV-NEXT: store i32 50, ptr addrspace(1) [[PRINTF_ALLOC_FN]], align 4
+// BUFFERED-AMDGCNSPIRV-NEXT: [[TMP2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[PRINTF_ALLOC_FN]], i32 4
+// BUFFERED-AMDGCNSPIRV-NEXT: store i64 -8840842864239206427, ptr addrspace(1) [[TMP2]], align 8
+// BUFFERED-AMDGCNSPIRV-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[TMP2]], i32 8
+// BUFFERED-AMDGCNSPIRV-NEXT: br label %[[END_BLOCK]]
+//
__device__ int foo1() {
- // HOSTCALL: call i64 @__ockl_printf_begin
- // HOSTCALL-AMDGCNSPIRV: call addrspace(4) i64 @__ockl_printf_begin
- // BUFFERED: call ptr addrspace(1) @__printf_alloc
- // BUFFERED-AMDGCNSPIRV: call addrspace(4) ptr addrspace(1) @__printf_alloc
- // CHECK-NOT: call i32 (ptr, ...) @printf
- // CHECK-AMDGCNSPIRV-NOT: call i32 (ptr, ...) @printf
return __builtin_printf("Hello World\n");
}
-// CHECK-LABEL: @_Z4foo2v()
+// HOSTCALL-LABEL: define dso_local noundef i32 @_Z4foo2v(
+// HOSTCALL-SAME: ) #[[ATTR0]] {
+// HOSTCALL-NEXT: [[ENTRY:.*:]]
+// HOSTCALL-NEXT: [[CALL:%.*]] = call i32 (ptr, ...) @printf(ptr noundef addrspacecast (ptr addrspace(4) @.str to ptr)) #[[ATTR2:[0-9]+]]
+// HOSTCALL-NEXT: ret i32 [[CALL]]
+//
+// HOSTCALL-AMDGCNSPIRV-LABEL: define spir_func noundef i32 @_Z4foo2v(
+// HOSTCALL-AMDGCNSPIRV-SAME: ) addrspace(4) #[[ATTR0]] {
+// HOSTCALL-AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
+// HOSTCALL-AMDGCNSPIRV-NEXT: [[CALL:%.*]] = call spir_func addrspace(4) i32 (ptr addrspace(4), ...) @printf(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @.str to ptr addrspace(4))) #[[ATTR2:[0-9]+]]
+// HOSTCALL-AMDGCNSPIRV-NEXT: ret i32 [[CALL]]
+//
+// BUFFERED-LABEL: define dso_local noundef i32 @_Z4foo2v(
+// BUFFERED-SAME: ) #[[ATTR0]] {
+// BUFFERED-NEXT: [[ENTRY:.*:]]
+// BUFFERED-NEXT: [[CALL:%.*]] = call i32 (ptr, ...) @printf(ptr noundef addrspacecast (ptr addrspace(4) @.str to ptr)) #[[ATTR3:[0-9]+]]
+// BUFFERED-NEXT: ret i32 [[CALL]]
+//
+// BUFFERED-AMDGCNSPIRV-LABEL: define spir_func noundef i32 @_Z4foo2v(
+// BUFFERED-AMDGCNSPIRV-SAME: ) addrspace(4) #[[ATTR0]] {
+// BUFFERED-AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
+// BUFFERED-AMDGCNSPIRV-NEXT: [[CALL:%.*]] = call spir_func addrspace(4) i32 (ptr addrspace(4), ...) @printf(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @.str to ptr addrspace(4))) #[[ATTR3:[0-9]+]]
+// BUFFERED-AMDGCNSPIRV-NEXT: ret i32 [[CALL]]
+//
__device__ int foo2() {
- // CHECK: call i32 (ptr, ...) @printf
- // CHECK-AMDGCNSPIRV: call spir_func addrspace(4) i32 (ptr addrspace(4), ...) @printf
return printf("Hello World\n");
}
+
+// HOSTCALL-LABEL: define dso_local noundef i32 @_Z16trailing_percentPKc(
+// HOSTCALL-SAME: ptr noundef [[S:%.*]]) #[[ATTR0]] {
+// HOSTCALL-NEXT: [[ENTRY:.*:]]
+// HOSTCALL-NEXT: [[S_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
+// HOSTCALL-NEXT: [[S_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[S_ADDR]] to ptr
+// HOSTCALL-NEXT: store ptr [[S]], ptr [[S_ADDR_ASCAST]], align 8
+// HOSTCALL-NEXT: [[TMP0:%.*]] = load ptr, ptr [[S_ADDR_ASCAST]], align 8
+// HOSTCALL-NEXT: [[CALL:%.*]] = call i32 (ptr, ...) @printf(ptr noundef addrspacecast (ptr addrspace(4) @.str.1 to ptr), ptr noundef [[TMP0]]) #[[ATTR2]]
+// HOSTCALL-NEXT: ret i32 [[CALL]]
+//
+// HOSTCALL-AMDGCNSPIRV-LABEL: define spir_func noundef i32 @_Z16trailing_percentPKc(
+// HOSTCALL-AMDGCNSPIRV-SAME: ptr addrspace(4) noundef [[S:%.*]]) addrspace(4) #[[ATTR0]] {
+// HOSTCALL-AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
+// HOSTCALL-AMDGCNSPIRV-NEXT: [[S_ADDR:%.*]] = alloca ptr addrspace(4), align 8
+// HOSTCALL-AMDGCNSPIRV-NEXT: [[S_ADDR_ASCAST:%.*]] = addrspacecast ptr [[S_ADDR]] to ptr addrspace(4)
+// HOSTCALL-AMDGCNSPIRV-NEXT: store ptr addrspace(4) [[S]], ptr addrspace(4) [[S_ADDR_ASCAST]], align 8
+// HOSTCALL-AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[S_ADDR_ASCAST]], align 8
+// HOSTCALL-AMDGCNSPIRV-NEXT: [[CALL:%.*]] = call spir_func addrspace(4) i32 (ptr addrspace(4), ...) @printf(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @.str.1 to ptr addrspace(4)), ptr addrspace(4) noundef [[TMP0]]) #[[ATTR2]]
+// HOSTCALL-AMDGCNSPIRV-NEXT: ret i32 [[CALL]]
+//
+// BUFFERED-LABEL: define dso_local noundef i32 @_Z16trailing_percentPKc(
+// BUFFERED-SAME: ptr noundef [[S:%.*]]) #[[ATTR0]] {
+// BUFFERED-NEXT: [[ENTRY:.*:]]
+// BUFFERED-NEXT: [[S_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
+// BUFFERED-NEXT: [[S_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[S_ADDR]] to ptr
+// BUFFERED-NEXT: store ptr [[S]], ptr [[S_ADDR_ASCAST]], align 8
+// BUFFERED-NEXT: [[TMP0:%.*]] = load ptr, ptr [[S_ADDR_ASCAST]], align 8
+// BUFFERED-NEXT: [[CALL:%.*]] = call i32 (ptr, ...) @printf(ptr noundef addrspacecast (ptr addrspace(4) @.str.1 to ptr), ptr noundef [[TMP0]]) #[[ATTR3]]
+// BUFFERED-NEXT: ret i32 [[CALL]]
+//
+// BUFFERED-AMDGCNSPIRV-LABEL: define spir_func noundef i32 @_Z16trailing_percentPKc(
+// BUFFERED-AMDGCNSPIRV-SAME: ptr addrspace(4) noundef [[S:%.*]]) addrspace(4) #[[ATTR0]] {
+// BUFFERED-AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
+// BUFFERED-AMDGCNSPIRV-NEXT: [[S_ADDR:%.*]] = alloca ptr addrspace(4), align 8
+// BUFFERED-AMDGCNSPIRV-NEXT: [[S_ADDR_ASCAST:%.*]] = addrspacecast ptr [[S_ADDR]] to ptr addrspace(4)
+// BUFFERED-AMDGCNSPIRV-NEXT: store ptr addrspace(4) [[S]], ptr addrspace(4) [[S_ADDR_ASCAST]], align 8
+// BUFFERED-AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[S_ADDR_ASCAST]], align 8
+// BUFFERED-AMDGCNSPIRV-NEXT: [[CALL:%.*]] = call spir_func addrspace(4) i32 (ptr addrspace(4), ...) @printf(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @.str.1 to ptr addrspace(4)), ptr addrspace(4) noundef [[TMP0]]) #[[ATTR3]]
+// BUFFERED-AMDGCNSPIRV-NEXT: ret i32 [[CALL]]
+//
+__device__ int trailing_percent(const char *s) { return printf("%s%", s); }
diff --git a/clang/test/CodeGenHIP/printf-trailing-percent.hip b/clang/test/CodeGenHIP/printf-trailing-percent.hip
deleted file mode 100644
index 1794b0f48ecaf..0000000000000
--- a/clang/test/CodeGenHIP/printf-trailing-percent.hip
+++ /dev/null
@@ -1,30 +0,0 @@
-// REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -emit-llvm -fcuda-is-device \
-// RUN: -Wno-format -mprintf-kind=hostcall -o - %s \
-// RUN: | FileCheck --check-prefix=HOSTCALL %s
-// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -emit-llvm -fcuda-is-device \
-// RUN: -Wno-format -mprintf-kind=buffered -o - %s \
-// RUN: | FileCheck --check-prefix=BUFFERED %s
-
-// A format string ending in an incomplete specifier must not be read past its
-// end, and the specifiers before it must still be found.
-
-#define __device__ __attribute__((device))
-
-extern "C" __device__ int printf(const char *format, ...);
-
-__device__ int trailing_percent(const char *s) { return printf("%s%", s); }
-
-// The %s argument is printed as a string, not as a pointer.
-
-// HOSTCALL-LABEL: define {{.*}} @_Z16trailing_percentPKc(
-// HOSTCALL: call i64 @__ockl_printf_begin(i64 0)
-// HOSTCALL: call i64 @__ockl_printf_append_string_n(i64 %{{.*}}, ptr {{.*}}@.str{{.*}}, i64 %{{.*}}, i32 0)
-// HOSTCALL: call i64 @__ockl_printf_append_string_n(i64 %{{.*}}, ptr %{{.*}}, i64 %{{.*}}, i32 1)
-// HOSTCALL-NOT: call i64 @__ockl_printf_append_args
-
-// BUFFERED-LABEL: define {{.*}} @_Z16trailing_percentPKc(
-// BUFFERED: call ptr addrspace(1) @__printf_alloc(
-// BUFFERED: call void @llvm.memcpy.p1.p0.i64(
-// BUFFERED-NOT: ptrtoint
-// BUFFERED: !{!"0:0:{{[0-9a-f]+}},%s%"}
More information about the cfe-commits
mailing list