[OpenGL] Basic OpenGL backend #492

[archibate]

Related issue id = #492

[archibate]

I could't wait to share you more about the work!

[yuanming-hu]

Great! I have one small suggestion (which will make PR/code maintenance much easier): keep each PR small. This will also make your life easier since you have fewer conflicts and get quicker reviews. The PR process of the Metal backend #396 is a great example.

I'm super excited to witness the birth of a new backend! At the same time, we should also pay more attention to code maintainability, as we are getting more and more backends :-) Maybe we should consider extracting the common parts of the OpenGL and Metal (and potentially OpenCL in the future) codegen into a base class.

[k-ye]

I can share some experiences when working on the Metal one, or any other big features. I think it’s better to have a working solution in your local repo before opening PRs. The local solution doesn’t have to be clean or anything, just a proof of concept. Then we can start breaking down the solution into smaller PRs, possibly doing some cleanups during this process. This way you can save a lot of the reviewer’s bandwidth :-) We also don’t run into the risk that, in case there is a major design flaw, the main codebase gets polluted by incomplete solutions.

[yuanming-hu]

Yeah, I strongly agree on this two-stage strategy: writing a new backend inevitably has a lot of uncertainty, so it's great to have a minimal working prototype (e.g. one that runs mpm99.py, which is already pretty challenging IMO). Then clean up the code and break down the changes into small (and preferably testable, if I'm not asking for too much) PRs, each with clear meaning.

As @k-ye said, doing so allows you to get quick design feedback and thereby minimizes the risk of high-level design errors. It also minimizes the number of hours you spend before you get a working codegen & demos (and everybody gets super excited so you get more helping hands).

Last but not least, following this strategy makes the reviewer's job much easier :-)

[archibate]

Thank for your advises! Managed to use git rebase to reduce reviewer's pressure.

[archibate]

Btw, I found https://github.com/9ballsyndrome/WebGL_Compute_shader!

[archibate]

How to deal with globals? How can I get raw data from snodes? I can pass in/out data to/from GL now.

[archibate]

You probably want to tell me here, so I'll see it tomorrow. Good night~~

[yuanming-hu]

We need to allocate something like a pixel buffer (or something related, my OpenGL knowledge is rusty), which we use as a byte array to store the global data structure :-)

[yuanming-hu]

Btw, I found https://github.com/9ballsyndrome/WebGL_Compute_shader!

It will be great if we can do compute shader on WebGL, but my read is that it might take a couple more years for WebGL compute shader to become mature :-) For example, it's still not yet supported by default in my browser.

[archibate]

I need a deeper understand of snode now. When we say

x = ti.var(ti.i32, shape=())
@ti.kernel
def func():
  x[None] = 233

x isn't allocated this time until kernel called, right?
Then, does the backend responsible for allocation?
In fact, SSBOs aren't so precious, they don't need to be allocated like cudaMallocManaged this way. Simply malloc can works - any part of CPU memory can be bound with SSBO.

[yuanming-hu]

x isn't allocated this time until kernel called, right?

Yes, it will be allocated the first time one kernel is called, or when you access the data structure in Python scope (e.g. x[1, 2, 3]=3 outside Taichi kernels)

Then, does the backend responsible for allocation?

The struct compiler does allocation:

taichi/taichi/backends/struct_llvm.cpp

Line 292 in 46490e7

creator = [=]() {

The struct_metal file also worth checking out.

[archibate]

or when you access the data structure in Python scope (e.g. x[1, 2, 3]=3 outside Taichi kernels)

I also noticed that x[1] = 2 or print(x[1]) is translated into a kernel with a argument somehow...

The struct_metal file also worth checking out.

Oh, I almost forget that!

taichi/taichi/backends/struct_llvm.cpp

Line 300 in 46490e7

prog->memory_pool->set_queue((MemRequestQueue *)mem_req_queue);

Seems you have a memory pool there working for LLVM backend? Maybe GL backend can use it later too, as lone as it doesn't use CUDA API.

[yuanming-hu]

I also noticed that x[1] = 2 or print(x[1]) is translated into a kernel with an argument somehow...

Yeah, you'll have to launch a compute shader task for single element accesses like these. A more efficient way is to use numpy interfaces to/from_numpy() - I would suggest start with the numpy interface first since that's easier to implement.

Seems you have a memory pool there working for LLVM backend? Maybe GL backend can use it later too, as lone as it doesn't use CUDA API.

The memory pool stuff relies on unified memory (which to my knowledge is not supported by OpenGL), so I would suggest starting with just dense snodes, whose storage can be directly pre-allocated.

[archibate]

to_numpy need struct_for kernel, which is not implemented for now.

[archibate]

Memory allocate done. Next event is to figure out how x[0] (read_int) works, wanna see print shows 233 on screen.

[archibate]

We're making history! After this and stage 2: struct_for_kernel & parallel process done, I think we're going to release our GPGPU programming language on both Linux, OS X and Windows with a even wider usersphere!!! I've also thought about the stage 3 title: nested for & func~

[archibate]

Ok! Will fix the bug tomorrow. Check examples/opengl_example.py for test, and good night.

[archibate]

NOTE: 976d387 may have conflicts with #527.

[archibate]

TODO: atan2 -> atan.

       genType atan(genType y, genType x);
       genType atan(genType y_over_x);

[archibate]

I simply skipped OpenGL from test_types for now, since it doesn't support lots of types involved in test_types like u8. All types supported by GLSL firmware: i1(bool), i32, u32, f32, f64.

[k-ye]

I saw that you requested a review. If you think this is ready, as suggested before, could you break this down into smaller PRs? Otherwise it's very hard for reviewing.. You may need to create new branches and open new PRs, whereas archibate:opengl will be an experimental branch containing a working solution for reference ;)

[archibate]

Yes, I will, after everything done. But the actual reason why I requested you for review is, I'm afraid my changes in test_types.py could broke yours.

[archibate]

ALL TESTS PASSED!!!

[k-ye]

I'm afraid my changes in test_types.py could broke yours.

Ha thanks. Don't worry too much until it actually gets merged... And I think it's fine not to worry too much about breaking the metal backend. After all, it's hard to test the code if you don't even have a way to run it..

[archibate]

Break down into fragments, closing here.