Support for printf

[Kerilk]

Hello,

I am trying to add support for missing HIP intrisics to HIPCL (see here for a WIP https://github.com/Kerilk/hipcl/tree/fence).
I encountered an issue when trying to add support for printf:

diff --git a/include/hip/hipcl.hh b/include/hip/hipcl.hh
index 16f9355..4b59078 100644
--- a/include/hip/hipcl.hh
+++ b/include/hip/hipcl.hh
@@ -65,6 +65,8 @@ typedef int hipLaunchParm;
 #pragma push_macro("__DEVICE__")
 #define __DEVICE__ static __device__ __forceinline__
 
+extern "C" __device__ int printf(const char * format, ...);
+
 extern "C" __device__ size_t _Z12get_local_idj(uint);
 __DEVICE__ uint __hip_get_thread_idx_x() { return _Z12get_local_idj(0); }
 __DEVICE__ uint __hip_get_thread_idx_y() { return _Z12get_local_idj(1); }

This "almost" works. The only issue is being unable to declare the function as:

extern "C" __device__ int printf(__constant__ char * format, ...);

as pointers in function arguments cannot be in the constant address space. This causes the OpenCL kernel build to fail due to the printf argument format usually being in the constant address space (as it should), but the format then needs to be cast to the generic address space which is forbidden.

I am no expert in llvm pass and how they are written, but I think I see a couple ways to get around this:

• define a new address space (__cl_constant__) in hip for the printf format causing the OpenCL counterpart to be decorated with the constant address space
• add a quirk for printf automatically adding the constant address space to it's first argument. There are already several quirks for printf support in the various part of the tool chain, as it is the only allowed var_arg function.

What would seem to you the most reasonable (maybe a different one of course) approach to solve this issue?

Thanks