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 + +// 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 #include +/* 6.12.13 Printf */ +#include + /* 6.11.13 Image Read and Write Functions */ #include #include 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