Index: libclc/amdgpu/lib/SOURCES
===================================================================
--- libclc/amdgpu/lib/SOURCES
+++ libclc/amdgpu/lib/SOURCES
@@ -12,3 +12,4 @@
 math/half_sqrt.cl
 math/nextafter.cl
 math/sqrt.cl
+misc/printf.cl
Index: libclc/amdgpu/lib/misc/printf.cl
===================================================================
--- /dev/null
+++ libclc/amdgpu/lib/misc/printf.cl
@@ -0,0 +1,29 @@
+
+// AMD target compiler replace printf call
+int printf(__constant const char* st, ...) __attribute__((format(printf, 1, 2)));
+
+//reserves space to the printf data buffer and returns a pointer to it
+__global char *
+__printf_alloc(uint bytes) {
+    __global char *ptr = (__global char *)(((__constant size_t *)__builtin_amdgcn_implicitarg_ptr())[2]);
+    if (!ptr)
+       return NULL;
+
+    // A printf buffer store the current offset and its size using 2 uint
+    // It's expected that the runtime store a initial offset of 8 (sizeof(cl_uint) * 2)
+    uint buffer_offset = ((__global uint*)ptr)[0];
+    uint buffer_size = ((__global uint*)ptr)[1];
+
+    for (;;) {
+        if (buffer_offset + bytes > buffer_size)
+            break;
+
+        buffer_offset = atomic_add((__global uint*)ptr, bytes);
+        if (buffer_offset + bytes <= buffer_size)
+            return ptr + buffer_offset;
+        else
+           atomic_sub((__global uint*)ptr, bytes);
+    }
+
+    return NULL;
+}