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

[scott.linder]

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.

[scott.linder]

For reference, the error I'm changing to a warning was added as part of https://reviews.llvm.org/D70038, and the issue that made me consider this approach appears when building OCL conformance tests at -O0

Another approach would be to do some more work in Clang or early in the backend to strip out the unused device-libs functions which are calling __ockl_hostcall_internal. I started with that thought, but backed off as it seemed the less we could do at -O0 the better. The limitation is also being resolved in later HSA ABI versions, so this would be a legacy path almost as soon as we added it.

[scott.linder]

llvm/lib/Target/AMDGPU/AMDGPUPrintfRuntimeBinding.cpp

565–566

I'm also not certain this is the the best option to classify the diagnostic, but it is what was used for the error so I just left it as-is. Suggestions welcome!

565–566

I will land this separately, as it is unrelated and NFC

[arsenm]

A spurious warning still isn't ideal but this is an improvement

[sameerds]

The check for "__ockl_hostcall_internal" is not longer relevant with the new hostcall attribute. Can we simply remove this check? What is the situation where the warning is still useful?

[scott.linder]

In D121951#3391472, @sameerds wrote:

The check for "__ockl_hostcall_internal" is not longer relevant with the new hostcall attribute. Can we simply remove this check? What is the situation where the warning is still useful?

I wasn't aware of the new attribute, I'm happy to just delete it.

[sameerds]

This comment was removed by sameerds.

[sameerds]

In D121951#3392470, @scott.linder wrote:

In D121951#3391472, @sameerds wrote:

The check for "__ockl_hostcall_internal" is not longer relevant with the new hostcall attribute. Can we simply remove this check? What is the situation where the warning is still useful?

I wasn't aware of the new attribute, I'm happy to just delete it.

Yeah, it happened in D119216.

But I am still curious about the check itself. Do we need to worry about a situation where it is important to check whether OpenCL printf (non-hostcall) and some other hostcall service are active in the same application?

[scott.linder]

In D121951#3393928, @sameerds wrote:

In D121951#3392470, @scott.linder wrote:

In D121951#3391472, @sameerds wrote:

The check for "__ockl_hostcall_internal" is not longer relevant with the new hostcall attribute. Can we simply remove this check? What is the situation where the warning is still useful?

I wasn't aware of the new attribute, I'm happy to just delete it.

Yeah, it happened in D119216.

But I am still curious about the check itself. Do we need to worry about a situation where it is important to check whether OpenCL printf (non-hostcall) and some other hostcall service are active in the same application?

I don't know how important it is to notify the user, but my understanding is pre-code-object-v5 we will use the same kernarg for both the printf and hostcall pointers. I don't know what actually happens (does the runtime notice it can't actually construct the kernarg and fail? does the code just fail when one of the two uses of the kernarg are incorrect?). Should I try to nail down exactly what currently happens?

[scott.linder]

Eliminate diagnostic, and instead add function attribute in Clang for pre-COV_5
OpenCL codegen.

[sameerds]

clang/lib/CodeGen/TargetInfo.cpp

9381 ↗

(On Diff #418135)

The frontend does not need to worry about this attribute. See the comment in the MetadataStreamer. A worthwhile check would be to generate an error if we are able to detect that some hostcall service is being used in OpenCL on code-object-v4 or lower. None exists right now, but we should add the check if such services show up. But those checks are likely to be in a different place. For example, enabling asan on OpenCL for code-object-v4 should result in an error in the place where asan commandline options are parsed.

llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp

408–410

I would structure this differently: If this is code-object-v4 or lower, then if "llvm.printf.fmts" is present, then this kernel clearly contains OpenCL bits, and cannot use hostcall. So it's okay to just assume that the no-hostcall-ptr attribute is always present in this situation, which means the only metadata generated is for ValueKind::HiddenPrintfBuffer. Else if this is code-object-v5, then proceed to emit both metadata.

[arsenm]

clang/lib/CodeGen/TargetInfo.cpp

9381 ↗

(On Diff #418135)

Should be all opencl, not just kernels. Also < instead of !=?

[scott.linder]

clang/lib/CodeGen/TargetInfo.cpp

9381 ↗

(On Diff #418135)

The frontend does not need to worry about this attribute. See the comment in the MetadataStreamer.

The reason I went this route is that otherwise the MetadataStreamer can only go off of the presence of the OCL printf metadata to infer that hostcall argument metadata is invalid and will be ignored by the runtime. Ideally, the MetadataStreamer or Verifier or something would actually require that a module does not have both OCL printf metadata and functions which are not "amdgpu-no-hostcall-ptr", but I didn't go that far as it would break old IR/bitcode that relies on the implicit behavior.

I'm OK with removing this code, and the rest of the patch remains essentially unchanged. We will still have an implicit rule based on code-object-version and presence of printf metadata, and at least conceptually you will still be able to compile OCL for pre-V5 and end up with hostcall argument metadata. That case is benign if the runtime just ignores it, but it still seems needlessly relaxed when we can just add the correct attribute in Clang codegen.

Should be all opencl, not just kernels. Also < instead of !=?

Yes, my mistake, I'll update this

llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp

408–410

I'm not sure I follow; doesn't the code as-is implement what you're describing?

If the printf metadata is present, this will only ever emit the HiddenPrintfBuffer, irrespective of the hostcall attribute. Otherwise, this respects amdgpu-no-hostcall-ptr (i.e. for HIP and other languages).

The "if this is code-object-v4 or lower" portion is implicit, as this is just the MetadataStreamerV2 impl. The MetadataStreamerV3 (below) and MetadataStreamerV4 (inherits from V3) impls below are similar. The MetadataStreamerV5 impl is already correct for V5 (i.e. supports both argument kinds for the same kernel).

[yaxunl]

clang/lib/CodeGen/TargetInfo.cpp

9381 ↗

(On Diff #418135)

+1 for using < instead of !=

However, device library is written in OpenCL language. There are functions using hostcall buffer. We cannot mark all OpenCL functions with this attribute.

[scott.linder]

Do not add the no-hostcall attribute for "COV_None" modules, to allow for
the OpenCL implementation of hostcall in the device-libs.

[scott.linder]

@yaxunl Does excluding device-libs via COV_None make sense?

@sameerds Would you still rather we just not add this attribute in the frontend at all? I'm OK with it if that is the consensus

[yaxunl]

In D121951#3411856, @scott.linder wrote:

@yaxunl Does excluding device-libs via COV_None make sense?

That should work. Hopefully, it can cover most spurious warnings.

[sameerds]

In D121951#3411856, @scott.linder wrote:

@yaxunl Does excluding device-libs via COV_None make sense?

@sameerds Would you still rather we just not add this attribute in the frontend at all? I'm OK with it if that is the consensus

Yes, I still think there is no need to emit that attribute in the frontend. It will always be inferred by the Attributor when optimization is enabled. This also eliminates the check for COV_None and there seems to be some uncertainty about COV_None anyway. This also eliminates the updates to all the tests where the no-hostcall-ptr attribute does not actually matter. If ever we need to check if hostcall is being used on OpenCL and COV < 5, we should do it per feature and inform the user appropriately.

llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp

408–410

Your last para about different streamers cleared the confusion. So, this looks good.

[scott.linder]

Remove frontend changes, and just drop the diagnostic in the backend. Update
some comments and naming.

[scott.linder]

In D121951#3414271, @sameerds wrote:

In D121951#3411856, @scott.linder wrote:

@yaxunl Does excluding device-libs via COV_None make sense?

@sameerds Would you still rather we just not add this attribute in the frontend at all? I'm OK with it if that is the consensus

Yes, I still think there is no need to emit that attribute in the frontend. It will always be inferred by the Attributor when optimization is enabled. This also eliminates the check for COV_None and there seems to be some uncertainty about COV_None anyway. This also eliminates the updates to all the tests where the no-hostcall-ptr attribute does not actually matter. If ever we need to check if hostcall is being used on OpenCL and COV < 5, we should do it per feature and inform the user appropriately.

OK, that makes sense to me. I think I was not interpreting the purpose of the attributes correctly, but if they are essentially just optimization remarks then I don't have any issue with not ensuring they are present at -O0.

I updated the patch, let me know what you think.

[sameerds]

LGTM! But not sure if @yaxunl and/or @arsenm should have another look at the final version.

[yaxunl]

LGTM.

clang/test/CodeGenOpenCL/amdgpu-printf.cl

+// 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

"ValueKind"       string         Required  Kernel argument kind that
                                                   An argument that is not used
                                                   by the kernel. Space needs to
                                                   be left for it, but it does
                                                   not need to be set up.
 
                                                 "HiddenPrintfBuffer"
                                                   A global address space pointer
                                                   to the runtime printf buffer
-                                                  is passed in kernarg.
+                                                  is passed in kernarg. Mutually
+                                                  exclusive with
+                                                  "HiddenHostcallBuffer".
 
                                                 "HiddenHostcallBuffer"
                                                   A global address space pointer
                                                   to the runtime hostcall buffer
-                                                  is passed in kernarg.
+                                                  is passed in kernarg. Mutually
+                                                  exclusive with
+                                                  "HiddenPrintfBuffer".
 
                                                 "HiddenDefaultQueue"
                                                   A global address space pointer
                                                   to the OpenCL device enqueue
                                                   queue that should be used by
                                                   the kernel by default is
                                                   passed in the kernarg.
 
".value_kind"          string         Required  Kernel argument kind that
                                                        An argument that is not used
                                                        by the kernel. Space needs to
                                                        be left for it, but it does
                                                        not need to be set up.
 
                                                      "hidden_printf_buffer"
                                                        A global address space pointer
                                                        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"
                                                        A global address space pointer
                                                        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"
                                                        A global address space pointer
                                                        to the OpenCL device enqueue
                                                        queue that should be used by
                                                        the kernel by default is
                                                        passed in the kernarg.
 

llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp

void MetadataStreamerV2::emitHiddenKernelArgs(const Function &Func,
   if (HiddenArgNumBytes >= 16)
     emitKernelArg(DL, Int64Ty, Align(8), ValueKind::HiddenGlobalOffsetY);
   if (HiddenArgNumBytes >= 24)
     emitKernelArg(DL, Int64Ty, Align(8), ValueKind::HiddenGlobalOffsetZ);
 
   auto Int8PtrTy = Type::getInt8PtrTy(Func.getContext(),
                                       AMDGPUAS::GLOBAL_ADDRESS);
 
-  // Emit "printf buffer" argument if printf is used, otherwise emit dummy
-  // "none" argument.
   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"))
       emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenPrintfBuffer);
-    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"));
+    else if (!Func.hasFnAttribute("amdgpu-no-hostcall-ptr"))
       emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenHostcallBuffer);
-    } else
+    else
       emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenNone);
   }
 
   // Emit "default queue" and "completion action" arguments if enqueue kernel is
   // used, otherwise emit dummy "none" arguments.
   if (HiddenArgNumBytes >= 48) {
     if (Func.hasFnAttribute("calls-enqueue-kernel")) {
       emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenDefaultQueue);
emitKernelArg(DL, Int64Ty, Align(8), "hidden_global_offset_y", Offset,
                   Args);
   if (HiddenArgNumBytes >= 24)
     emitKernelArg(DL, Int64Ty, Align(8), "hidden_global_offset_z", Offset,
                   Args);
 
   auto Int8PtrTy =
       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) {
+    // 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"))
       emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_printf_buffer", Offset,
                     Args);
-    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"));
+    else if (MFI.hasHostcallPtr())
       emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_hostcall_buffer", Offset,
                     Args);
-    } else
+    else
       emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_none", Offset, Args);
   }
 
   // Emit "default queue" and "completion action" arguments if enqueue kernel is
   // used, otherwise emit dummy "none" arguments.
   if (HiddenArgNumBytes >= 48) {
     if (Func.hasFnAttribute("calls-enqueue-kernel")) {
       emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_default_queue", Offset,

llvm/lib/Target/AMDGPU/AMDGPUPrintfRuntimeBinding.cpp

for (auto &U : PrintfFunction->uses()) {
     if (auto *CI = dyn_cast<CallInst>(U.getUser())) {
       if (CI->isCallee(&U))
         Printfs.push_back(CI);
     }
   }
 
   if (Printfs.empty())
     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();
 
   return lowerPrintfForGpu(M);
 }
 
 bool AMDGPUPrintfRuntimeBinding::runOnModule(Module &M) {
   auto GetDT = [this](Function &F) -> DominatorTree & {
     return this->getAnalysis<DominatorTreeWrapperPass>(F).getDomTree();
   };

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

+; 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

-; 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