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