<div class="gmail_quote">On Thu, Oct 20, 2011 at 1:42 PM, Denis Steckelmacher <span dir="ltr"><<a href="mailto:steckdenis@yahoo.fr">steckdenis@yahoo.fr</a>></span> wrote:<br><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex;">
Hello,<br>
<br>
I am the developer of Clover, and so much activity about OpenCL these days is really exciting. Here is my point of view, mainly on Clover and how the projects could use each other.<br>
<br>
Clover is made in a way that allow a certain level of modularity. Although POCL would be very difficult to merge into Clover (or Clover into POCL), as these two projects are nearly exactly doing the same things (an OpenCL platform layer), Libclc and the German driver are very interesting.<br>
<br>
I'll begin with the German driver, maybe the project the easiest to integrate or bind to Clover. Clover uses a hardware abstraction layer, a set of virtual classes a driver has to implement in order to be usable by Clover. They are DeviceInterface, DeviceBuffer, DeviceKernel and DeviceProgram. You can see their prototype here : <a href="http://cgit.freedesktop.org/~steckdenis/clover/tree/src/core/deviceinterface.h" target="_blank">http://cgit.freedesktop.org/~steckdenis/clover/tree/src/core/deviceinterface.h</a> . Currently, I have developed a small driver, normally not too slow, that simply runs the LLVM IR produced by Clang using the LLVM JIT. The kernels are split in work-groups, split in work-items in such a way that multithreading is efficiently used.<br>
<br>
What would be interesting is to try to integrate the German driver into Clover using this interface, or to have this driver built as a library on which Clover links (if there is a problem of license, but Clover is BSD and it seems that the driver will be the same). I would personally be very excited to see how another driver would perform in Clover, feature-wise and performance-wise.<br>
<br>
Libclc could also be useful to Clover, but less likely. The goal of this project is to implement all the OpenCL built-in functions. It's good, but Clover already does the same, using a different technique. Libclc is very elegant (I think), it seems to use custom LLVM intrinsics, and is built around pure C macros. </blockquote>
<div><br></div><div>libclc only uses LLVM intrinsics (currently) for back-end specific functionality. For example, the get_local_id() function is implemented separately for each target, and uses LLVM PTX intrinsics if compiling for the PTX back-end. This is not something you could implement in a generic way without back-end hooks (at least not without dirty hacks in the back-end).</div>
<div> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex;">Clover uses a slightly more complex system, involving a Python script "compiling" a set of built-ins into four files. For example, this declaration (REPL is a macro that does a simple for()) :<br>
</blockquote><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex;"><br>
----<br>
def vecf : float2 float3 float4 float8 float16<br>
<br>
native $type acospi $vecf : x:$type<br>
REPL($vecdim)<br>
result[i] = std::acos(x[i]) / M_PI;<br>
end<br>
----<br>
<br>
Is compiled to these fragments, one for each vector type (float2, float3, etc) :<br>
<br>
----<br>
// In stdlib_def.h : what the OpenCL C kernel sees<br>
float2 OVERLOAD acospi(float2 x);<br>
<br>
// In stdlib_impl.h : what gets compiled to LLVM IR at Clover compile time, and then linked to each kernel<br>
void __cpu_float2_acospi_float2(float *result, float *x);<br>
float2 OVERLOAD acospi(float2 x)<br>
{<br>
float2 result;<br>
<br>
__cpu_float2_acospi_float2((float *)&result, (float *)&x);<br>
<br>
return result;<br>
}<br>
<br>
// __cpu_float2_acospi_float2 is a function implemented in the Clover .so library, using llvm::JIT::registerLasyFunctionCreator<br>
// In builtins_impl.h : the actual C++ implementation, included in src/core/cpu/builtins.cpp<br>
static void float2_acospi_float2(float *result, float *x)<br>
{<br>
REPL(2)<br>
result[i] = std::acos(x[i]) / M_PI;<br>
<br>
}<br>
<br>
// And then a small else if in the lazy function creator, in order to bind everything together<br>
else if (name == "__cpu_float2_acospi_float2")<br>
return (void *)&float2_acospi_float2;<br>
----<br></blockquote><div><br></div><div>If the LLVM JIT picks up these functions at run-time, then there is no chance of inlining these math functions. This is not good for performance.</div><div> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex;">
<br>
The system works fairly well, and I was able to implement a dozen of built-in functions in only two hours. It's very fast to simply declare "native" functions using STL or Boost math functions, and hardware drivers simply can replace the LLVM "call" statements with what they need to accelerate the functions on the GPU.<br>
</blockquote><div><br></div><div>But then the hardware driver layer has to have GPU implementations for all of these functions.</div><div> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex;">
<br>
So, libclc would only be useful to Clover if it is developed by so much people that its development becomes way faster than Clover, and if it provides an easy and efficient way to natively implement functions, without needing to have a LLVM pass turning LLVM intrinsics to native function calls.<br>
<br>
Here is my personal point of view, and I hope a solution will be found not to have three or four different projects working on the same things.<br>
<br>
Best regards,<br>
<font color="#888888">Denis Steckelmacher.<br>
</font><div><div></div><div class="h5"><br>
_______________________________________________<br>
LLVM Developers mailing list<br>
<a href="mailto:LLVMdev@cs.uiuc.edu">LLVMdev@cs.uiuc.edu</a> <a href="http://llvm.cs.uiuc.edu" target="_blank">http://llvm.cs.uiuc.edu</a><br>
<a href="http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev" target="_blank">http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev</a><br>
</div></div></blockquote></div><br><br clear="all"><div><br></div>-- <br><br><div>Thanks,</div><div><br></div><div>Justin Holewinski</div><br>