[SYCL][HIP] Implement buffered printf for the AMD backend

clang/lib/CodeGen/CodeGenModule.cpp

@@ -1255,6 +1255,17 @@ void CodeGenModule::Release() {
                                 : "buffered");
       getModule().addModuleFlag(llvm::Module::Error, "amdgpu_printf_kind",
                                 MDStr);
+    } else if (LangOpts.SYCLIsDevice) {
+      // SYCL device code on AMDGPU always uses the buffered printf scheme: the
+      // hostcall scheme relies on ROCm's hostcall runtime service (and PCIe
+      // atomics), which the SYCL/UR HIP runtime does not set up. The buffered
+      // scheme only needs the printf buffer implicit kernarg, which the HIP
+      // runtime allocates and parses automatically from the code object's
+      // printf metadata. This flag tells AMDGPUPrintfRuntimeBinding to lower
+      // printf using the buffered layout expected by the runtime.
+      getModule().addModuleFlag(llvm::Module::Error, "amdgpu_printf_kind",
+                                llvm::MDString::get(getLLVMContext(),
+                                                    "buffered"));
     }
   }
 

libclc/libspirv/lib/amdgpu/CMakeLists.txt

@@ -65,3 +65,12 @@ libclc_add_sources(${LIBCLC_LIBSPIRV_TARGET} FILES
   async/wait_group_events.cl
   assert/__assert_fail.cl
 )
+
+# Provide the AMDGPU buffered-printf allocator helper. The AMDGPUPrintfRuntimeBinding
+# backend pass lowers printf calls to __printf_alloc + a device print buffer; the
+# definition is shared with the OpenCL libclc build.
+libclc_add_sources(${LIBCLC_LIBSPIRV_TARGET}
+  BASE_DIR ${CMAKE_CURRENT_SOURCE_DIR}/../../../opencl/lib/amdgpu
+  FILES
+  printf/__printf_alloc.cl
+)

libclc/opencl/lib/amdgpu/printf/__printf_alloc.cl

@@ -11,8 +11,13 @@
 #define OFFSET 8
 
 // Atomically reserves space to the printf data buffer and returns a pointer to
-// it
-__global char *__printf_alloc(uint bytes) {
+// it.
+//
+// This helper is referenced by the AMDGPUPrintfRuntimeBinding backend pass,
+// which only introduces the reference late in the pipeline (after the device
+// bitcode libraries have already been internalized and dead-stripped). Mark it
+// as used so the definition is retained and available when the pass runs.
+__attribute__((used)) __global char *__printf_alloc(uint bytes) {
   __constant amdhsa_implicit_kernarg_v5 *args =
       (__constant amdhsa_implicit_kernarg_v5 *)
           __builtin_amdgcn_implicitarg_ptr();

llvm/lib/SYCLPostLink/ModuleSplitter.cpp

@@ -295,7 +295,11 @@ void collectFunctionsAndGlobalVariablesToExtract(
 
 static bool isIntrinsicOrBuiltin(const Function &F) {
   return F.isIntrinsic() || F.getName().starts_with("__") ||
-         isSpirvSyclBuiltin(F.getName()) || isESIMDBuiltin(F.getName());
+         // printf is a standard C library builtin. On AMDGPU it is left as an
+         // undefined declaration until the AMDGPUPrintfRuntimeBinding backend
+         // pass lowers it, so it is not an undefined *user* function.
+         F.getName() == "printf" || isSpirvSyclBuiltin(F.getName()) ||
+         isESIMDBuiltin(F.getName());
 }
 
 // Checks for use of undefined user functions and emits a warning message.

llvm/lib/Target/AMDGPU/AMDGPUPrintfRuntimeBinding.cpp

@@ -29,6 +29,7 @@
 #include "llvm/InitializePasses.h"
 #include "llvm/Support/DataExtractor.h"
 #include "llvm/TargetParser/Triple.h"
+#include "llvm/Transforms/Utils/AMDGPUEmitPrintf.h"
 #include "llvm/Transforms/Utils/BasicBlockUtils.h"
 
 using namespace llvm;
@@ -59,6 +60,14 @@ class AMDGPUPrintfRuntimeBindingImpl {
 
   bool lowerPrintfForGpu(Module &M);
 
+  // Lowers printf calls using the buffered scheme shared with clang's HIP
+  // -mprintf-kind=buffered lowering (llvm::emitAMDGPUPrintfCall). This produces
+  // a device print buffer layout that the HIP/OpenCL runtime parses. Unlike
+  // lowerPrintfForGpu(), it can be run after inlining, which is required for
+  // SYCL where printf calls are wrapped in a helper and the format string is
+  // only a resolvable constant once the wrapper has been inlined.
+  bool lowerPrintfForGpuBuffered(Module &M);
+
   const DataLayout *TD;
   SmallVector<CallInst *, 32> Printfs;
 };
@@ -423,6 +432,36 @@ bool AMDGPUPrintfRuntimeBindingImpl::lowerPrintfForGpu(Module &M) {
   return true;
 }
 
+bool AMDGPUPrintfRuntimeBindingImpl::lowerPrintfForGpuBuffered(Module &M) {
+  for (auto *CI : Printfs) {
+    // Collect the printf call arguments (format string followed by the
+    // C-vararg-promoted arguments) and hand them off to the shared buffered
+    // printf emitter.
+    SmallVector<Value *, 8> Args(CI->args());
+
+    // Split the call's basic block so the hammock created by
+    // emitAMDGPUPrintfCall has a well-defined continuation.
+    BasicBlock *Tail = CI->getParent()->splitBasicBlock(CI, "printf.split");
+    BasicBlock *Head = Tail->getSinglePredecessor();
+    assert(Head && "expected split to create a single predecessor");
+    Head->getTerminator()->eraseFromParent();
+
+    IRBuilder<> Builder(Head);
+    Builder.SetCurrentDebugLocation(CI->getDebugLoc());
+    Value *Result = emitAMDGPUPrintfCall(Builder, Args, /*IsBuffered=*/true);
+    Builder.CreateBr(Tail);
+
+    if (!CI->use_empty())
+      CI->replaceAllUsesWith(Result);
+  }
+
+  for (auto *CI : Printfs)
+    CI->eraseFromParent();
+
+  Printfs.clear();
+  return true;
+}
+
 bool AMDGPUPrintfRuntimeBindingImpl::run(Module &M) {
   auto *PrintfFunction = M.getFunction("printf");
   if (!PrintfFunction || !PrintfFunction->isDeclaration() ||
@@ -452,6 +491,16 @@ bool AMDGPUPrintfRuntimeBindingImpl::run(Module &M) {
 
   TD = &M.getDataLayout();
 
+  // When the module requests the buffered printf scheme (e.g. SYCL device code
+  // targeting the HIP runtime), use the buffered emitter which produces the
+  // print buffer layout expected by the HIP/OpenCL runtime and correctly
+  // handles %s once the format string is a resolvable constant.
+  if (auto *MD = M.getModuleFlag("amdgpu_printf_kind")) {
+    if (auto *MDS = dyn_cast<MDString>(MD))
+      if (MDS->getString() == "buffered")
+        return lowerPrintfForGpuBuffered(M);
+  }
+
   return lowerPrintfForGpu(M);
 }
 

sycl/test-e2e/DeviceLib/built-ins/printf.cpp

@@ -1,6 +1,3 @@
-// UNSUPPORTED: target-amd
-// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/22300
-// HIP doesn't support printf.
 // CUDA doesn't support vector format specifiers ("%v").
 // XFAIL: run-mode && spirv-backend
 // XFAIL-TRACKER: https://github.com/intel/llvm/issues/21618

sycl/test-e2e/Printf/char.cpp

@@ -4,8 +4,6 @@
 // The test is written using conversion specifiers table from cppreference [1]
 // [1]: https://en.cppreference.com/w/cpp/io/c/fprintf
 //
-// UNSUPPORTED: target-amd
-// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/22300
 //
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out | FileCheck %s

sycl/test-e2e/Printf/double.cpp

@@ -5,8 +5,6 @@
 // [1]: https://en.cppreference.com/w/cpp/io/c/fprintf
 //
 // REQUIRES: aspect-fp64
-// UNSUPPORTED: target-amd
-// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/22300
 //
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out | FileCheck %s

sycl/test-e2e/Printf/float.cpp

@@ -4,8 +4,6 @@
 // The test is written using conversion specifiers table from cppreference [1]
 // [1]: https://en.cppreference.com/w/cpp/io/c/fprintf
 //
-// UNSUPPORTED: target-amd
-// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/22300
 //
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out | FileCheck %s

sycl/test-e2e/Printf/int.cpp

@@ -4,8 +4,6 @@
 // The test is written using conversion specifiers table from cppreference [1]
 // [1]: https://en.cppreference.com/w/cpp/io/c/fprintf
 //
-// UNSUPPORTED: target-amd
-// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/22300
 // XFAIL: spirv-backend && gpu && run-mode
 // XFAIL-TRACKER: https://github.com/intel/llvm/issues/21314
 // RUN: %{build} -o %t.out

sycl/test-e2e/Printf/mixed-address-space.cpp

@@ -1,7 +1,6 @@
 // This test is written with an aim to check that experimental::printf versions
 // for constant and generic address space can be used in the same module.
 //
-// UNSUPPORTED: target-amd
 // XFAIL: cuda && windows
 // XFAIL-TRACKER: https://github.com/intel/llvm/issues/14733
 // FIXME: Drop the test once generic AS support is considered stable and the

sycl/test-e2e/Printf/percent-symbol.cpp

@@ -4,8 +4,6 @@
 // The test is written using conversion specifiers table from cppreference [1]
 // [1]: https://en.cppreference.com/w/cpp/io/c/fprintf
 //
-// UNSUPPORTED: target-amd
-// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/22300
 // XFAIL: cuda && windows
 // XFAIL-TRACKER: https://github.com/intel/llvm/issues/14733
 // RUN: %{build} -o %t.out