This is an archive of the discontinued LLVM Phabricator instance.

Differential D122669

[AMDGPU][OpenCL] Remove "printf and hostcall" diagnostic
AbandonedPublic

Authored by scott.linder on Mar 29 2022, 10:25 AM.

Details

Reviewers
None
Summary

The diagnostic is unreliable, and triggers even for dead uses of
hostcall that may exist when linking the device-libs at lower
optimization levels.

Eliminate the diagnostic, and directly document the limitation for
OpenCL before code object V5.

Make some NFC changes to clarify the related code in the
MetadataStreamer.

Add a clang test to tie OCL sources containing printf to the backend IR
tests for this situation.

Diff Detail

Event Timeline

scott.linder created this revision.Mar 29 2022, 10:25 AM
Herald added a project: Restricted Project. · View Herald TranscriptMar 29 2022, 10:25 AM
Herald added subscribers: hsmhsm, foad, Naghasan and 12 others. · View Herald Transcript
scott.linder requested review of this revision.Mar 29 2022, 10:25 AM
Herald added projects: Restricted Project, Restricted Project. · View Herald TranscriptMar 29 2022, 10:25 AM
Herald added subscribers: llvm-commits, cfe-commits, wdng. · View Herald Transcript
Harbormaster completed remote builds in B156794: Diff 418921.Mar 29 2022, 10:26 AM
scott.linder abandoned this revision.Mar 29 2022, 10:26 AM

Revision Contents

PathSize
clang/
test/
CodeGenOpenCL/
46 lines
llvm/
docs/
18 lines
lib/
Target/
AMDGPU/
24 lines
9 lines
test/
CodeGen/
AMDGPU/
19 lines

Diff 418921

clang/test/CodeGenOpenCL/amdgpu-printf.cl

  • This file was added.
// 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
int printf(__constant const char* st, ...) __attribute__((format(printf, 1, 2)));
// CHECK-LABEL: @test_printf_noargs(
// CHECK-NEXT: entry:
// CHECK-NEXT: [[CALL:%.*]] = call i32 (i8 addrspace(4)*, ...) @printf(i8 addrspace(4)* noundef getelementptr inbounds ([1 x i8], [1 x i8] addrspace(4)* @.str, i64 0, i64 0)) #[[ATTR4:[0-9]+]]
// CHECK-NEXT: ret void
//
__kernel void test_printf_noargs() {
printf("");
}
// CHECK-LABEL: @test_printf_int(
// CHECK-NEXT: entry:
// CHECK-NEXT: [[I_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK-NEXT: store i32 [[I:%.*]], i32 addrspace(5)* [[I_ADDR]], align 4, !tbaa [[TBAA8:![0-9]+]]
// CHECK-NEXT: [[TMP0:%.*]] = load i32, i32 addrspace(5)* [[I_ADDR]], align 4, !tbaa [[TBAA8]]
// CHECK-NEXT: [[CALL:%.*]] = call i32 (i8 addrspace(4)*, ...) @printf(i8 addrspace(4)* noundef getelementptr inbounds ([3 x i8], [3 x i8] addrspace(4)* @.str.1, i64 0, i64 0), i32 noundef [[TMP0]]) #[[ATTR4]]
// CHECK-NEXT: ret void
//
__kernel void test_printf_int(int i) {
printf("%d", i);
}
// CHECK-LABEL: @test_printf_str_int(
// CHECK-NEXT: entry:
// CHECK-NEXT: [[I_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK-NEXT: [[S:%.*]] = alloca [4 x i8], align 1, addrspace(5)
// CHECK-NEXT: store i32 [[I:%.*]], i32 addrspace(5)* [[I_ADDR]], align 4, !tbaa [[TBAA8]]
// CHECK-NEXT: [[TMP0:%.*]] = bitcast [4 x i8] addrspace(5)* [[S]] to i8 addrspace(5)*
// CHECK-NEXT: call void @llvm.lifetime.start.p5i8(i64 4, i8 addrspace(5)* [[TMP0]]) #[[ATTR5:[0-9]+]]
// CHECK-NEXT: [[TMP1:%.*]] = bitcast [4 x i8] addrspace(5)* [[S]] to i8 addrspace(5)*
// CHECK-NEXT: call void @llvm.memcpy.p5i8.p4i8.i64(i8 addrspace(5)* align 1 [[TMP1]], i8 addrspace(4)* align 1 getelementptr inbounds ([4 x i8], [4 x i8] addrspace(4)* @__const.test_printf_str_int.s, i32 0, i32 0), i64 4, i1 false)
// CHECK-NEXT: [[ARRAYDECAY:%.*]] = getelementptr inbounds [4 x i8], [4 x i8] addrspace(5)* [[S]], i64 0, i64 0
// CHECK-NEXT: [[TMP2:%.*]] = load i32, i32 addrspace(5)* [[I_ADDR]], align 4, !tbaa [[TBAA8]]
// CHECK-NEXT: [[CALL:%.*]] = call i32 (i8 addrspace(4)*, ...) @printf(i8 addrspace(4)* noundef getelementptr inbounds ([6 x i8], [6 x i8] addrspace(4)* @.str.2, i64 0, i64 0), i8 addrspace(5)* noundef [[ARRAYDECAY]], i32 noundef [[TMP2]]) #[[ATTR4]]
// CHECK-NEXT: [[TMP3:%.*]] = bitcast [4 x i8] addrspace(5)* [[S]] to i8 addrspace(5)*
// CHECK-NEXT: call void @llvm.lifetime.end.p5i8(i64 4, i8 addrspace(5)* [[TMP3]]) #[[ATTR5]]
// CHECK-NEXT: ret void
//
__kernel void test_printf_str_int(int i) {
char s[] = "foo";
printf("%s:%d", s, i);
}

llvm/docs/AMDGPUUsage.rst

  • This file is larger than 256 KB, so syntax highlighting is disabled by default.
Show First 20 Lines • Show All 2,821 Lines • ▼ Show 20 Lines "ValueKind" string Required Kernel argument kind that
An argument that is not used An argument that is not used
by the kernel. Space needs to by the kernel. Space needs to
be left for it, but it does be left for it, but it does
not need to be set up. not need to be set up.
"HiddenPrintfBuffer" "HiddenPrintfBuffer"
A global address space pointer A global address space pointer
to the runtime printf buffer to the runtime printf buffer
is passed in kernarg. is passed in kernarg. Mutually
exclusive with
"HiddenHostcallBuffer".
"HiddenHostcallBuffer" "HiddenHostcallBuffer"
A global address space pointer A global address space pointer
to the runtime hostcall buffer to the runtime hostcall buffer
is passed in kernarg. is passed in kernarg. Mutually
exclusive with
"HiddenPrintfBuffer".
"HiddenDefaultQueue" "HiddenDefaultQueue"
A global address space pointer A global address space pointer
to the OpenCL device enqueue to the OpenCL device enqueue
queue that should be used by queue that should be used by
the kernel by default is the kernel by default is
passed in the kernarg. passed in the kernarg.
▲ Show 20 Lines • Show All 503 Lines • ▼ Show 20 Lines ".value_kind" string Required Kernel argument kind that
An argument that is not used An argument that is not used
by the kernel. Space needs to by the kernel. Space needs to
be left for it, but it does be left for it, but it does
not need to be set up. not need to be set up.
"hidden_printf_buffer" "hidden_printf_buffer"
A global address space pointer A global address space pointer
to the runtime printf buffer to the runtime printf buffer
is passed in kernarg. is passed in kernarg. Mutually
exclusive with
"hidden_hostcall_buffer"
before Code Object V5.
"hidden_hostcall_buffer" "hidden_hostcall_buffer"
A global address space pointer A global address space pointer
to the runtime hostcall buffer to the runtime hostcall buffer
is passed in kernarg. is passed in kernarg. Mutually
exclusive with
"hidden_printf_buffer"
before Code Object V5.
"hidden_default_queue" "hidden_default_queue"
A global address space pointer A global address space pointer
to the OpenCL device enqueue to the OpenCL device enqueue
queue that should be used by queue that should be used by
the kernel by default is the kernel by default is
passed in the kernarg. passed in the kernarg.
▲ Show 20 Lines • Show All 11,652 Lines • Show Last 20 Lines

llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp

Show First 20 Lines • Show All 395 Lines • ▼ Show 20 Linesvoid MetadataStreamerV2::emitHiddenKernelArgs(const Function &Func,
if (HiddenArgNumBytes >= 16) if (HiddenArgNumBytes >= 16)
emitKernelArg(DL, Int64Ty, Align(8), ValueKind::HiddenGlobalOffsetY); emitKernelArg(DL, Int64Ty, Align(8), ValueKind::HiddenGlobalOffsetY);
if (HiddenArgNumBytes >= 24) if (HiddenArgNumBytes >= 24)
emitKernelArg(DL, Int64Ty, Align(8), ValueKind::HiddenGlobalOffsetZ); emitKernelArg(DL, Int64Ty, Align(8), ValueKind::HiddenGlobalOffsetZ);
auto Int8PtrTy = Type::getInt8PtrTy(Func.getContext(), auto Int8PtrTy = Type::getInt8PtrTy(Func.getContext(),
AMDGPUAS::GLOBAL_ADDRESS); AMDGPUAS::GLOBAL_ADDRESS);
// Emit "printf buffer" argument if printf is used, otherwise emit dummy
// "none" argument.
if (HiddenArgNumBytes >= 32) { if (HiddenArgNumBytes >= 32) {
// We forbid the use of features requiring hostcall when compiling OpenCL
// before code object V5, which makes the mutual exclusion between the
// "printf buffer" and "hostcall buffer" here sound.
if (Func.getParent()->getNamedMetadata("llvm.printf.fmts")) if (Func.getParent()->getNamedMetadata("llvm.printf.fmts"))
emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenPrintfBuffer); emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenPrintfBuffer);
else if (!Func.hasFnAttribute("amdgpu-no-hostcall-ptr")) { else if (!Func.hasFnAttribute("amdgpu-no-hostcall-ptr"))
// The printf runtime binding pass should have ensured that hostcall and
// printf are not used in the same module.
assert(!Func.getParent()->getNamedMetadata("llvm.printf.fmts"));
emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenHostcallBuffer); emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenHostcallBuffer);
} else else
emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenNone); emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenNone);
} }
// Emit "default queue" and "completion action" arguments if enqueue kernel is // Emit "default queue" and "completion action" arguments if enqueue kernel is
// used, otherwise emit dummy "none" arguments. // used, otherwise emit dummy "none" arguments.
if (HiddenArgNumBytes >= 48) { if (HiddenArgNumBytes >= 48) {
if (Func.hasFnAttribute("calls-enqueue-kernel")) { if (Func.hasFnAttribute("calls-enqueue-kernel")) {
emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenDefaultQueue); emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenDefaultQueue);
▲ Show 20 Lines • Show All 392 Lines • ▼ Show 20 Lines emitKernelArg(DL, Int64Ty, Align(8), "hidden_global_offset_y", Offset,
Args); Args);
if (HiddenArgNumBytes >= 24) if (HiddenArgNumBytes >= 24)
emitKernelArg(DL, Int64Ty, Align(8), "hidden_global_offset_z", Offset, emitKernelArg(DL, Int64Ty, Align(8), "hidden_global_offset_z", Offset,
Args); Args);
auto Int8PtrTy = auto Int8PtrTy =
Type::getInt8PtrTy(Func.getContext(), AMDGPUAS::GLOBAL_ADDRESS); Type::getInt8PtrTy(Func.getContext(), AMDGPUAS::GLOBAL_ADDRESS);
// Emit "printf buffer" argument if printf is used, emit "hostcall buffer"
// if "hostcall" module flag is set, otherwise emit dummy "none" argument.
if (HiddenArgNumBytes >= 32) { if (HiddenArgNumBytes >= 32) {
// We forbid the use of features requiring hostcall when compiling OpenCL
// before code object V5, which makes the mutual exclusion between the
// "printf buffer" and "hostcall buffer" here sound.
if (M->getNamedMetadata("llvm.printf.fmts")) if (M->getNamedMetadata("llvm.printf.fmts"))
emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_printf_buffer", Offset, emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_printf_buffer", Offset,
Args); Args);
else if (MFI.hasHostcallPtr()) { else if (MFI.hasHostcallPtr())
// The printf runtime binding pass should have ensured that hostcall and
// printf are not used in the same module.
assert(!M->getNamedMetadata("llvm.printf.fmts"));
emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_hostcall_buffer", Offset, emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_hostcall_buffer", Offset,
Args); Args);
} else else
emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_none", Offset, Args); emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_none", Offset, Args);
} }
// Emit "default queue" and "completion action" arguments if enqueue kernel is // Emit "default queue" and "completion action" arguments if enqueue kernel is
// used, otherwise emit dummy "none" arguments. // used, otherwise emit dummy "none" arguments.
if (HiddenArgNumBytes >= 48) { if (HiddenArgNumBytes >= 48) {
if (Func.hasFnAttribute("calls-enqueue-kernel")) { if (Func.hasFnAttribute("calls-enqueue-kernel")) {
emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_default_queue", Offset, emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_default_queue", Offset,
▲ Show 20 Lines • Show All 220 Lines • Show Last 20 Lines

llvm/lib/Target/AMDGPU/AMDGPUPrintfRuntimeBinding.cpp

Show First 20 Lines • Show All 557 Lines • ▼ Show 20 Lines if (auto *CI = dyn_cast<CallInst>(U.getUser())) {
if (CI->isCallee(&U)) if (CI->isCallee(&U))
Printfs.push_back(CI); Printfs.push_back(CI);
} }
} }
if (Printfs.empty()) if (Printfs.empty())
return false; return false;
if (auto HostcallFunction = M.getFunction("__ockl_hostcall_internal")) {
for (auto &U : HostcallFunction->uses()) {
if (auto *CI = dyn_cast<CallInst>(U.getUser())) {
M.getContext().emitError(
CI, "Cannot use both printf and hostcall in the same module");
}
}
}
TD = &M.getDataLayout(); TD = &M.getDataLayout();
return lowerPrintfForGpu(M); return lowerPrintfForGpu(M);
} }
bool AMDGPUPrintfRuntimeBinding::runOnModule(Module &M) { bool AMDGPUPrintfRuntimeBinding::runOnModule(Module &M) {
auto GetDT = [this](Function &F) -> DominatorTree & { auto GetDT = [this](Function &F) -> DominatorTree & {
return this->getAnalysis<DominatorTreeWrapperPass>(F).getDomTree(); return this->getAnalysis<DominatorTreeWrapperPass>(F).getDomTree();
Show All 21 Lines

llvm/test/CodeGen/AMDGPU/opencl-printf-and-hostcall.ll

  • This file was added.
; RUN: opt -S -mtriple=amdgcn-unknown-unknown -amdgpu-printf-runtime-binding < %s 2>&1 | FileCheck %s
@.str = private unnamed_addr addrspace(4) constant [6 x i8] c"%s:%d\00", align 1
define amdgpu_kernel void @test_kernel(i32 %n) {
entry:
%str = alloca [9 x i8], align 1, addrspace(5)
%arraydecay = getelementptr inbounds [9 x i8], [9 x i8] addrspace(5)* %str, i32 0, i32 0
%call1 = call i32 (i8 addrspace(4)*, ...) @printf(i8 addrspace(4)* getelementptr inbounds ([6 x i8], [6 x i8] addrspace(4)* @.str, i32 0, i32 0), i8 addrspace(5)* %arraydecay, i32 %n)
%call2 = call <2 x i64> (i8*, i32, i64, i64, i64, i64, i64, i64, i64, i64) @__ockl_hostcall_internal(i8* undef, i32 1, i64 2, i64 3, i64 4, i64 5, i64 6, i64 7, i64 8, i64 9)
ret void
}
declare i32 @printf(i8 addrspace(4)*, ...)
declare <2 x i64> @__ockl_hostcall_internal(i8*, i32, i64, i64, i64, i64, i64, i64, i64, i64)
; CHECK-NOT: error:
; CHECK-NOT: warning:

llvm/test/CodeGen/AMDGPU/opencl-printf-no-hostcall.ll

  • This file was deleted.
; RUN: not opt -S -mtriple=amdgcn-unknown-unknown -amdgpu-printf-runtime-binding < %s 2>&1 | FileCheck %s
@.str = private unnamed_addr addrspace(2) constant [6 x i8] c"%s:%d\00", align 1
define amdgpu_kernel void @test_kernel(i32 %n) {
entry:
%str = alloca [9 x i8], align 1
%arraydecay = getelementptr inbounds [9 x i8], [9 x i8]* %str, i32 0, i32 0
%call1 = call i32 (i8 addrspace(2)*, ...) @printf(i8 addrspace(2)* getelementptr inbounds ([6 x i8], [6 x i8] addrspace(2)* @.str, i32 0, i32 0), i8* %arraydecay, i32 %n)
%call2 = call <2 x i64> (i8*, i32, i64, i64, i64, i64, i64, i64, i64, i64) @__ockl_hostcall_internal(i8* undef, i32 1, i64 2, i64 3, i64 4, i64 5, i64 6, i64 7, i64 8, i64 9)
ret void
}
declare i32 @printf(i8 addrspace(2)*, ...)
declare <2 x i64> @__ockl_hostcall_internal(i8*, i32, i64, i64, i64, i64, i64, i64, i64, i64)
; CHECK: error: Cannot use both printf and hostcall in the same module
llvm/test/CodeGen/AMDGPU/opencl-printf-and-hostcall.ll