[libclc] libclc: Add __printf_alloc implementation (PR #185231)

[Matt Arsenault via cfe-commits]

https://github.com/arsenm created https://github.com/llvm/llvm-project/pull/185231

AMDGPU OpenCL printf implementation emits a call to this helper
function.

>From 737019030edbfe51f34139d592adb318e453132c Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Sat, 7 Mar 2026 18:40:22 +0100
Subject: [PATCH] libclc: Add __printf_alloc implementation

AMDGPU OpenCL printf implementation emits a call to this helper
function.
---
 libclc/opencl/lib/amdgcn/SOURCES              |  1 +
 .../lib/amdgcn/printf/__printf_alloc.cl       | 36 +++++++++++++++++++
 2 files changed, 37 insertions(+)
 create mode 100644 libclc/opencl/lib/amdgcn/printf/__printf_alloc.cl

diff --git a/libclc/opencl/lib/amdgcn/SOURCES b/libclc/opencl/lib/amdgcn/SOURCES
index e52f54789bfab..bb94e3fe3b403 100644
--- a/libclc/opencl/lib/amdgcn/SOURCES
+++ b/libclc/opencl/lib/amdgcn/SOURCES
@@ -1,4 +1,5 @@
 mem_fence/fence.cl
+printf/__printf_alloc.cl
 subgroup/subgroup.cl
 synchronization/sub_group_barrier.cl
 workitem/get_global_offset.cl
diff --git a/libclc/opencl/lib/amdgcn/printf/__printf_alloc.cl b/libclc/opencl/lib/amdgcn/printf/__printf_alloc.cl
new file mode 100644
index 0000000000000..f4a7f46cbbcac
--- /dev/null
+++ b/libclc/opencl/lib/amdgcn/printf/__printf_alloc.cl
@@ -0,0 +1,36 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include <amdhsa_abi.h>
+
+#define OFFSET 8
+
+// Atomically reserves space to the printf data buffer and returns a pointer to
+// it
+__global char *__printf_alloc(uint bytes) {
+  __constant amdhsa_implicit_kernarg_v5 *args =
+      (__constant amdhsa_implicit_kernarg_v5 *)
+          __builtin_amdgcn_implicitarg_ptr();
+  __global char *ptr = (__global char *)args->printf_buffer;
+
+  uint size = ((__global uint *)ptr)[1];
+  uint offset = __opencl_atomic_load((__global atomic_uint *)ptr,
+                                     memory_order_relaxed, memory_scope_device);
+
+  for (;;) {
+    if (OFFSET + offset + bytes > size)
+      return NULL;
+
+    if (__opencl_atomic_compare_exchange_strong(
+            (__global atomic_uint *)ptr, &offset, offset + bytes,
+            memory_order_relaxed, memory_order_relaxed, memory_scope_device))
+      break;
+  }
+
+  return ptr + OFFSET + offset;
+}