Index: libclc/CMakeLists.txt
===================================================================
--- libclc/CMakeLists.txt
+++ libclc/CMakeLists.txt
@@ -261,8 +261,8 @@
 			"generic/include" )
 		target_compile_definitions( builtins.link.${arch_suffix} PRIVATE
 			"__CLC_INTERNAL" )
-		target_compile_options( builtins.link.${arch_suffix} PRIVATE  -target
-			${t} ${mcpu} -fno-builtin -nostdlib )
+		target_compile_options( builtins.link.${arch_suffix} PRIVATE  -cl-std=CL1.2
+			 -target ${t} ${mcpu} -fno-builtin -nostdlib )
 		set_target_properties( builtins.link.${arch_suffix} PROPERTIES
 			LINKER_LANGUAGE CLC )
 
Index: libclc/amdgcn/lib/SOURCES
===================================================================
--- libclc/amdgcn/lib/SOURCES
+++ libclc/amdgcn/lib/SOURCES
@@ -4,6 +4,7 @@
 math/fmin.cl
 math/ldexp.cl
 mem_fence/fence.cl
+printf/printf_alloc.cl
 synchronization/barrier.cl
 workitem/get_global_offset.cl
 workitem/get_group_id.cl
Index: libclc/amdgcn/lib/printf/printf_alloc.cl
===================================================================
--- /dev/null
+++ libclc/amdgcn/lib/printf/printf_alloc.cl
@@ -0,0 +1,32 @@
+#include <clc/clc.h>
+
+// AMD target compiler replaces printf calls and adds a call to __printf_alloc
+
+//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 stores 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)
+    // On each call the driver store an id (cl_uint) linked to a format litteral and the printf args values
+    // | buffer_offset | buffer_size | (format id | value, ...) ... (format id | value, ...) |
+    uint buffer_offset = ((__global uint*)ptr)[0];
+    uint buffer_size = ((__global uint*)ptr)[1];
+
+    // The function return the buffer ptr to write to (ptr + offset)
+    // and increment the offset to reserve the memory (offset + bytes)
+    // NULL mean that the buffer is full
+    __global uint* buffer_offset_ptr = (__global uint*)ptr;
+    if (buffer_offset + bytes <= buffer_size) {
+        buffer_offset = atomic_add(buffer_offset_ptr, bytes);
+        if (buffer_offset + bytes <= buffer_size)
+            return ptr + buffer_offset;
+        else
+           atomic_sub(buffer_offset_ptr, bytes);
+    }
+
+    return NULL;
+}
Index: libclc/generic/include/clc/clc.h
===================================================================
--- libclc/generic/include/clc/clc.h
+++ libclc/generic/include/clc/clc.h
@@ -281,6 +281,9 @@
 #include <clc/misc/shuffle.h>
 #include <clc/misc/shuffle2.h>
 
+/* 6.12.13 Printf */
+#include <clc/printf/printf.h>
+
 /* 6.11.13 Image Read and Write Functions */
 #include <clc/image/image_defines.h>
 #include <clc/image/image.h>
Index: libclc/generic/include/clc/printf/printf.h
===================================================================
--- /dev/null
+++ libclc/generic/include/clc/printf/printf.h
@@ -0,0 +1,4 @@
+#if __OPENCL_VERSION__ >= 120
+__attribute__((format(printf, 1, 2)))
+_CLC_DECL int printf(__constant const char* st, ...);
+#endif