diff --git a/docs/gui.rst b/docs/gui.rst index b45c6b53610f4..cd5783d0de674 100644 --- a/docs/gui.rst +++ b/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:: + + If possible, consider enabling this option, especially when ``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. diff --git a/examples/fullscreen.py b/examples/fullscreen.py new file mode 100644 index 0000000000000..d13f1eab36daa --- /dev/null +++ b/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() diff --git a/misc/demo_fullscreen.py b/misc/demo_fullscreen.py deleted file mode 100644 index f29d4a15ac538..0000000000000 --- a/misc/demo_fullscreen.py +++ /dev/null @@ -1,22 +0,0 @@ -import taichi as ti -import numpy as np - -ti.init() - -res = (1920, 1080) -pixels = ti.Vector.field(3, dtype=float, shape=res) - - -@ti.kernel -def paint(): - for i, j in pixels: - u = i / res[0] - v = j / res[1] - pixels[i, j] = [u, v, 0] - - -gui = ti.GUI('UV', res, fullscreen=True) -while not gui.get_event(ti.GUI.ESCAPE): - paint() - gui.set_image(pixels) - gui.show() diff --git a/python/taichi/lang/meta.py b/python/taichi/lang/meta.py index 243c27153f8bc..1fc85986f0333 100644 --- a/python/taichi/lang/meta.py +++ b/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): diff --git a/python/taichi/misc/gui.py b/python/taichi/misc/gui.py index 5945cf0ae2047..f12196145bb19 100644 --- a/python/taichi/misc/gui.py +++ b/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 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 diff --git a/taichi/gui/gui.h b/taichi/gui/gui.h index 5885f5464027a..9ff4e0fa337ce 100644 --- a/taichi/gui/gui.h +++ b/taichi/gui/gui.h @@ -506,6 +506,8 @@ class GUI : public GUIBase { std::vector> widget_values; bool show_gui; bool fullscreen; + bool fast_gui; + uintptr_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, + uintptr_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, + uintptr_t fast_buf = 0, bool normalized_coord = true) : GUI(window_name, res[0], res[1], show_gui, fullscreen, + fast_gui, + fast_buf, normalized_coord) { } diff --git a/taichi/gui/x11.cpp b/taichi/gui/x11.cpp index 5f221c6a66b29..15b5ebf401167 100644 --- a/taichi/gui/x11.cpp +++ b/taichi/gui/x11.cpp @@ -20,7 +20,9 @@ class CXImage { public: XImage *image; std::vector 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 &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); } diff --git a/taichi/python/export_visual.cpp b/taichi/python/export_visual.cpp index f621dd412e4b0..ae959d83371c7 100644 --- a/taichi/python/export_visual.cpp +++ b/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_(m, "GUI") - .def(py::init()) + .def(py::init()) .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)