[Nouveau] Towards OpenCL/Cuda Support, Part 2
What Happened in Six Months?

My initial and last post on the subject was almost half a year ago, day for day. But this silence didn't meant the project had stalled and I did make some progress! (Well, not as much as what I initially hoped, to be honest.) Somewhere around mid-September or mid-October I did manage to get the following kernel (from Tom Stellar's OpenCL examples) to run witout Nouveau complaining and to return the expected value:

__kernel void pi(__global float* out)
  out[0] = 3.14159f;

Later in November, I added some additionnal SPIR-V support along with get_local_id(), get_local_size() and their global counterpart. All of those progress were made on Tesla cards, since Fermi+ handle input parameters differently and Kepler+ only uses 64-bit pointer rather than 32-bit ones. Support for Fermi+ was hacked at multiple occasions, but never ended in a working state (nor were the last round of patches tested :-D).

Then came end of November, and with it the release by Khronos of SPIR-V's final spec and of SPIR-V Tools and SPIR-V LLVM. Those releases marked the end of writing and reading SPIR-V binaries by hand! \o/ (And the beginning of producing correct SPIR-V binaries. :-D) It also marked the point when I started rewriting all my SPIR-V related work to get rid of my two-pass translation (first generate some structures out of the SPIR-V code, and then iterate over those structures to generate the corresponding NV50 IR) and my use of dynamic_cast<>(). This rewriting work is still ongoing and hasn't caught up yet with the previous branch (still missing support for get_local_id() & co.), but it should have support for Tesla+, has almost support for vectors, arrays and structures; it can be found on the branch spirv_1.0 of rMESA.

What is being worked on?

  • I am hoping to soon finish support for vectors and arrays and to get get_local_id() & co. running again.
  • Plug the SPIR-V generator (a modified version of clang and LLVM) into Mesa to directly compile the OpenCL code to SPIR-V without having to do that manually and load the resulting file from Mesa. I am having some issues to compile clover with clang 3.6.1 (same code works fine with clang from SVN), but hopefully they will soon be solved.

What are the next steps later on?

  • Beginning to work on control flow: if statements, loops and so on. This will probably take a long time since the generated NV50 IR needs to be out of SSA, and SPIR-V is in SSA (SSA = Static Single Assignment, I'll talk more about it once I know more about it, but basically each variable in SSA form is only assigned once for the whole program).
  • Add support for textures and surfaces.
  • Add support for atomics.
  • Figure out what is needed for OpenCL 1.0 support, what are the frequently used OpenCL features that should be supported.

Timeframe (or the lack of it)

I can't give any timeframe since I'm still discovering many new things and don't know how much time I will be able to allocate to this project, but I'm pretty sure I can say: do not expect any experimental OpenCL 1.0 support before summer 2016!

Regarding initial submission for upstreaming, this one is even harder to predict. I won't be able to upstream anything before the SPIR-V generator has been upstreamed as well, or has even been released, which will definitely take quite some time but I have no idea how much.


Hans de Goede and I will give a talk at FOSDEM 2016 (end of January) about OpenCL for Nouveau. He will be talking about his work on translating LLVM IR to TGSI, whereas I will be talking about translating SPIR-V to NV50 IR. Hopefully we will have some nice progress to report.

Written by pmoreau on Dec 30 2015, 4:35 PM.

Event Timeline