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 <[email protected]> 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 <[email protected]> --- .../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 <[email protected]> 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 <[email protected]> --- 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%"} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
