[GUI] [linux] Support "ti.GUI(fast_gui=True)" for zero-copy while gui.set_image(img)

docs/gui.rst

@@ -17,6 +17,7 @@ Create a window
     :parameter background_color: (optional, RGB hex) background color of the window
     :parameter show_gui: (optional, bool) see the note below
     :parameter fullscreen: (optional, bool) ``True`` for fullscreen window
+    :parameter fast_gui: (optional, bool) see :ref:`fast_gui`
     :return: (GUI) an object represents the window
 
     Create a window.
@@ -558,3 +559,44 @@ Image I/O
     If ``h`` is not specified, it will be equal to ``w`` by default.
 
     The output image shape is: ``(w, h, *img.shape[2:])``.
+
+
+.. _fast_gui:
+
+Zero-copying frame buffer
+-------------------------
+
+Sometimes when the GUI resolution (window size) is large, we find it impossible
+to reach 60 FPS even without any kernel invocations between each frame.
+
+This is mainly due to the copy overhead when Taichi GUI is copying image buffer
+from a place to another place, to make high-level painting APIs like
+``gui.circles`` functional. The larger the image, the larger the overhead.
+
+However, in some cases we only need ``gui.set_image`` alone. Then we may turn
+on the ``fast_gui`` mode for better performance.
+
+It will directly write the image specified in ``gui.set_image`` to frame buffer
+without hesitation, results in a much better FPS when resolution is huge.
+
+To do so, simply initialize your GUI with ``fast_gui=True``:
+
+.. code-block:: python
+
+   gui = ti.GUI(res, title, fast_gui=True)
+
+
+.. note::
+
+   This option is recommended to be used together with ``fullscreen=True``.
+
+
+.. warning::
+
+    Despite the performance boost, it has many limitations as trade off:
+
+    ``gui.set_image`` is the only available paint API in this mode. **All other
+    APIs like ``gui.circles``, ``gui.rect``, ``gui.triangles``, etc., won't work**.
+
+    ``gui.set_image`` will only takes Taichi 3D or 4D vector fields (RGB or RGBA)
+    as input.

examples/fullscreen.py

@@ -0,0 +1,21 @@
+import taichi as ti
+import numpy as np
+
+ti.init(ti.gpu)
+
+res = (1920, 1080)
+img = ti.Vector.field(3, float, res)
+
+
+@ti.kernel
+def render(t: float):
+    for i, j in img:
+        a = ti.Vector([i / res[0], j / res[1] + 2, i / res[0] + 4])
+        img[i, j] = ti.cos(a + t) * 0.5 + 0.5
+
+
+gui = ti.GUI('UV', res, fullscreen=True, fast_gui=True)
+while not gui.get_event(ti.GUI.ESCAPE):
+    render(gui.frame * 0.04)
+    gui.set_image(img)
+    gui.show()

python/taichi/lang/meta.py

@@ -21,6 +21,16 @@ def cook_image_type(x):
     return x
 
 
+@ti.kernel
+def vector_to_fast_image(img: ti.template(), out: ti.ext_arr()):
+    # FIXME: Why is ``for i, j in img:`` slower than:
+    for i, j in ti.ndrange(*img.shape):
+        u, v, w = min(255, max(0, int(img[i, img.shape[1] - 1 - j] * 255)))
+        # We use i32 for |out| since OpenGL and Metal doesn't support u8 types
+        # TODO: treat Cocoa and Big-endian machines, with XOR logic
+        out[j * img.shape[0] + i] = w + (v << 8) + (u << 16)
+
+
 @ti.kernel
 def tensor_to_image(tensor: ti.template(), arr: ti.ext_arr()):
     for I in ti.grouped(tensor):

python/taichi/misc/gui.py

@@ -40,19 +40,31 @@ def __init__(self,
                  res=512,
                  background_color=0x0,
                  show_gui=True,
-                 fullscreen=False):
+                 fullscreen=False,
+                 fast_gui=False):
         if 'TI_GUI_SHOW' in os.environ:
             show_gui = bool(int(os.environ['TI_GUI_SHOW']))
         if 'TI_GUI_FULLSCREEN' in os.environ:
             fullscreen = bool(int(os.environ['TI_GUI_FULLSCREEN']))
+        if 'TI_GUI_FAST' in os.environ:
+            fast_gui = bool(int(os.environ['TI_GUI_FAST']))
 
         self.name = name
         if isinstance(res, numbers.Number):
             res = (res, res)
         self.res = res
-        # The GUI canvas uses RGBA for storage, therefore we need NxMx4 for an image.
-        self.img = np.ascontiguousarray(np.zeros(self.res + (4, ), np.float32))
-        self.core = ti_core.GUI(name, core_veci(*res), show_gui, fullscreen)
+        self.fast_gui = fast_gui
+        if fast_gui:
+            self.img = np.ascontiguousarray(
+                np.zeros(self.res[0] * self.res[1], dtype=np.uint32))
+            fast_buf = self.img.ctypes.data
+        else:
+            # The GUI canvas uses RGBA for storage, therefore we need NxMx4 for an image.
+            self.img = np.ascontiguousarray(
+                np.zeros(self.res + (4, ), np.float32))
+            fast_buf = 0
+        self.core = ti_core.GUI(name, core_veci(*res), show_gui, fullscreen,
+                                fast_gui, fast_buf)
         self.canvas = self.core.get_canvas()
         self.background_color = background_color
         self.key_pressed = set()
@@ -143,6 +155,18 @@ def set_image(self, img):
         import numpy as np
         import taichi as ti
 
+        if self.fast_gui:
+            assert isinstance(img, ti.Matrix), \
+                    "Only ti.Vector.field is supported in GUI.set_image when fast_gui=True"
+            assert img.shape == self.res, \
+                    "Image resolution does not match GUI resolution"
+            assert img.n in [3, 4] and img.m == 1, \
+                    "Only RGB or RGBA images are supported in GUI.set_image when fast_gui=True"
+
+            from taichi.lang.meta import vector_to_fast_image
+            vector_to_fast_image(img, self.img)
+            return
+
         if isinstance(img, ti.Expr):
             if ti.core.is_integral(img.dtype) or len(img.shape) != 2:
                 # Images of uint is not optimized by xxx_to_image
@@ -196,7 +220,7 @@ def circles(self, pos, color=0xFFFFFF, radius=1):
 
         if isinstance(color, np.ndarray):
             assert color.shape == (n, )
-            color = np.ascontiguousarray(color.astype(np.uint32))
+            color = np.ascontiguousarray(color.astype(np.int32))
             color_array = int(color.ctypes.data)
             color_single = 0
         elif isinstance(color, int):

taichi/gui/gui.h

@@ -506,6 +506,8 @@ class GUI : public GUIBase {
   std::vector<std::unique_ptr<float>> widget_values;
   bool show_gui;
   bool fullscreen;
+  bool fast_gui;
+  size_t fast_buf;
 
   void set_mouse_pos(int x, int y) {
     cursor_pos = Vector2i(x, y);
@@ -763,13 +765,17 @@ class GUI : public GUIBase {
                int height = 800,
                bool show_gui = true,
                bool fullscreen = true,
+               bool fast_gui = false,
+               size_t fast_buf = 0,
                bool normalized_coord = true)
       : window_name(window_name),
         width(width),
         height(height),
         key_pressed(false),
         show_gui(show_gui),
-        fullscreen(fullscreen) {
+        fullscreen(fullscreen),
+        fast_gui(fast_gui),
+        fast_buf(fast_buf) {
     memset(button_status, 0, sizeof(button_status));
     start_time = taichi::Time::get_time();
     buffer.initialize(Vector2i(width, height));
@@ -788,12 +794,16 @@ class GUI : public GUIBase {
                Vector2i res,
                bool show_gui,
                bool fullscreen = true,
+               bool fast_gui = false,
+               size_t fast_buf = 0,
                bool normalized_coord = true)
       : GUI(window_name,
             res[0],
             res[1],
             show_gui,
             fullscreen,
+            fast_gui,
+            fast_buf,
             normalized_coord) {
   }
 

taichi/gui/x11.cpp

@@ -20,7 +20,9 @@ class CXImage {
  public:
   XImage *image;
   std::vector<uint8> image_data;
+  void *fast_data{nullptr};
   int width, height;
+
   CXImage(Display *display, Visual *visual, int width, int height)
       : width(width), height(height) {
     image_data.resize(width * height * 4);
@@ -29,6 +31,17 @@ class CXImage {
     TI_ASSERT((void *)image->data == image_data.data());
   }
 
+  CXImage(Display *display,
+          Visual *visual,
+          void *fast_data,
+          int width,
+          int height)
+      : width(width), height(height) {
+    image = XCreateImage(display, visual, 24, ZPixmap, 0, (char *)fast_data,
+                         width, height, 32, 0);
+    TI_ASSERT((void *)image->data == fast_data);
+  }
+
   void set_data(const Array2D<Vector4> &color) {
     auto p = image_data.data();
     for (int j = 0; j < height; j++) {
@@ -164,11 +177,16 @@ void GUI::create_window() {
   XSetWMProtocols((Display *)display, window, (Atom *)wmDeleteMessage.data(),
                   1);
   XMapWindow((Display *)display, window);
-  img = new CXImage((Display *)display, (Visual *)visual, width, height);
+  if (!fast_gui)
+    img = new CXImage((Display *)display, (Visual *)visual, width, height);
+  else
+    img = new CXImage((Display *)display, (Visual *)visual, (void *)fast_buf,
+                      width, height);
 }
 
 void GUI::redraw() {
-  img->set_data(buffer);
+  if (!fast_gui)
+    img->set_data(buffer);
   XPutImage((Display *)display, window, DefaultGC(display, 0), img->image, 0, 0,
             0, 0, width, height);
 }

taichi/python/export_visual.cpp

@@ -25,7 +25,7 @@ void export_visual(py::module &m) {
       .value("Press", Type::press)
       .value("Release", Type::release);
   py::class_<GUI>(m, "GUI")
-      .def(py::init<std::string, Vector2i, bool, bool>())
+      .def(py::init<std::string, Vector2i, bool, bool, bool, size_t>())
       .def_readwrite("frame_delta_limit", &GUI::frame_delta_limit)
       .def_readwrite("should_close", &GUI::should_close)
       .def("get_canvas", &GUI::get_canvas, py::return_value_policy::reference)