taichi Documentation

Productive, portable, and performant GPU programming in Python.

28,286 stars C++ C++ 1 file ยท ~26,647 tokens #computer-graphics#differentiable-programming#gpu#gpu-programming#sparse-computation#taichi
RAW Doc

Docker & Self-Hosted Deployment Verified OCI

Deploy taichi on any server, homelab, or cloud container platform.

1-Click Terminal Command
bash (docker run) 
docker run -d --name taichi --restart always taichi-dev/taichi:latest

File: README.md





shell
pip install taichi  # Install Taichi Lang
ti gallery          # Launch demo gallery

What is Taichi Lang?

Taichi Lang is an open-source, imperative, parallel programming language for high-performance numerical computation. It is embedded in Python and uses just-in-time (JIT) compiler frameworks, for example LLVM, to offload the compute-intensive Python code to the native GPU or CPU instructions.

The language has broad applications spanning real-time physical simulation, numerical computation, augmented reality, artificial intelligence, vision and robotics, visual effects in films and games, general-purpose computing, and much more.





...More

Why Taichi Lang?

  • Built around Python: Taichi Lang shares almost the same syntax with Python, allowing you to write algorithms with minimal language barrier. It is also well integrated into the Python ecosystem, including NumPy and PyTorch.
  • Flexibility: Taichi Lang provides a set of generic data containers known as SNode (/หˆsnoสŠd/), an effective mechanism for composing hierarchical, multi-dimensional fields. This can cover many use patterns in numerical simulation (e.g. spatially sparse computing).
  • Performance: With the @ti.kernel decorator, Taichi Lang's JIT compiler automatically compiles your Python functions into efficient GPU or CPU machine code for parallel execution.
  • Portability: Write your code once and run it everywhere. Currently, Taichi Lang supports most mainstream GPU APIs, such as CUDA and Vulkan.
  • ... and many more features! A cross-platform, Vulkan-based 3D visualizer, differentiable programming, quantized computation (experimental), etc.

Getting Started

Installation

Prerequisites

  • Operating systems
    • Windows
    • Linux
    • macOS
  • Python: 3.6 ~ 3.10 (64-bit only)
  • Compute backends
    • x64/ARM CPUs
    • CUDA
    • Vulkan
    • OpenGL (4.3+)
    • Apple Metal
    • WebAssembly (experiemental)

Use Python's package installer pip to install Taichi Lang:

bash
pip install --upgrade taichi

We also provide a nightly package. Note that nightly packages may crash because they are not fully tested. We cannot guarantee their validity, and you are at your own risk trying out our latest, untested features. The nightly packages can be installed from our self-hosted PyPI (Using self-hosted PyPI allows us to provide more frequent releases over a longer period of time)

bash
pip install -i https://pypi.taichi.graphics/simple/ taichi-nightly

Run your "Hello, world!"

Here is how you can program a 2D fractal in Taichi:

py
# python/taichi/examples/simulation/fractal.py

import taichi as ti

ti.init(arch=ti.gpu)

n = 320
pixels = ti.field(dtype=float, shape=(n * 2, n))


@ti.func
def complex_sqr(z):
    return ti.Vector([z[0]**2 - z[1]**2, z[1] * z[0] * 2])


@ti.kernel
def paint(t: float):
    for i, j in pixels:  # Parallelized over all pixels
        c = ti.Vector([-0.8, ti.cos(t) * 0.2])
        z = ti.Vector([i / n - 1, j / n - 0.5]) * 2
        iterations = 0
        while z.norm() < 20 and iterations < 50:
            z = complex_sqr(z) + c
            iterations += 1
        pixels[i, j] = 1 - iterations * 0.02


gui = ti.GUI("Julia Set", res=(n * 2, n))

for i in range(1000000):
    paint(i * 0.03)
    gui.set_image(pixels)
    gui.show()

If Taichi Lang is properly installed, you should get the animation below ๐ŸŽ‰:

See Get started for more information.

Build from source

If you wish to try our experimental features or build Taichi Lang for your own environments, see Developer installation.

Documentation

Community activity [](https://repography.com)




Contributing

Kudos to all of our amazing contributors! Taichi Lang thrives through open-source. In that spirit, we welcome all kinds of contributions from the community. If you would like to participate, check out the Contribution Guidelines first.

Contributor avatars are randomly shuffled.

License

Taichi Lang is distributed under the terms of Apache License (Version 2.0).

See Apache License for details.

Community

For more information about the events or community, please refer to this page

Join our discussions

Report an issue

Contact us

Reference

Demos

AOT deployment

Lectures & talks

Citations

If you use Taichi Lang in your research, please cite the corresponding papers:

File: benchmarks/README.md

Install a few of extra requirements:

bash
python3 -m pip install -r requirements.txt

Run

To run all benchmarks:

bash
python3 run.py

Result

The benchmark results will be stored in the results folder in your current directory.
If you wish to save the results as a single json file (./results/results.json):

bash
python3 deserialize.py

Or you can specify the input and output path:

bash
python3 deserialize.py --folder PATH_OF_RESULTS_FOLDER --output_path PATH_YOU_WIHS_TO_STORE

Tools

After getting benchmark results (./results), you can use a visualization tool to profile performance problems:

bash
python3 visualization.py

You can specify the results file path:

bash
python3 visualization.py --folder PATH_OF_RESULTS_FOLDER

The default host and port is localhost:5006\visualization.
If you want to enable remote access, take the following steps:

bash
python3 visualization.py --host YOUR_IP_ADDRESS --port PORT_YOU_WISH_TO_USE

File: c_api/docs/taichi/taichi_core.h.md

Core Functionality

Taichi Core exposes all necessary interfaces for offloading the AOT modules to Taichi. The following is a list of features that are available regardless of your backend. The corresponding APIs are still under development and subject to change.

Availability

Taichi C-API intends to support the following backends:

Backend Offload Target Maintenance Tier Stabilized?
Vulkan GPU Tier 1 Yes
Metal GPU (macOS, iOS) Tier 2 No
CUDA (LLVM) GPU (NVIDIA) Tier 2 No
CPU (LLVM) CPU Tier 2 No
OpenGL GPU Tier 2 No
OpenGL ES GPU Tier 2 No
DirectX 11 GPU (Windows) N/A No

The backends with tier-1 support are being developed and tested more intensively. And most new features will be available on Vulkan first because it has the most outstanding cross-platform compatibility among all the tier-1 backends.
For the backends with tier-2 support, you should expect a delay in the fixes to minor issues.

For convenience, in the following text and other C-API documents, the term host refers to the user of the C-API; the term device refers to the logical (conceptual) compute device, to which Taichi's runtime offloads its compute tasks. A device may not be a physical discrete processor other than the CPU and the host may not be able to access the memory allocated on the device.

Unless otherwise specified, device, backend, offload target, and GPU are interchangeable; host, user code, user procedure, and CPU are interchangeable.

HowTo

The following section provides a brief introduction to the Taichi C-API.

Create and destroy a Runtime Instance

You must create a runtime instance before working with Taichi, and only one runtime per thread. Currently, we do not officially claim that multiple runtime instances can coexist in a process, but please feel free to file an issue with us if you run into any problem with runtime instance coexistence.

cpp
// Create a Taichi Runtime on Vulkan device at index 0.
TiRuntime runtime = ti_create_runtime(TI_ARCH_VULKAN, 0);

When your program runs to the end, ensure that:

  • You destroy the runtime instance,
  • All related resources are destroyed before the handle.runtime itself.
cpp
ti_destroy_runtime(runtime);

Allocate and free memory

Allocate a piece of memory that is visible only to the device. On the GPU backends, it usually means that the memory is located in the graphics memory (GRAM).

cpp
TiMemoryAllocateInfo mai {};
mai.size = 1024; // Size in bytes.
mai.usage = TI_MEMORY_USAGE_STORAGE_BIT;
TiMemory memory = ti_allocate_memory(runtime, &mai);

Allocated memory is automatically freed when the related handle.runtime is destroyed. You can also manually free the allocated memory.

cpp
ti_free_memory(runtime, memory);

Allocate host-accessible memory

By default, memory allocations are physically or conceptually local to the offload target for performance reasons. You can configure the structure.memory_allocate_info to enable host access to memory allocations. But please note that host-accessible allocations may slow down computation on GPU because of the limited bus bandwidth between the host memory and the device.

You must set host_write to definition.true to allow zero-copy data streaming to the memory.

cpp
TiMemoryAllocateInfo mai {};
mai.size = 1024; // Size in bytes.
mai.host_write = TI_TRUE;
mai.usage = TI_MEMORY_USAGE_STORAGE_BIT;
TiMemory steaming_memory = ti_allocate_memory(runtime, &mai);

// ...

std::vector src = some_random_data_source();

void* dst = ti_map_memory(runtime, steaming_memory);
std::memcpy(dst, src.data(), src.size());
ti_unmap_memory(runtime, streaming_memory);

To read data back to the host, host_read must be set to definition.true.

cpp
TiMemoryAllocateInfo mai {};
mai.size = 1024; // Size in bytes.
mai.host_read = TI_TRUE;
mai.usage = TI_MEMORY_USAGE_STORAGE_BIT;
TiMemory read_back_memory = ti_allocate_memory(runtime, &mai);

// ...

std::vector dst(1024);
void* src = ti_map_memory(runtime, read_back_memory);
std::memcpy(dst.data(), src, dst.size());
ti_unmap_memory(runtime, read_back_memory);

ti_free_memory(runtime, read_back_memory);

You can set host_read and host_write at the same time.

Load and destroy a Taichi AOT module

You can load a Taichi AOT module from the filesystem.

cpp
TiAotModule aot_module = ti_load_aot_module(runtime, "/path/to/aot/module");

/path/to/aot/module should point to the directory that contains a metadata.json.

You can destroy an unused AOT module, but please ensure that there is no kernel or compute graph related to it pending to ti_flush.

cpp
ti_destroy_aot_module(aot_module);

Launch kernels and compute graphs

You can extract kernels and compute graphs from an AOT module. Kernel and compute graphs are a part of the module, so you don't have to destroy them.

cpp
TiKernel kernel = ti_get_aot_module_kernel(aot_module, "foo");
TiComputeGraph compute_graph = ti_get_aot_module_compute_graph(aot_module, "bar");

You can launch a kernel with positional arguments. Please ensure the types, the sizes and the order matches the source code in Python.

cpp
TiNdArray ndarray{};
ndarray.memory = get_some_memory();
ndarray.shape.dim_count = 1;
ndarray.shape.dims[0] = 16;
ndarray.elem_shape.dim_count = 2;
ndarray.elem_shape.dims[0] = 4;
ndarray.elem_shape.dims[1] = 4;
ndarray.elem_type = TI_DATA_TYPE_F32;

std::array<TiArgument, 3> args{};

TiArgument& arg0 = args[0];
arg0.type = TI_ARGUMENT_TYPE_I32;
arg0.value.i32 = 123;

TiArgument& arg1 = args[1];
arg1.type = TI_ARGUMENT_TYPE_F32;
arg1.value.f32 = 123.0f;

TiArgument& arg2 = args[2];
arg2.type = TI_ARGUMENT_TYPE_NDARRAY;
arg2.value.ndarray = ndarray;

ti_launch_kernel(runtime, kernel, args.size(), args.data());

You can launch a compute graph in a similar way. But additionally please ensure the argument names matches those in the Python source.

cpp
std::array<TiNamedArgument, 3> named_args{};
TiNamedArgument& named_arg0 = named_args[0];
named_arg0.name = "foo";
named_arg0.argument = args[0];
TiNamedArgument& named_arg1 = named_args[1];
named_arg1.name = "bar";
named_arg1.argument = args[1];
TiNamedArgument& named_arg2 = named_args[2];
named_arg2.name = "baz";
named_arg2.argument = args[2];

ti_launch_compute_graph(runtime, compute_graph, named_args.size(), named_args.data());

When you have launched all kernels and compute graphs for this batch, you should function.flush and function.wait for the execution to finish.

cpp
ti_flush(runtime);
ti_wait(runtime);

WARNING This part is subject to change. We will introduce multi-queue in the future.

API Reference

alias.bool

A boolean value. Can be either definition.true or definition.false. Assignment with other values could lead to undefined behavior.

definition.true

A condition or a predicate is satisfied; a statement is valid.

definition.false

A condition or a predicate is not satisfied; a statement is invalid.

alias.flags

A bit field that can be used to represent 32 orthogonal flags. Bits unspecified in the corresponding flag enum are ignored.

Enumerations and bit-field flags in the C-API have a TI_XXX_MAX_ENUM case to ensure the enum has a 32-bit range and in-memory size. It has no semantical impact and can be safely ignored.

definition.null_handle

A sentinal invalid handle that will never be produced from a valid call to Taichi C-API.

handle.runtime

Taichi runtime represents an instance of a logical backend and its internal dynamic state. The user is responsible to synchronize any use of handle.runtime. The user must not manipulate multiple handle.runtimes in the same thread.

handle.aot_module

An ahead-of-time (AOT) compiled Taichi module, which contains a collection of kernels and compute graphs.

handle.event

A synchronization primitive to manage device execution flows in multiple queues.

handle.memory

A contiguous allocation of device memory.

handle.image

A contiguous allocation of device image.

handle.sampler

An image sampler. definition.null_handle represents a default image sampler provided by the runtime implementation. The filter modes and address modes of default samplers depend on backend implementation.

handle.kernel

A Taichi kernel that can be launched on the offload target for execution.

handle.compute_graph

A collection of Taichi kernels (a compute graph) to launch on the offload target in a predefined order.

enumeration.error

Errors reported by the Taichi C-API.

  • enumeration.error.success: The Taichi C-API invocation finished gracefully.
  • enumeration.error.not_supported: The invoked API, or the combination of parameters is not supported by the Taichi C-API.
  • enumeration.error.corrupted_data: Provided data is corrupted.
  • enumeration.error.name_not_found: Provided name does not refer to any existing item.
  • enumeration.error.invalid_argument: One or more function arguments violate constraints specified in C-API documents, or kernel arguments mismatch the kernel argument list defined in the AOT module.
  • enumeration.error.argument_null: One or more by-reference (pointer) function arguments point to null.
  • enumeration.error.argument_out_of_range: One or more function arguments are out of its acceptable range; or enumeration arguments have undefined value.
  • enumeration.error.argument_not_found: One or more kernel arguments are missing.
  • enumeration.error.invalid_interop: The intended interoperation is not possible on the current arch. For example, attempts to export a Vulkan object from a CUDA runtime are not allowed.
  • enumeration.error.invalid_state: The Taichi C-API enters an unrecoverable invalid state. Related Taichi objects are potentially corrupted. The users should release the contaminated resources for stability. Please feel free to file an issue if you encountered this error in a normal routine.
  • enumeration.error.incompatible_module: The AOT module is not compatible with the current runtime.

enumeration.arch

Types of backend archs.

  • enumeration.arch.vulkan: Vulkan GPU backend.
  • enumeration.arch.metal: Metal GPU backend.
  • enumeration.arch.cuda: NVIDIA CUDA GPU backend.
  • enumeration.arch.x64: x64 native CPU backend.
  • enumeration.arch.arm64: Arm64 native CPU backend.
  • enumeration.arch.opengl: OpenGL GPU backend.
  • enumeration.arch.gles: OpenGL ES GPU backend.

enumeration.capability

Device capabilities.

structure.capability_level_info

An integral device capability level. It currently is not guaranteed that a higher level value is compatible with a lower level value.

enumeration.data_type

Elementary (primitive) data types. There might be vendor-specific constraints on the available data types so it's recommended to use 32-bit data types if multi-platform distribution is desired.

  • enumeration.data_type.f16: 16-bit IEEE 754 half-precision floating-point number.
  • enumeration.data_type.f32: 32-bit IEEE 754 single-precision floating-point number.
  • enumeration.data_type.f64: 64-bit IEEE 754 double-precision floating-point number.
  • enumeration.data_type.i8: 8-bit one's complement signed integer.
  • enumeration.data_type.i16: 16-bit one's complement signed integer.
  • enumeration.data_type.i32: 32-bit one's complement signed integer.
  • enumeration.data_type.i64: 64-bit one's complement signed integer.
  • enumeration.data_type.u8: 8-bit unsigned integer.
  • enumeration.data_type.u16: 16-bit unsigned integer.
  • enumeration.data_type.u32: 32-bit unsigned integer.
  • enumeration.data_type.u64: 64-bit unsigned integer.

enumeration.argument_type

Types of kernel and compute graph argument.

  • enumeration.argument_type.i32: 32-bit one's complement signed integer.
  • enumeration.argument_type.f32: 32-bit IEEE 754 single-precision floating-point number.
  • enumeration.argument_type.ndarray: ND-array wrapped around a handle.memory.
  • enumeration.argument_type.texture: Texture wrapped around a handle.image.
  • enumeration.argument_type.scalar: Typed scalar.
  • enumeration.argument_type.tensor: Typed tensor.

bit_field.memory_usage

Usages of a memory allocation. Taichi requires kernel argument memories to be allocated with bit_field.memory_usage.storage.

  • bit_field.memory_usage.storage: The memory can be read/write accessed by any kernel.
  • bit_field.memory_usage.uniform: The memory can be used as a uniform buffer in graphics pipelines.
  • bit_field.memory_usage.vertex: The memory can be used as a vertex buffer in graphics pipelines.
  • bit_field.memory_usage.index: The memory can be used as an index buffer in graphics pipelines.

structure.memory_allocate_info

Parameters of a newly allocated memory.

  • structure.memory_allocate_info.size: Size of the allocation in bytes.
  • structure.memory_allocate_info.host_write: True if the host needs to write to the allocated memory.
  • structure.memory_allocate_info.host_read: True if the host needs to read from the allocated memory.
  • structure.memory_allocate_info.export_sharing: True if the memory allocation needs to be exported to other backends (e.g., from Vulkan to CUDA).
  • structure.memory_allocate_info.usage: All possible usage of this memory allocation. In most cases, bit_field.memory_usage.storage is enough.

structure.memory_slice

A subsection of a memory allocation. The sum of structure.memory_slice.offset and structure.memory_slice.size cannot exceed the size of structure.memory_slice.memory.

  • structure.memory_slice.memory: The subsectioned memory allocation.
  • structure.memory_slice.offset: Offset from the beginning of the allocation.
  • structure.memory_slice.size: Size of the subsection.

structure.nd_shape

Multi-dimensional size of an ND-array. Dimension sizes after structure.nd_shape.dim_count are ignored.

  • structure.nd_shape.dim_count: Number of dimensions.
  • structure.nd_shape.dims: Dimension sizes.

structure.nd_array

Multi-dimensional array of dense primitive data.

  • structure.nd_array.memory: Memory bound to the ND-array.
  • structure.nd_array.shape: Shape of the ND-array.
  • structure.nd_array.elem_shape: Shape of the ND-array elements. It must not be empty for vector or matrix ND-arrays.
  • structure.nd_array.elem_type: Primitive data type of the ND-array elements.

bit_field.image_usage

Usages of an image allocation. Taichi requires kernel argument images to be allocated with bit_field.image_usage.storage and bit_field.image_usage.sampled.

  • bit_field.image_usage.storage: The image can be read/write accessed by any kernel.
  • bit_field.image_usage.sampled: The image can be read-only accessed by any kernel.
  • bit_field.image_usage.attachment: The image can be used as a color or depth-stencil attachment depending on its format.

enumeration.image_dimension

Dimensions of an image allocation.

  • enumeration.image_dimension.1d: The image is 1-dimensional.
  • enumeration.image_dimension.2d: The image is 2-dimensional.
  • enumeration.image_dimension.3d: The image is 3-dimensional.
  • enumeration.image_dimension.1d_array: The image is 1-dimensional and it has one or more layers.
  • enumeration.image_dimension.2d_array: The image is 2-dimensional and it has one or more layers.
  • enumeration.image_dimension.cube: The image is 2-dimensional and it has 6 layers for the faces towards +X, -X, +Y, -Y, +Z, -Z in sequence.

enumeration.image_layout

  • enumeration.image_layout.undefined: Undefined layout. An image in this layout does not contain any semantical information.
  • enumeration.image_layout.shader_read: Optimal layout for read-only access, including sampling.
  • enumeration.image_layout.shader_write: Optimal layout for write-only access.
  • enumeration.image_layout.shader_read_write: Optimal layout for read/write access.
  • enumeration.image_layout.color_attachment: Optimal layout as a color attachment.
  • enumeration.image_layout.color_attachment_read: Optimal layout as an input color attachment.
  • enumeration.image_layout.depth_attachment: Optimal layout as a depth attachment.
  • enumeration.image_layout.depth_attachment_read: Optimal layout as an input depth attachment.
  • enumeration.image_layout.transfer_dst: Optimal layout as a data copy destination.
  • enumeration.image_layout.transfer_src: Optimal layout as a data copy source.
  • enumeration.image_layout.present_src: Optimal layout as a presentation source.

enumeration.format

Texture formats. The availability of texture formats depends on runtime support.

structure.image_offset

Offsets of an image in X, Y, Z, and array layers.

  • structure.image_offset.x: Image offset in the X direction.
  • structure.image_offset.y: Image offset in the Y direction. Must be 0 if the image has a dimension of enumeration.image_dimension.1d or enumeration.image_dimension.1d_array.
  • structure.image_offset.z: Image offset in the Z direction. Must be 0 if the image has a dimension of enumeration.image_dimension.1d, enumeration.image_dimension.2d, enumeration.image_dimension.1d_array, enumeration.image_dimension.2d_array or enumeration.image_dimension.cube_array.
  • structure.image_offset.array_layer_offset: Image offset in array layers. Must be 0 if the image has a dimension of enumeration.image_dimension.1d, enumeration.image_dimension.2d or enumeration.image_dimension.3d.

structure.image_extent

Extents of an image in X, Y, Z, and array layers.

  • structure.image_extent.width: Image extent in the X direction.
  • structure.image_extent.height: Image extent in the Y direction. Must be 1 if the image has a dimension of enumeration.image_dimension.1d or enumeration.image_dimension.1d_array.
  • structure.image_extent.depth: Image extent in the Z direction. Must be 1 if the image has a dimension of enumeration.image_dimension.1d, enumeration.image_dimension.2d, enumeration.image_dimension.1d_array, enumeration.image_dimension.2d_array or enumeration.image_dimension.cube_array.
  • structure.image_extent.array_layer_count: Image extent in array layers. Must be 1 if the image has a dimension of enumeration.image_dimension.1d, enumeration.image_dimension.2d or enumeration.image_dimension.3d. Must be 6 if the image has a dimension of enumeration.image_dimension.cube_array.

structure.image_allocate_info

Parameters of a newly allocated image.

  • structure.image_allocate_info.dimension: Image dimension.
  • structure.image_allocate_info.extent: Image extent.
  • structure.image_allocate_info.mip_level_count: Number of mip-levels.
  • structure.image_allocate_info.format: Image texel format.
  • structure.image_allocate_info.export_sharing: True if the memory allocation needs to be exported to other backends (e.g., from Vulkan to CUDA).
  • structure.image_allocate_info.usage: All possible usages of this image allocation. In most cases, bit_field.image_usage.storage and bit_field.image_usage.sampled enough.

structure.image_slice

A subsection of a memory allocation. The sum of structure.image_slice.offset and structure.image_slice.extent in each dimension cannot exceed the size of structure.image_slice.image.

  • structure.image_slice.image: The subsectioned image allocation.
  • structure.image_slice.offset: Offset from the beginning of the allocation in each dimension.
  • structure.image_slice.extent: Size of the subsection in each dimension.
  • structure.image_slice.mip_level: The subsectioned mip-level.

structure.texture

Image data bound to a sampler.

  • structure.texture.image: Image bound to the texture.
  • structure.texture.sampler: The bound sampler that controls the sampling behavior of structure.texture.image.
  • structure.texture.dimension: Image Dimension.
  • structure.texture.extent: Image extent.
  • structure.texture.format: Image texel format.

union.scalar_value

Scalar value represented by a power-of-two number of bits.

NOTE The unsigned integer types merely hold the number of bits in memory and doesn't reflect any type of the underlying data. For example, a 32-bit floating-point scalar value is assigned by *(float*)&scalar_value.x32 = 0.0f; a 16-bit signed integer is assigned by *(int16_t)&scalar_vaue.x16 = 1. The actual type of the scalar is hinted via structure.scalar.type.

  • union.scalar_value.x8: Scalar value that fits into 8 bits.
  • union.scalar_value.x16: Scalar value that fits into 16 bits.
  • union.scalar_value.x32: Scalar value that fits into 32 bits.
  • union.scalar_value.x64: Scalar value that fits into 64 bits.

structure.scalar

A typed scalar value.

union.tensor_value

Tensor value represented by a power-of-two number of bits.

  • union.tensor_value.x8: Tensor value that fits into 8 bits.
  • union.tensor_value.x16: Tensor value that fits into 16 bits.
  • union.tensor_value.x32: Tensor value that fits into 32 bits.
  • union.tensor_value.x64: Tensor value that fits into 64 bits.

structure.tensor_value_with_length

A tensor value with a length.

structure.tensor

A typed tensor value.

union.argument_value

A scalar or structured argument value.

  • union.argument_value.i32: Value of a 32-bit one's complement signed integer. This is equivalent to union.scalar_value.x32 with enumeration.data_type.i32.
  • union.argument_value.f32: Value of a 32-bit IEEE 754 single-precision floating-poing number. This is equivalent to union.scalar_value.x32 with enumeration.data_type.f32.
  • union.argument_value.ndarray: An ND-array to be bound.
  • union.argument_value.texture: A texture to be bound.
  • union.argument_value.scalar: An scalar to be bound.
  • union.argument_value.tensor: A tensor to be bound.

structure.argument

An argument value to feed kernels.

  • structure.argument.type: Type of the argument.
  • structure.argument.value: Value of the argument.

structure.named_argument

A named argument value to feed compute graphs.

  • structure.named_argument.name: Name of the argument.
  • structure.named_argument.argument: Argument body.

function.get_version

Get the current taichi version. It has the same value as TI_C_API_VERSION as defined in taichi_core.h.

function.get_available_archs

Gets a list of available archs on the current platform. An arch is only available if:

  1. The Runtime library is compiled with its support;
  2. The current platform is installed with a capable hardware or an emulation software.

An available arch has at least one device available, i.e., device index 0 is always available. If an arch is not available on the current platform, a call to function.create_runtime with that arch is guaranteed failing.

WARNING Please also note that the order or returned archs is undefined.

function.get_last_error

Gets the last error raised by Taichi C-API invocations. Returns the semantical error code.

  • function.get_last_error.message_size: Size of textual error message in function.get_last_error.message
  • function.get_last_error.message: Text buffer for the textual error message. Ignored when message_size is 0.

function.set_last_error

Sets the provided error as the last error raised by Taichi C-API invocations. It can be useful in extended validation procedures in Taichi C-API wrappers and helper libraries.

  • function.set_last_error.error: Semantical error code.
  • function.set_last_error.message: A null-terminated string of the textual error message or nullptr for empty error message.

function.create_runtime

Creates a Taichi Runtime with the specified enumeration.arch.

  • function.create_runtime.arch: Arch of Taichi Runtime.
  • function.create_runtime.device_index: The index of device in function.create_runtime.arch to create Taichi Runtime on.

function.destroy_runtime

Destroys a Taichi Runtime.

function.set_runtime_capabilities

Force override the list of available capabilities in the runtime instance.

function.get_runtime_capabilities

Gets all capabilities available on the runtime instance.

  • function.get_runtime_capabilities.capability_count: The total number of capabilities available.
  • function.get_runtime_capabilities.capabilities: Returned capabilities.

function.allocate_memory

Allocates a contiguous device memory with provided parameters.

function.free_memory

Frees a memory allocation.

function.map_memory

Maps a device memory to a host-addressable space. You must ensure that the device is not being used by any device command before the mapping.

function.unmap_memory

Unmaps a device memory and makes any host-side changes about the memory visible to the device. You must ensure that there is no further access to the previously mapped host-addressable space.

function.allocate_image

Allocates a device image with provided parameters.

function.free_image

Frees an image allocation.

function.copy_memory_device_to_device

Copies the data in a contiguous subsection of the device memory to another subsection. The two subsections must not overlap.

function.copy_image_device_to_device

Copies the image data in a contiguous subsection of the device image to another subsection. The two subsections must not overlap.

function.track_image

Tracks the device image with the provided image layout. Because Taichi tracks image layouts internally, it is only useful to inform Taichi that the image is transitioned to a new layout by external procedures.

function.transition_image

Transitions the image to the provided image layout. Because Taichi tracks image layouts internally, it is only useful to enforce an image layout for external procedures to use.

function.launch_kernel

Launches a Taichi kernel with the provided arguments. The arguments must have the same count and types in the same order as in the source code.

function.launch_compute_graph

Launches a Taichi compute graph with provided named arguments. The named arguments must have the same count, names, and types as in the source code.

function.flush

Submits all previously invoked device commands to the offload device for execution.

function.wait

Waits until all previously invoked device commands are executed. Any invoked command that has not been submitted is submitted first.

function.load_aot_module

Loads a pre-compiled AOT module from the file system.
Returns definition.null_handle if the runtime fails to load the AOT module from the specified path.

function.create_aot_module

Creates a pre-compiled AOT module from TCM data.
Returns definition.null_handle if the runtime fails to create the AOT module from TCM data.

function.destroy_aot_module

Destroys a loaded AOT module and releases all related resources.

function.get_aot_module_kernel

Retrieves a pre-compiled Taichi kernel from the AOT module.
Returns definition.null_handle if the module does not have a kernel of the specified name.

function.get_aot_module_compute_graph

Retrieves a pre-compiled compute graph from the AOT module.
Returns definition.null_handle if the module does not have a compute graph of the specified name.

File: c_api/docs/taichi/taichi_vulkan.h.md

Vulkan Backend Features

Taichi's Vulkan API gives you further control over the Vulkan version and extension requirements and allows you to interop with external Vulkan applications with shared resources.

API Reference

structure.vulkan_runtime_interop_info

Necessary detail to share the same Vulkan runtime between Taichi and external procedures.

  • structure.vulkan_runtime_interop_info.get_instance_proc_addr: Pointer to Vulkan loader function vkGetInstanceProcAddr.
  • structure.vulkan_runtime_interop_info.api_version: Target Vulkan API version.
  • structure.vulkan_runtime_interop_info.instance: Vulkan instance handle.
  • structure.vulkan_runtime_interop_info.physical_device: Vulkan physical device handle.
  • structure.vulkan_runtime_interop_info.device: Vulkan logical device handle.
  • structure.vulkan_runtime_interop_info.compute_queue: Vulkan queue handle created in the queue family at structure.vulkan_runtime_interop_info.compute_queue_family_index.
  • structure.vulkan_runtime_interop_info.compute_queue_family_index: Index of a Vulkan queue family with the VK_QUEUE_COMPUTE_BIT set.
  • structure.vulkan_runtime_interop_info.graphics_queue: Vulkan queue handle created in the queue family at structure.vulkan_runtime_interop_info.graphics_queue_family_index.
  • structure.vulkan_runtime_interop_info.graphics_queue_family_index: Index of a Vulkan queue family with the VK_QUEUE_GRAPHICS_BIT set.

NOTE structure.vulkan_runtime_interop_info.compute_queue and structure.vulkan_runtime_interop_info.graphics_queue can be the same if the queue family have VK_QUEUE_COMPUTE_BIT and VK_QUEUE_GRAPHICS_BIT set at the same tiem.

structure.vulkan_memory_interop_info

Necessary detail to share the same piece of Vulkan buffer between Taichi and external procedures.

  • structure.vulkan_memory_interop_info.buffer: Vulkan buffer.
  • structure.vulkan_memory_interop_info.size: Size of the piece of memory in bytes.
  • structure.vulkan_memory_interop_info.usage: Vulkan buffer usage. In most of the cases, Taichi requires the VK_BUFFER_USAGE_STORAGE_BUFFER_BIT.
  • structure.vulkan_memory_interop_info.memory: Device memory binded to the Vulkan buffer.
  • structure.vulkan_memory_interop_info.offset: Offset in VkDeviceMemory object to the beginning of this allocation, in bytes.

structure.vulkan_image_interop_info

Necessary detail to share the same piece of Vulkan image between Taichi and external procedures.

  • structure.vulkan_image_interop_info.image: Vulkan image.
  • structure.vulkan_image_interop_info.image_type: Vulkan image allocation type.
  • structure.vulkan_image_interop_info.format: Pixel format.
  • structure.vulkan_image_interop_info.extent: Image extent.
  • structure.vulkan_image_interop_info.mip_level_count: Number of mip-levels of the image.
  • structure.vulkan_image_interop_info.array_layer_count: Number of array layers.
  • structure.vulkan_image_interop_info.sample_count: Number of samples per pixel.
  • structure.vulkan_image_interop_info.tiling: Image tiling.
  • structure.vulkan_image_interop_info.usage: Vulkan image usage. In most cases, Taichi requires the VK_IMAGE_USAGE_STORAGE_BIT and the VK_IMAGE_USAGE_SAMPLED_BIT.

structure.vulkan_event_interop_info

Necessary detail to share the same Vulkan event synchronization primitive between Taichi and the user application.

  • structure.vulkan_event_interop_info.event: Vulkan event handle.

function.create_vulkan_runtime

Creates a Vulkan Taichi runtime with user-controlled capability settings.

function.import_vulkan_runtime

Imports the Vulkan runtime owned by Taichi to external procedures.

function.export_vulkan_runtime

Exports a Vulkan runtime from external procedures to Taichi.

function.import_vulkan_memory

Imports the Vulkan buffer owned by Taichi to external procedures.

function.export_vulkan_memory

Exports a Vulkan buffer from external procedures to Taichi.

function.import_vulkan_image

Imports the Vulkan image owned by Taichi to external procedures.

function.export_vulkan_image

Exports a Vulkan image from external procedures to Taichi.

function.import_vulkan_event

Imports the Vulkan event owned by Taichi to external procedures.

function.export_vulkan_event

Exports a Vulkan event from external procedures to Taichi.

File: c_api/version_scripts/README.md

Taichi C-API uses version script to control symbol exports from the shared library (https://github.com/taichi-dev/taichi/issues/6722)

Note that the two version scripts export_symbols_linux.ld and export_symbols_mac.ld should have consistent contents otherwise will cause symbols mismatch on different platforms.

CI will check for symbol leakage automatically, the regular expression rule of which should also stay consistent with the version scripts. CI scripts can be found at "check-c-api-export-symbols" function in .github/workflows/scripts/aot-demo.sh

File: conda/README.md

Installation

:::note
The scripts provided here only work for Linux/macOS.

We need an equivalent bat script for Windows.
:::

Download and install the conda matching your OS: https://docs.conda.io/en/latest/miniconda.html.

Run the following commands:

bash
# Assuming you're in the root directory
$ cd conda

$ conda env create -f conda_env.yaml
$ conda activate taichi-dev

# This will only install the dependencies to the 'taichi-dev' environment.
(taichi-dev) $ python3 -m pip install -r ../requirements_dev.txt

Update `taichi-dev`'s environment variables

bash
# Change scripts/activate_env_vars.sh

$ conda activate taichi-dev

(taichi-dev) $ ./update_env_vars.sh

File: docs/design/llvm_sparse_runtime.md

How is Taichi's LLVM Sparse Runtime Implemented?

Last update: 2022-04-21

The Taichi LLVM sparse runtime can be found in the taichi/runtime/llvm directory of the Taichi repository on GitHub. It consists of one header file for each SNode type and a single common source file, runtime.cpp, which applies to all SNode types. The URL for the Taichi repository on GitHub is: https://github.com/taichi-dev/taichi/tree/master/taichi/runtime/llvm/runtime_module

SNode

There are four types of SNodes in Taichi: dense, bitmasked, pointer and dynamic. dynamic or pointer SNodes are spatially sparse in nature: Cells in them are not necessarily stored contiguously. Therefore, when this doc talks about sparse SNodes hereafter, it refers to dynamic and pointer SNodes.

Each SNode type comes with an XMeta struct derived from StructMeta. StructMeta has the following properties:

  • snode_id: SNode ID.
  • i32 X_get_num_elements(Ptr meta, Ptr node): Returns the capacity this SNode can hold. Note that it is not the current number of active cells, but the maximum.
  • void X_activate(Ptr meta, Ptr node, int i): Activates cell i.
  • u1 X_is_active(Ptr meta, Ptr node, int i): Checks whether cell i is active.
  • Ptr X_lookup_element(Ptr meta, Ptr node, int i): Returns the pointer to cell i. The result can be nullptr for sparse SNodes.

Here, Ptr is an alias for uint8_t*, and i32 for int32_t. As the name suggests, meta points to the corresponding meta struct, while node points to the SNode instance.

For sparse SNodes, they also implement an additional API:

  • void X_deactivate(Ptr meta, Ptr node, int i): Deactivates cell i.

However, this additional API is subject to change to make all SNodes share the same set of APIs in the future.

`dense` SNode

dense is the simplest form of SNode. It is just an array of cells living in a chunk of contiguous memory, or std::array<Cell, N> for users with a C++ background. Its header file is in node_dense.h.

  • Dense_get_num_elements: This is just max_num_elements stored in DenseMeta.
  • Dense_activate: No-op, because in a dense SNode cells are always activated.
  • Dense_is_active: Always returns 1.
  • Dense_lookup_element: Returns the address of the i-th cell. That is, node + element_size * i.

Layout of a dense SNode:

sh
+- node
|
+------------+------------+------------+------------+
|            |            |            |            |
|   cell-0   |   cell-1   |   cell-2   |   cell-3   |
|            |            |            |            |
+------------+------------+------------+------------+

`pointer` SNode

The pointer SNode is a popular choice for sparse computation. It dynamically allocates memory only for activated cells and recycles it back into a memory pool once the cell is deactivated, conserving memory resources in large-scale grid computation. The pointer SNode can be thought of as an std::array<Cell*, N>.

Upon initialization, Taichi preallocates a chunk of memory space named ambient_elements, which is shared across all inactive sparse SNodes. Therefore, dereferencing an inactive sparse SNode generates the default value (usually zero) stored in ambient_elements.

Layout of a pointer SNode:

sh
+- node
|
+------------+------------+------------+------------+------------+------------+------------+------------+
|  c0-lock   |  c1-lock   |  c2-lock   |  c3-lock   |  nullptr   |  *cell-1   |  *cell-2   |  nullptr   |
+------------+------------+------------+------------+------------+------------+------------+------------+
                                                                 |            |
                                                                 |            +> +------------+
                                                                 |               |            |
                                                                 |               |   cell-2   |
                                                                 |               |            |
                                                                 |               +------------+
                                                                 |
                                                                 +-------------> +------------+
                                                                                 |            |
                                                                                 |   cell-1   |
                                                                                 |            |
                                                                                 +------------+

We can follow Pointer_activate to see how the pointer SNode is implemented using the sparse runtime infrastructure.

cpp
void Pointer_activate(Ptr meta_, Ptr node, int i) {
  auto meta = (StructMeta *)meta_;
  auto num_elements = Pointer_get_num_elements(meta_, node);
  // 1
  volatile Ptr lock = node + 8 * i;
  volatile Ptr *data_ptr = (Ptr *)(node + 8 * (num_elements + i));
  // 2
  if (*data_ptr == nullptr) {
    // 3
    // The cuda_ calls will return 0 or do noop on CPUs
    u32 mask = cuda_active_mask();
    if (is_representative(mask, (u64)lock)) {
      // 4
      locked_task(
          lock,
          [&] {
            // 5
            auto rt = meta->context->runtime;
            auto alloc = rt->node_allocators[meta->snode_id];
            auto allocated = (u64)alloc->allocate();
            atomic_exchange_u64((u64 *)data_ptr, allocated);
          },
          [&]() { return *data_ptr == nullptr; });
    }
    warp_barrier(mask);
  }
}
  1. Retrieves both the lock and the pointer for the i-th cell. Note that the pointer size is assumed to be 8-byte wide. Locks are simply 64-bit integers. According to the layout, there are max_num_elements number of locks, followed by max_num_elements number of pointers to cells.
  2. Checks whether the content of data_ptr is nullptr without any locking. This is the classical double-checked locking pattern.
  3. If 2 is true, pick one thread within a CUDA warp to acquire the lock. This is a small optimization to prevent lock contention. On pre-Volta devices without independent thread scheduling, this also prevents deadlocking.
  4. The winning thread tries to acquire the lock using locked_task.
  5. Retrieves the memory allocator (node_allocators) for this particular SNode, allocates a new cell, and stores the address of the allocated cell into data_ptr. Because the cell size of each SNode is different, the runtime has a dedicated allocator for each SNode, which knows how much space to allocate per cell.

The procedures for deactivating and checking the active status of cells are similar, so they have been omitted for brevity.

`dynamic` SNode

dynamic is a special kind of SNodes in a few ways:

  • It must be a 1-D, terminating SNode. By terminating, it means dynamic can only be followed by place-ing a Taichi field.
  • The axis of dynamic must be different from those of all its predecessors. For example:
    • dense(ti.ij, (2, 4)).dynamic(ti.k, 8): This is OK, because dynamic's axis is unique.
    • dense(ti.ij, (2, 4)).dynamic(ti.j, 8): This results in an error, because dynamic's axis and dense's overlaps on axis j.

Below shows the layout of a dynamic SNode. Logically speaking, dynamic SNode can be viewed as std::vector. However, dynamic is implemented as a singly linked list of chunks.

sh
# `n` stores the number of elements in this dynammic SNode cell

+- node
|
+------------+------------+------------+
|    lock    |     n      |    ptr     |
+------------+------------+------------+
                          |
                          +-+------------+  # chunk-0, chunk_start = 0
                            |     >------|--+
                            +------------+  |
                            |     0      |  |
                            +------------+  |
                            |     1      |  |
                            +------------+  |
                            |     2      |  |
                            +------------+  |
                            |     3      |  |
                            +------------+  |
                                            |
                                            +-+------------+  # chunk-1, chunk_start = 4
                                              |   nullptr  |
                                              +------------+
                                              |     4      |
                                              +------------+
                                              |     5      |
                                              +------------+
                                              |     6      |
                                              +------------+
                                              |     7      |
                                              +------------+

The activation/deactivation process is no different from that of pointer. We can trace through Dynamic_append to better understand the layout.

cpp
i32 Dynamic_append(Ptr meta_, Ptr node_, i32 data) {
  auto meta = (DynamicMeta *)(meta_);
  auto node = (DynamicNode *)(node_);
  auto chunk_size = meta->chunk_size;
  // 1
  auto i = atomic_add_i32(&node->n, 1);
  // 2
  int chunk_start = 0;
  auto p_chunk_ptr = &node->ptr;
  while (true) {
    // 3
    if (*p_chunk_ptr == nullptr) {
      locked_task(Ptr(&node->lock), [&] {
        if (*p_chunk_ptr == nullptr) {
          auto rt = meta->context->runtime;
          auto alloc = rt->node_allocators[meta->snode_id];
          *p_chunk_ptr = alloc->allocate();
        }
      });
    }
    // 4
    if (i < chunk_start + chunk_size) {
      // 4-1
      *(i32 *)(*p_chunk_ptr + sizeof(Ptr) +
               (i - chunk_start) * meta->element_size) = data;
      break;
    }
    // 4-2
    p_chunk_ptr = (Ptr *)(*p_chunk_ptr);
    chunk_start += chunk_size;
  }
  return i;
}
  1. Uses the current length n as the index (i) to store data.
  2. chunk_start tracks the starting index of a given chunk, and always starts at 0. p_chunk_ptr is initialized to the pointer to the first chunk.
  3. Inside the while loop, checks if the given chunk slot is empty first, and allocates a new chunk if so.
  4. Compares if the determined index i falls within the current chunk.
    1. If so, stores data into the corresponding slot in this chunk. Note that the first sizeof(Ptr) bytes are skipped: they are reserved to store the address of the next chunk.
    2. Otherwise, jumps to the next chunk.

Runtime

The runtime code for the LLVM backends can be found at this URL: https://github.com/taichi-dev/taichi/blob/master/taichi/runtime/llvm/runtime_module/runtime.cpp. It is important to note that this file is not linked into the core C++ library of Taichi, but instead is compiled into an LLVM byte code file (.bc). When Taichi starts, the .bc file is loaded into memory, deserialized into an llvm::Module, and linked with the JIT compiled Taichi kernels. This design allows the runtime code to be written once and shared between all LLVM backends, such as CPU and CUDA. Furthermore, the sparse runtime can be implemented in any language with sufficient abstraction (e.g., C++) rather than in raw LLVM IR.

The core data structure of this runtime is LLVMRuntime, which holds a handful of data:

  • All the root SNodes info (roots and root_mem_sizes).
  • SNode memory allocators.
  • Random states (rand_states) used by ti.random.
  • Print and error message buffers.
  • ...

The SNode memory allocator is the bedrock of sparse SNodes. The following sections explain how it is implemented.

`NodeManager`: a recycling memory allocator

Each SNode is associated with its own memory allocator. These allocators are stored in an array, node_allocators.

The allocator is of type NodeManager. It contains three linked lists of type ListManager:

  • data_list: A list of fixed-sized memory chunks. Each chunk can store chunk_num_elements SNode cells. The aforementioned SNode chunks are from this list.
  • free_list: Indices of the free SNode cells. Each node in this list is an int32_t (or list_data_type). When allocating, the runtime will first try to reuse a cell in the free list if there is one available, before requesting extra space from the memory allocator. (More details below.)
  • recycled_list: Indices of the released SNode cells. After a GC execution, items in this list will be transferred into free_list for reuse. Each node in this list is also an int32_t.

Here's how allocate() is implemented:

cpp
Ptr allocate() {
  // 1
  int old_cursor = atomic_add_i32(&free_list_used, 1);
  i32 l;
  if (old_cursor >= free_list->size()) {
    // 2
    l = data_list->reserve_new_element();
  } else {
    // 3
    l = free_list->get<list_data_type>(old_cursor);
  }
  // 4
  return data_list->get_element_ptr(l);
}
  1. Reads the (possible) index of the next free cell.
  2. If running out of the indices in the free list, allocate a new chunk from data_list.
  3. Otherwise, re-use the index from free_list.
  4. Either way, index l points to a memory slot in data_list. Returns that slot.

So far, we have analyzed the allocate procedure. For the recycle procedure, the code recycle() is self-explanatory.

Before jumping into the garbage collection (GC) system, we will also take look at the lower-level ListManager.

`ListManager`: a CPU/GPU linked list

The way ListManager gets implemented is described as follows:

A simple list data structure that is infinitely long. Data are organized in chunks, where each chunk is allocated on demand.

Calling it a linked list can be a bit misleading. In fact, it is more like std::deque in C++. ListManager holds an array of memory chunks in chunks, whereas each chunk can hold max_num_elements_per_chunk elements.

ListManager's constructor shows the necessary member variables:

cpp
ListManager(LLVMRuntime *runtime,
            std::size_t element_size,
            std::size_t num_elements_per_chunk)
    : element_size(element_size),
      max_num_elements_per_chunk(num_elements_per_chunk),
      runtime(runtime) {
  taichi_assert_runtime(runtime, is_power_of_two(max_num_elements_per_chunk),
                        "max_num_elements_per_chunk must be POT.");
  lock = 0;
  num_elements = 0;
  log2chunk_num_elements = taichi::log2int(num_elements_per_chunk);
}

The important method for allocating a new element is reserve_new_element():

cpp
i32 reserve_new_element() {
  auto i = atomic_add_i32(&num_elements, 1);
  auto chunk_id = i >> log2chunk_num_elements;
  touch_chunk(chunk_id);
  return i;
}

It increments num_elements to get the index of this new element, and calculates the belonging chunk ID. Then it ensures that the chunk is actually allocated using touch_chunk():

cpp
void ListManager::touch_chunk(int chunk_id) {
  taichi_assert_runtime(runtime, chunk_id < max_num_chunks,
                        "List manager out of chunks.");
  if (!chunks[chunk_id]) {
    locked_task(&lock, [&] {
      // may have been allocated during lock contention
      if (!chunks[chunk_id]) {
        grid_memfence();
        auto chunk_ptr = runtime->request_allocate_aligned(
            max_num_elements_per_chunk * element_size, 4096);
        atomic_exchange_u64((u64 *)&chunks[chunk_id], (u64)chunk_ptr);
      }
    });
  }
}

By now, you should be quite familiar with this double-checked locking pattern. One thing noteworthy is that the fundamental memory allocation function is provided by LLVMRuntime::request_allocate_aligned().

get_element_ptr() illustrates the way to lookup the address of an element from a given index:

cpp
Ptr get_element_ptr(i32 i) {
  return chunks[i >> log2chunk_num_elements] +                       // chunk base
          element_size * (i & ((1 << log2chunk_num_elements) - 1));  // slot within the chunk
}

There is also a reverse operation, ptr2index(). It iterates over each chunk to see if the passed-in ptr falls into that memory range.

With these primitive helpers, allocate() is trivial to implement:

cpp
Ptr ListManager::allocate() {
  auto i = reserve_new_element();
  return get_element_ptr(i);
}

So is append():

cpp
void ListManager::append(void *data_ptr) {
  auto ptr = allocate();
  std::memcpy(ptr, data_ptr, element_size);
}

clear() simply resets num_elements to zero, without doing anything to the list contents.

Garbage collection (GC)

GC will happen (in a parallel way on GPUs), after an offloaded task with possible sparse SNode deactivations.
The GC process for a given SNode is divided into three stages:

  1. gc_parallel_0: Moves the remaining, unused indices in free_list to its head.

Normally this can be done via a simple for loop that copies data. However, we need to run it in parallel on GPUs and the data source and destination may overlap, so special care is needed.
Therefore, the code differentiates between the cases where destination and source ranges overlap or not:

sh
# Legend: "\\\" for already re-used cell indices, "   " (empty) for cell indices that are still available.
# src and dst range do not overlap:  src=[6, 8), dst=[0, 2)
0   1   2   3   4   5   6   7   8        0   1   2
+---+---+---+---+---+---+---+---+        +---+---+
|\\\|\\\|\\\|\\\|\\\|\\\|   |   |  --->  |   |   |
+---+---+---+---+---+---+---+---+        +---+---+
                        src              dst

# src and dst range overlap: src=[3, 8), dst=[0, 5)
0   1   2   3   4   5   6   7   8        0   1   2   3   4   5
+---+---+---+---+---+---+---+---+        +---+---+---+---+---+
|\\\|\\\|\\\|   |   |   |   |   |  --->  |   |   |   |   |   |
+---+---+---+---+---+---+---+---+        +---+---+---+---+---+
            src                          dst

Note cell indices 3 and 4 are in both the source and the destination regions.

  1. gc_parallel_1: Does some bookkeeping for free_list and recycled_list.
cpp
void gc_parallel_1(RuntimeContext *context, int snode_id) {
  LLVMRuntime *runtime = context->runtime;
  auto allocator = runtime->node_allocators[snode_id];
  auto free_list = allocator->free_list;

  const i32 num_unused =
      max_i32(free_list->size() - allocator->free_list_used, 0);
  free_list->resize(num_unused);

  allocator->free_list_used = 0;
  // 1
  allocator->recycle_list_size_backup = allocator->recycled_list->size();
  // 2
  allocator->recycled_list->clear();
}

It is important to understand that this stage must run on a single thread in order to avoid data races on recycled_list. During GC, recycled_list must be cleared, with all of its contents transferred into free_list. If this were done in the third stage (which runs in parallel), it would be difficult to coordinate the event sequence of clearing and transferring among the GPU threads. As a result, this serial stage is created, when the following steps are implemented to handle the storage and clearing of the elements in the recycled_list:

  1. Back up the number of elements in the recycled_list into recycle_list_size_backup.
  2. Clear the recycled_list.

Note: The computation task involved in these steps is relatively simple and does not pose a significant burden in terms of processing time, even when executed in serial.

  1. gc_parallel_2: Replenishes free_list with the indices in recycled_list, then zero-fills all the recycled memory locations.
cpp
void gc_parallel_2(RuntimeContext *context, int snode_id) {
  LLVMRuntime *runtime = context->runtime;
  auto allocator = runtime->node_allocators[snode_id];
  // 1
  auto elements = allocator->recycle_list_size_backup;
  auto free_list = allocator->free_list;
  auto recycled_list = allocator->recycled_list;
  auto data_list = allocator->data_list;
  auto element_size = allocator->element_size;
  using T = NodeManager::list_data_type;
  // 2
  auto i = block_idx();
  while (i < elements) {
    // 3
    auto idx = recycled_list->get<T>(i);
    auto ptr = data_list->get_element_ptr(idx);
    // 4
    if (thread_idx() == 0) {
      free_list->push_back(idx);
    }
    // memset
    // 5
    auto ptr_stop = ptr + element_size;
    if ((uint64)ptr % 4 != 0) {
      auto new_ptr = ptr + 4 - (uint64)ptr % 4;
      if (thread_idx() == 0) {
        for (uint8 *p = ptr; p < new_ptr; p++) {
          *p = 0;
        }
      }
      ptr = new_ptr;
    }
    // now ptr is a multiple of 4
    // 6
    ptr += thread_idx() * sizeof(uint32);
    while (ptr + sizeof(uint32) <= ptr_stop) {
      *(uint32 *)ptr = 0;
      ptr += sizeof(uint32) * block_dim();
    }
    while (ptr < ptr_stop) {
      *ptr = 0;
      ptr++;
    }
    // 7
    i += grid_dim();
  }
}

  1. Reads back the size of recycled_list before clearing stored in stage 2.

  2. i is set to the CUDA block index. Note that all threads within a CUDA block cooperatively zero-fills the memory region of a SNode cell.

  3. Reads the SNode cell index stored in the i-th element in recycled_list. Recall that ListManager::clear() doesn't really modify its content, so it's still safe to read this list.

  4. The first thread of the block gets to append this cell index into free_list.

  5. This begins the zero-fill process. First thing first, thread 0 clears the first few bytes, in case the cell is not 4-byte aligned.

  6. Each thread clears its corresponding 4 bytes, then jumps to the next location, with stride being sizeof(uint32) * block_dim().

    sh
    # A SNode cell of 40 bytes, zero-filled by 4 threads:
+----+----+----+----+----+----+----+----+----+----+
| t0 | t1 | t2 | t3 | t0 | t1 | t2 | t3 | t0 | t1 |
+----+----+----+----+----+----+----+----+----+----+

  7. The entire thread block jumps to the next recycled SNode cell.

File: docs/lang/articles/visualization/export_results.md

Export Your Results

Taichi has functions that help you export visual results to images or
videos. This tutorial demonstrates how to use them step by step.

Export images

  • There are two ways to export visual results of your program to
    images.
  • The first and easier way is to make use of ti.GUI.
  • The second way is to call some Taichi functions such as
    ti.tools.imwrite.

Export images using `ti.GUI.show`

  • ti.GUI.show(filename) can not only display the GUI canvas on your
    screen, but also save the image to your specified filename.
  • Note that the format of the image is fully determined by the suffix
    of filename.
  • Taichi now supports saving to png, jpg, and bmp formats.
  • We recommend using png format. For example:
python {23}
import taichi as ti
import os

ti.init()

pixels = ti.field(ti.u8, shape=(512, 512, 3))

@ti.kernel
def paint():
    for i, j, k in pixels:
        pixels[i, j, k] = ti.random() * 255

iterations = 1000
gui = ti.GUI("Random pixels", res=512)

# mainloop
for i in range(iterations):
    paint()
    gui.set_image(pixels)

    filename = f'frame_{i:05d}.png'   # create filename with suffix png
    print(f'Frame {i} is recorded in {filename}')
    gui.show(filename)  # export and show in GUI

After running the code above, you will get a series of images in the current folder.

Export images using `ti.tools.imwrite`

To save images without invoking ti.GUI.show(filename), use
ti.tools.imwrite(filename). For example:

python {14}
import taichi as ti

ti.init()

pixels = ti.field(ti.u8, shape=(512, 512, 3))

@ti.kernel
def set_pixels():
    for i, j, k in pixels:
        pixels[i, j, k] = ti.random() * 255

set_pixels()
filename = f'imwrite_export.png'
ti.tools.imwrite(pixels.to_numpy(), filename)
print(f'The image has been saved to {filename}')
  • ti.tools.imwrite can export Taichi fields (ti.Matrix.field,
    ti.Vector.field, ti.field) and numpy arrays np.ndarray.
  • Same as above ti.GUI.show(filename), the image format (png,
    jpg and bmp) is also controlled by the suffix of filename in
    ti.tools.imwrite(filename).
  • Meanwhile, the resulted image type (grayscale, RGB, or RGBA) is
    determined by the number of channels in the input field, i.e.,
    the length of the third dimension (field.shape[2]).
  • In other words, a field that has shape (w, h) or (w, h, 1) will
    be exported as a grayscale image.
  • If you want to export RGB or RGBA images instead, the input
    field should have a shape (w, h, 3) or (w, h, 4) respectively.

:::note
All Taichi fields have their own data types, such as ti.u8 and
ti.f32. Different data types can lead to different behaviors of
ti.tools.imwrite. Please check out GUI system for
more details.
:::

  • Taichi offers other helper functions that read and show images in
    addition to ti.tools.imwrite. They are also demonstrated in
    GUI system.

Convert PNGs to video

Sometimes it's convenient to convert a series of png files into a
single video when showing your result to others.

For example, suppose you have 000000.png, 000001.png, ... generated
according to Export your results in the
current working directory.

Then you could run ti video to create a file video.mp4 containing
all these images as frames (sorted by file name).

Use ti video -f40 for creating a video with 40 FPS.

Export videos

:::note
The video export utilities of Taichi depend on ffmpeg. If ffmpeg is
not installed on your machine, please follow the installation
instructions of ffmpeg at the end of this page.
:::

  • ti.tools.VideoManager can help you export results in mp4 or gif
    format. For example,
python {13,24}
import taichi as ti

ti.init()

pixels = ti.field(ti.u8, shape=(512, 512, 3))

@ti.kernel
def paint():
    for i, j, k in pixels:
        pixels[i, j, k] = ti.random() * 255

result_dir = "./results"
video_manager = ti.tools.VideoManager(output_dir=result_dir, framerate=24, automatic_build=False)

for i in range(50):
    paint()

    pixels_img = pixels.to_numpy()
    video_manager.write_frame(pixels_img)
    print(f'\rFrame {i+1}/50 is recorded', end='')

print()
print('Exporting .mp4 and .gif videos...')
video_manager.make_video(gif=True, mp4=True)
print(f'MP4 video is saved to {video_manager.get_output_filename(".mp4")}')
print(f'GIF video is saved to {video_manager.get_output_filename(".gif")}')

After running the code above, you will find the output videos in the
./results/ folder.

Convert video to GIF

Sometimes you may need gif images to post a result on forums.

To do so, run ti gif -i video.mp4, where video.mp4 is the
mp4 video (generated with instructions above).

Use ti gif -i video.mp4 -f40 to create a GIF at 40 FPS.

Install ffmpeg

Install ffmpeg on Windows

  • Download the ffmpeg archive(named ffmpeg-2020xxx.zip) from
    ffmpeg.
  • Unzip this archive to a folder, such as D:/YOUR_FFMPEG_FOLDER.
  • Important: add D:/YOUR_FFMPEG_FOLDER/bin to the PATH
    environment variable.
  • Open the Windows cmd or PowerShell and type the line of code
    below to test your installation. If ffmpeg is set up properly, the
    version information will be printed.
bash
ffmpeg -version

Install `ffmpeg` on Linux

  • Most Linux distribution came with ffmpeg natively, so you do not
    need to read this part if the ffmpeg command is already there on
    your machine.
  • Install ffmpeg on Ubuntu
bash
sudo apt-get update
sudo apt-get install ffmpeg
  • Install ffmpeg on CentOS and RHEL
bash
sudo yum install ffmpeg ffmpeg-devel
  • Install ffmpeg on Arch Linux:
bash
pacman -S ffmpeg
  • Test your installation using
bash
ffmpeg -h

Install `ffmpeg` on macOS

  • ffmpeg can be installed on macOS using homebrew:
bash
brew install ffmpeg

Export PLY files

  • ti.tools.PLYWriter can help you export results in the ply format.
    Below is a short example of exporting 10 frames of a moving cube
    with vertices randomly colored,
python
import taichi as ti
import numpy as np

ti.init(arch=ti.cpu)

num_vertices = 1000
pos = ti.Vector.field(3, dtype=ti.f32, shape=(10, 10, 10))
rgba = ti.Vector.field(4, dtype=ti.f32, shape=(10, 10, 10))


@ti.kernel
def place_pos():
    for i, j, k in pos:
        pos[i, j, k] = 0.1 * ti.Vector([i, j, k])


@ti.kernel
def move_particles():
    for i, j, k in pos:
        pos[i, j, k] += ti.Vector([0.1, 0.1, 0.1])


@ti.kernel
def fill_rgba():
    for i, j, k in rgba:
        rgba[i, j, k] = ti.Vector(
            [ti.random(), ti.random(), ti.random(), ti.random()])


place_pos()
series_prefix = "example.ply"
for frame in range(10):
    move_particles()
    fill_rgba()
    # now adding each channel only supports passing individual np.array
    # so converting into np.ndarray, reshape
    # remember to use a temp var to store so you dont have to convert back
    np_pos = np.reshape(pos.to_numpy(), (num_vertices, 3))
    np_rgba = np.reshape(rgba.to_numpy(), (num_vertices, 4))
    # create a PLYWriter
    writer = ti.tools.PLYWriter(num_vertices=num_vertices)
    writer.add_vertex_pos(np_pos[:, 0], np_pos[:, 1], np_pos[:, 2])
    writer.add_vertex_rgba(
        np_rgba[:, 0], np_rgba[:, 1], np_rgba[:, 2], np_rgba[:, 3])
    writer.export_frame_ascii(frame, series_prefix)

After running the code above, you will find the output sequence of ply
files in the current working directory. Next, we will break down the
usage of ti.tools.PLYWriter into 4 steps and show some examples.

  • Setup ti.tools.PLYWriter
python
# num_vertices must be a positive int
# num_faces is optional, default to 0
# face_type can be either "tri" or "quad", default to "tri"

# in our previous example, a writer with 1000 vertices and 0 triangle faces is created
num_vertices = 1000
writer = ti.tools.PLYWriter(num_vertices=num_vertices)

# in the below example, a writer with 20 vertices and 5 quadrangle faces is created
writer2 = ti.tools.PLYWriter(num_vertices=20, num_faces=5, face_type="quad")
  • Add required channels
python
# A 2D grid with quad faces
#     y
#     |
# z---/
#    x
#         19---15---11---07---03
#         |    |    |    |    |
#         18---14---10---06---02
#         |    |    |    |    |
#         17---13---19---05---01
#         |    |    |    |    |
#         16---12---08---04---00

writer = ti.tools.PLYWriter(num_vertices=20, num_faces=12, face_type="quad")

# For the vertices, the only required channel is the position,
# which can be added by passing 3 np.array x, y, z into the following function.

x = np.zeros(20)
y = np.array(list(np.arange(0, 4))*5)
z = np.repeat(np.arange(5), 4)
writer.add_vertex_pos(x, y, z)

# For faces (if any), the only required channel is the list of vertex indices that each face contains.
indices = np.array([0, 1, 5, 4]*12)+np.repeat(
    np.array(list(np.arange(0, 3))*4)+4*np.repeat(np.arange(4), 3), 4)
writer.add_faces(indices)
  • Add optional channels
python
# Add custom vertex channel, the input should include a key, a supported datatype and, the data np.array
vdata = np.random.rand(20)
writer.add_vertex_channel("vdata1", "double", vdata)

# Add custom face channel
foo_data = np.zeros(12)
writer.add_face_channel("foo_key", "foo_data_type", foo_data)
# error! because "foo_data_type" is not a supported datatype. Supported ones are
# ['char', 'uchar', 'short', 'ushort', 'int', 'uint', 'float', 'double']

# PLYwriter already defines several useful helper functions for common channels
# Add vertex color, alpha, and rgba
# using float/double r g b alpha to reprent color, the range should be 0 to 1
r = np.random.rand(20)
g = np.random.rand(20)
b = np.random.rand(20)
alpha = np.random.rand(20)
writer.add_vertex_color(r, g, b)
writer.add_vertex_alpha(alpha)
# equivilantly
# add_vertex_rgba(r, g, b, alpha)

# vertex normal
writer.add_vertex_normal(np.ones(20), np.zeros(20), np.zeros(20))

# vertex index, and piece (group id)
writer.add_vertex_id()
writer.add_vertex_piece(np.ones(20))

# Add face index, and piece (group id)
# Indexing the existing faces in the writer and add this channel to face channels
writer.add_face_id()
# Set all the faces is in group 1
writer.add_face_piece(np.ones(12))
  • Export files
python
series_prefix = "example.ply"
series_prefix_ascii = "example_ascii.ply"
# Export a single file
# use ascii so you can read the content
writer.export_ascii(series_prefix_ascii)

# alternatively, use binary for a bit better performance
# writer.export(series_prefix)

# Export a sequence of files, ie in 10 frames
for frame in range(10):
    # write each frame as i.e. "example_000000.ply" in your current running folder
    writer.export_frame_ascii(frame, series_prefix_ascii)
    # alternatively, use binary
    # writer.export_frame(frame, series_prefix)

    # update location/color
    x = x + 0.1*np.random.rand(20)
    y = y + 0.1*np.random.rand(20)
    z = z + 0.1*np.random.rand(20)
    r = np.random.rand(20)
    g = np.random.rand(20)
    b = np.random.rand(20)
    alpha = np.random.rand(20)
    # re-fill
    writer = ti.tools.PLYWriter(num_vertices=20, num_faces=12, face_type="quad")
    writer.add_vertex_pos(x, y, z)
    writer.add_faces(indices)
    writer.add_vertex_channel("vdata1", "double", vdata)
    writer.add_vertex_color(r, g, b)
    writer.add_vertex_alpha(alpha)
    writer.add_vertex_normal(np.ones(20), np.zeros(20), np.zeros(20))
    writer.add_vertex_id()
    writer.add_vertex_piece(np.ones(20))
    writer.add_face_id()
    writer.add_face_piece(np.ones(12))

Import `ply` files into Houdini and Blender

Houdini supports importing a series of ply files sharing the same
prefix/post-fix. Our export_frame can achieve the requirement for you.

In Houdini, click File->Import->Geometry and navigate to the folder
containing your frame results, which should be collapsed into one single
entry like example_$F6.ply (0-9). Double-click this entry to finish
the importing process.

Blender requires an add-on called
Stop-motion-OBJ to
load the result sequences. Detailed
documentation is
provided by the author on how to install and use the add-on. If you're
using the latest version of Blender (2.80+), download and install the
latest
release
of Stop-motion-OBJ. For Blender 2.79 and older, use version v1.1.1 of
the add-on.

File: docs/lang/articles/visualization/ggui.md

A New UI system: GGUI

Category Prerequisites
OS Windows / Linux / Mac OS X
Backend x64 / CUDA / Vulkan

Starting from v0.8.0, Taichi adds a new UI system GGUI. The new system uses GPU for rendering, making it much faster to render 3D scenes. That is why this new system gets its name as GGUI. This document describes the APIs that it provides.

:::caution IMPORTANT
If you choose Vulkan as backend, ensure that you install the Vulkan environment.
:::

:::note
It is recommended that you familiarize yourself with GGUI through the examples in examples/ggui_examples.
:::

:::note
The variables referenced in code snippets below are define like this:

python as-prelude:vars
vertices         = ti.Vector.field(2, ti.f32, shape=200)
vertices_3d      = ti.Vector.field(3, ti.f32, shape=200)
indices          = ti.field(ti.i32, shape=200 * 3)
normals          = ti.Vector.field(3, ti.f32, shape=200)
per_vertex_color = ti.Vector.field(3, ti.f32, shape=200)

color  = (0.5, 0.5, 0.5)

:::

Create a window

ti.ui.Window(name, res) creates a window.

python preludes:vars
window = ti.ui.Window(name='Window Title', res = (640, 360), fps_limit=200, pos = (150, 150))
  • The name parameter sets the title of the window.
  • The res parameter specifies the resolution (width and height) of the window.
  • The fps_limit parameter sets the maximum frames per second (FPS) for the window.
  • The pos parameter specifies the position of the window with respect to the top-left corner of the main screen.

A ti.ui.Window can display three types of objects:

  • 2D Canvas, which is used to draw simple 2D geometries like circles and triangles.+ 3D Scene, which is used to render 3D meshes and particles, and provides configurable camera and light sources.
  • Immediate mode GUI components, such as buttons and textboxes.

2D Canvas

Create a canvas

The following code retrieves a Canvas object that covers the entire window.

python cont
canvas = window.get_canvas()

Draw on the canvas

python cont
canvas.set_background_color(color)
canvas.triangles(vertices, color, indices, per_vertex_color)

radius = 5
canvas.circles(vertices, radius, color, per_vertex_color)

width = 2
canvas.lines(vertices, width, indices, color, per_vertex_color)
canvas.set_image(window.get_image_buffer_as_numpy())

The arguments vertices, indices, per_vertex_color, and image must be Taichi fields. If per_vertex_color is provided, color is ignored.

The positions/centers of geometries are represented as floats between 0.0 and 1.0, which indicate the relative positions of the geometries on the canvas. For circles() and lines(), the radius and width arguments are relative to the height of the window.

The canvas is cleared after every frame. Always call these methods within the render loop.

3D Scene

Create a scene

python cont
scene = window.get_scene()

Configure camera

python cont
camera = ti.ui.Camera()
camera.position(1, 2, 3)  # x, y, z
camera.lookat(4, 5, 6)
camera.up(0, 1, 0)
camera.projection_mode(ti.ui.ProjectionMode.Perspective)
scene.set_camera(camera)

Configuring light sources

Add a point light

Call point_light() to add a point light to the scene.

python cont
scene.point_light(pos=(1, 2, 3), color=(0.5, 0.5, 0.5))

Note that you need to call point_light() for every frame. Similar to the canvas() methods, call this method within your render loop.

3D Geometries

python cont
scene.lines(vertices, width, indices, color, per_vertex_color)
scene.mesh(vertices_3d, indices, normals, color, per_vertex_color)
scene.particles(vertices_3d, radius, color, per_vertex_color)

The arguments vertices, indices, per_vertex_color, and image are all expected to be Taichi fields. If per_vertex_color is provided, color is ignored.

The positions/centers of geometries should be in the world-space coordinates.

:::note

If a mesh has num triangles, the indices should be a 1D scalar field with a shape (num * 3), not a vector field.

normals is an optional parameter for scene.mesh().

:::example

  1. An example of drawing 3d-lines
python
import taichi as ti

ti.init(arch=ti.cuda)

N = 10

particles_pos = ti.Vector.field(3, dtype=ti.f32, shape = N)
points_pos = ti.Vector.field(3, dtype=ti.f32, shape = N)

@ti.kernel
def init_points_pos(points : ti.template()):
    for i in range(points.shape[0]):
        points[i] = [i for j in ti.static(range(3))]

init_points_pos(particles_pos)
init_points_pos(points_pos)

window = ti.ui.Window("Test for Drawing 3d-lines", (768, 768))
canvas = window.get_canvas()
scene = window.get_scene()
camera = ti.ui.Camera()
camera.position(5, 2, 2)

while window.running:
    camera.track_user_inputs(window, movement_speed=0.03, hold_key=ti.ui.RMB)
    scene.set_camera(camera)
    scene.ambient_light((0.8, 0.8, 0.8))
    scene.point_light(pos=(0.5, 1.5, 1.5), color=(1, 1, 1))

    scene.particles(particles_pos, color = (0.68, 0.26, 0.19), radius = 0.1)
    # Draw 3d-lines in the scene
    scene.lines(points_pos, color = (0.28, 0.68, 0.99), width = 5.0)
    canvas.scene(scene)
    window.show()

Advanced 3d Geometries

python preludes:vars
window = ti.ui.Window(name='Advanced 3d Geometries', res = (720, 720))
scene = window.get_scene()
width = 2
radius = 5

scene.lines(vertices, width, indices, color, per_vertex_color, vertex_offset=0, vertex_count=10, index_offset=0, index_count=10)

scene.mesh(vertices_3d, indices, normals, color, per_vertex_color, vertex_offset=0, vertex_count=10, index_offset=0, index_count=10, show_wireframe=True)

scene.particles(vertices_3d, radius, color, per_vertex_color, index_offset=0, index_count=10)

scene.mesh_instance(vertices_3d, indices, normals, color, per_vertex_color, vertex_offset=0, vertex_count=10, index_offset=0, index_count=10, show_wireframe=True)

The additional arguments vertex_offset, vertex_count, index_offset and index_count control the visible part of the particles and mesh. For the mesh() and mesh_instance() methods, set whether to show wireframe mode through setting show_wireframe.

:::example

  1. Example of drawing a part of the mesh/particles
python cont
scene = window.get_scene()

center = ti.Vector.field(3, ti.f32, shape=10)

# For particles
# draw the 2-th to 7-th particles
scene.particles(center, radius=1, index_offset = 1, index_count = 6)

# For mesh
# 1. with indices
scene.mesh(
    vertices_3d, indices, index_offset=1, index_count=3,
    # vertex_offset is set to 0 by default, and it is not necessary
    # to assign vertex_offset a value that otherwise you must.
    vertex_offset = 1
    )

# usually used as below:
# draw the 11-th to 111-th mesh vertexes
scene.mesh(vertices_3d, indices, index_offset=10, index_count=100)

# 2. without indices (similar to the particles' example above)
scene.mesh(
    vertices_3d,
    vertex_offset=2,  # user defined first vertex index
    vertex_count=3,  # user defined vertex count
    )
  1. An example of drawing part of lines
text
/* Detailed source-code truncated for AI context efficiency. */
  1. Details of mesh instancing
python preludes:vars
window = ti.ui.Window("Display Instanced Mesh", (1024, 1024))
scene = window.get_scene()

num_instance = 100
m_transforms = ti.Matrix.field(4, 4, dtype = ti.f32, shape = num_instance)


# For example: An object is scaled by 2, rotated by rotMat, and translated by t = [1, 2, 3], then
#
# The ScaleMatrix is:
# 2, 0, 0, 0
# 0, 2, 0, 0
# 0, 0, 2, 0
# 0, 0, 0, 1
#
# The RotationMatrix is:
# https://en.wikipedia.org/wiki/Rotation_matrix#General_rotations
#
# The TranslationMatrix is:
# 1, 0, 0, 1
# 0, 1, 0, 2
# 0, 0, 1, 3
# 0, 0, 0, 1
#
# Let TransformMatrix = TranslationMatrix @ RotationMatrix @ ScaleMatrix, then the final TransformMatrix is:
#   2 * rotMat00,     rotMat01,       rotMat02, 1
#       rotMat10, 2 * rotMat11,       rotMat12, 2
#       rotMat20,     rotMat21,   2 * rotMat22, 3
#              0,            0,              0, 1
...

# Draw mesh instances (from the 1st instance)
scene.mesh_instance(vertices_3d, indices, transforms = m_transforms, instance_offset = 1)
  1. Example of setting wireframe mode
python preludes:vars
window = ti.ui.Window("Display Mesh", (1024, 1024), vsync=True)
canvas = window.get_canvas()
scene = window.get_scene()
camera = ti.ui.Camera()

# slider_int usage
some_int_type_value = 0
def show_options():
    global some_int_type_value

    window.GUI.begin("Display Panel", 0.05, 0.1, 0.2, 0.15)
    display_mode = window.GUI.slider_int("Value Range", some_int_type_value, 0, 5)
    window.GUI.end()

while window.running:

    ...
    # if to show wireframe
    scene.mesh_instance(vertices_3d, indices, instance_count = 100 , show_wireframe = True)

    canvas.scene(scene)
    show_options()
    window.show()

:::note

If indices is not provided, consider using like this:

python preludes:vars skip-ci:taichi-bug
scene = window.get_scene()
scene.mesh(vertices_3d, normals, color, per_vertex_color, vertex_offset=0, vertex_count=50, show_wireframe=True)

If indices is provided, consider using like this:

python cont
scene.mesh(vertices_3d, indices, normals, color, per_vertex_color, vertex_offset=0, index_offset=0, index_count=50, show_wireframe=True)

:::

Rendering the scene

You can render a scene on a canvas.

python cont
window = ti.ui.Window(name='Title', res=(640, 360))
canvas = window.get_canvas()
canvas.scene(scene)

Fetching Color/Depth information

python cont
img = window.get_image_buffer_as_numpy()
window.get_depth_buffer(scene_depth)
depth = window.get_depth_buffer_as_numpy()

After rendering the current scene, you can fetch the color and depth information of the current scene using get_image_buffer_as_numpy() and get_depth_buffer_as_numpy(), which copy the gpu data to a NumPy array(cpu).
get_depth_buffer() copies the GPU data to a Taichi field (depend on the arch you choose) or copies data from GPU to GPU.

:::example

  1. Example of fetching color information
python skip-ci:trivial
window = ti.ui.Window("Test for getting image buffer from ggui", (768, 768), vsync=True)
video_manager = ti.tools.VideoManager("OutputDir")

while window.running:
    # render_scene()
    img = window.get_image_buffer_as_numpy()
    video_manager.write_frame(img)
    window.show()

video_manager.make_video(gif=True, mp4=True)
  1. An example of fetching the depth data
python
window_shape = (720, 1080)
window = ti.ui.Window("Test for copy depth data", window_shape)
canvas = window.get_canvas()
scene = window.get_scene()
camera = ti.ui.Camera()

# Get the shape of the window
w, h = window.get_window_shape()
# The field/ndarray stores the depth information, and must be of the ti.f32 data type and have a 2d shape.
# or, in other words, the shape must equal the window's shape
scene_depth = ti.ndarray(ti.f32, shape = (w, h))
# scene_depth = ti.field(ti.f32, shape = (w, h))

while window.running:
    # render()
    canvas.scene(scene)
    window.get_depth_buffer(scene_depth)
    window.show()

GUI components

The design of GGUI's GUI components follows the Dear ImGui APIs.

python
window = ti.ui.Window("Test for GUI", res=(512, 512))
gui = window.get_gui()
value = 0
color = (1.0, 1.0, 1.0)
with gui.sub_window("Sub Window", x=10, y=10, width=300, height=100):
    gui.text("text")
    is_clicked = gui.button("name")
    value = gui.slider_float("name1", value, minimum=0, maximum=100)
    color = gui.color_edit_3("name2", color)

Show a window

Call show() to show a window.

python cont
window.show()

Call this method only at the end of the render loop for each frame.

User input processing

To retrieve the events that have occurred since the last method call:

python cont
events = window.get_events()

Each event in events is an instance of ti.ui.Event. It has the following properties:

  • event.action, which can be ti.ui.PRESS, ti.ui.RELEASE, or ti.ui.MOTION.
  • event.key: the key related to this event.

To retrieve the mouse position:

  • window.get_cursor_pos()

To check if a key is pressed:

  • window.is_pressed(key)

The following is a user input processing example from mpm128:

python cont
gravity = ti.Vector.field(2, ti.f32, shape=())
attractor_strength = ti.field(ti.f32, shape=())

while window.running:
    # keyboard event processing
    if window.get_event(ti.ui.PRESS):
        if window.event.key == 'r': reset()
        elif window.event.key in [ti.ui.ESCAPE]: break
    if window.event is not None: gravity[None] = [0, 0]  # if had any event
    if window.is_pressed(ti.ui.LEFT, 'a'): gravity[None][0] = -1
    if window.is_pressed(ti.ui.RIGHT, 'd'): gravity[None][0] = 1
    if window.is_pressed(ti.ui.UP, 'w'): gravity[None][1] = 1
    if window.is_pressed(ti.ui.DOWN, 's'): gravity[None][1] = -1

    # mouse event processing
    mouse = window.get_cursor_pos()
    # ...
    if window.is_pressed(ti.ui.LMB):
        attractor_strength[None] = 1
    if window.is_pressed(ti.ui.RMB):
        attractor_strength[None] = -1

    window.show()

Image I/O

To write the current frame in the window to an image file:

python cont
window.save_image('frame.png')

Note that you must call window.save_image() before calling window.show().

Off-screen rendering

GGUI supports saving frames to images without showing the window. This is also known as "headless" rendering. To enable this mode, set the argument show_window to False when initializing a window.

python
window = ti.ui.Window('Window Title', (640, 360), show_window = False)

Then you can call window.save_image() as normal and remove the window.show() call at the end.

File: docs/lang/articles/visualization/gui_system.md

GUI System

Taichi has a built-in GUI system for visualizing simulation data in data containers like Taichi fields or NumPy ndarrays. It also has limited support for drawing primitive geometries.

Create and display a window

The following code creates a 640x360 window with a "Hello World!" title:

python as-prelude:gui
gui = ti.GUI('Hello World!', (640, 360))

Displays it by calling gui.show():

python preludes:gui
while gui.running:
    gui.show()

:::note

Call gui.show() inside a while loop. Otherwise, the window would flash once and disappear.

:::

Close the window

You can set gui.running=False in the while loop to close the GUI:

python
gui = ti.GUI('Window Title', (640, 360))
some_events_happend = lambda: random.random() < 0.8

while gui.running:
    if some_events_happend():
        gui.running = False
    gui.show()

Coordinate system

Each window is built on a coordinate system: the origin is located in the lower-left corner, with the +x direction stretching to the right and the +y direction stretching upward.

Display a field or ndarray

To display a Taichi field or a NumPy ndarray, call gui.set_image(). The method accepts both types as input.

python
gui = ti.GUI('Set Image', (640, 480))
image = ti.Vector.field(3, ti.f32, shape=(640, 480))
while gui.running:
    gui.set_image(image)
    gui.show()

Because Taichi field is a global data container, if the vector field image is updated between the while loops, the GUI window refreshes to display the latest image.

:::caution IMPORTANT

Ensure that the shape of the input matches the resolution of the GUI window.

:::

Zero-copying frame buffer

In each loop of the gui.set_image() method call, the GUI system converts the image data to a displayable format and copies the result to the window buffer. This causes huge overload when the window size is large, making it hard to achieve high FPS (frames per second).

If you only need to call the set_image() method without using any drawing command, you can enable fast_gui mode for better performance. This mode allows Taichi GUI to write the image data directly to the frame buffer without additional copying, and significantly increases FPS.

python
gui = ti.GUI('Fast GUI', res=(400, 400), fast_gui=True)

For this mode to work, ensure that the data passed into gui.set_image() is in a display-compatible format. In other words, If it is a Taichi field, ensure that it is one of the following:

  • a vector field ti.field(3, dtype, shape) compatible with RGB format.
  • a vector field ti.field(4, dtype, shape) compatible with RGBA format.

Note that dtype must be ti.f32, ti.f64, or ti.u8.

Draw on a window

Taichi's GUI system supports drawing simple geometries, such as lines, circles, triangles, rectangles, arrows, and texts.

Single geometry

In Taichi, drawing basic geometric shapes on the GUI is very intuitive. In most cases, all we need to do is specify information such as the position and size of the geometry and call the corresponding APIs.

Line

You can draw a single line on a GUI canvas by specifying its begin and end points:

python
import numpy as np
gui = ti.GUI('Single Line', res=(400, 400))
begin = [0.1, 0.1]
end = [0.9, 0.9]
while gui.running:
    gui.line(begin, end, radius=1, color=0x068587)
    gui.show()

:::note

Coordinates such as begin and end for single geometry can be Python lists, Numpy arrays or ti.Vector, as long as it's subscriptable and its dimension is (2, ).

:::

Circle

You can draw a single circle on a GUI canvas by specifying its center poistion and its radius:

python
import numpy as np
gui = ti.GUI('Single Circle', res=(400, 400))
center = [0.5, 0.5]
while gui.running:
    gui.circle(pos=center, radius=30, color=0xED553B)
    gui.show()

Triangle

You can draw a single triangle on a GUI canvas by specifying its three end points:

python
import numpy as np
gui = ti.GUI('Single Triangle', res=(400, 400))
p1 = [0.5, 0.5]
p2 = [0.6, 0.5]
p3 = [0.5, 0.6]
while gui.running:
    gui.triangle(a=p1, b=p2, c=p3, color=0xEEEEF0)
    gui.show()

Rectangle

You can draw a single rectangle on a GUI canvas by specifying its topleft and bottomright points:

python
import numpy as np
gui = ti.GUI('Single Rectangle', res=(400, 400))
p1 = [0.3, 0.4]
p2 = [0.7, 0.6]
while gui.running:
    gui.rect(topleft=p1, bottomright=p2, color=0xFFFFFF)
    gui.show()

Arrow

You can draw a single arrow on a GUI canvas by specifying its start point and direction:

python
import numpy as np
gui = ti.GUI('Single Arrow', res=(400, 400))
begin = [0.3, 0.3]
increment = [0.5, 0.5]
while gui.running:
    gui.arrow(orig=begin, direction=increment, color=0xFFFFFF)
    gui.show()

Text

You can draw a single line of text on a GUI canvas by specifying its position and contents:

python
gui = ti.GUI('Text', res=(400, 400))
position = [0.3, 0.5]
while gui.running:
    gui.text(content='Hello Taichi', pos=position, font_size=34, color=0xFFFFFF)
    gui.show()

Multiple geometries

It's also possible to draw multiple geometries at once by providing a collection of their positions to the GUI. The pos parameter of every drawing method accepts Taichi fields or NumPy arrays, not Python primitive lists. Each element of the array is a pair of floats ranging from 0.0 to 1.0, which represent the relative positions of the geometries. For example:

  • (0.0, 0.0): the lower-left corner of the window.
  • (1.0, 1.0): the upper-right corner of the window.

Lines

The following code draws five blue line segments whose width is 2, with X and Y representing the five starting points and the five ending points.

python
import numpy as np
X = np.random.random((5, 2))
Y = np.random.random((5, 2))
gui = ti.GUI("lines", res=(400, 400))
while gui.running:
    gui.lines(begin=X, end=Y, radius=2, color=0x068587)
    gui.show()

Circles

The following code draws 50 circles with a radius of 5 and in three different colors randomly assigned by indices, an integer array of the same size as pos.

python
import numpy as np
pos = np.random.random((50, 2))
# Create an array of 50 integer elements whose values are randomly 0, 1, 2
# 0 corresponds to 0x068587
# 1 corresponds to 0xED553B
# 2 corresponds to 0xEEEEF0
indices = np.random.randint(0, 2, size=(50,))
gui = ti.GUI("circles", res=(400, 400))
while gui.running:
    gui.circles(pos, radius=5, palette=[0x068587, 0xED553B, 0xEEEEF0], palette_indices=indices)
    gui.show()

Triangles

The following code draws two orange triangles orange, with X, Y, and Z representing the three points of the triangles.

python
import numpy as np
X = np.random.random((2, 2))
Y = np.random.random((2, 2))
Z = np.random.random((2, 2))
gui = ti.GUI("triangles", res=(400, 400))
while gui.running:
    gui.triangles(a=X, b=Y, c=Z, color=0xED553B)
    gui.show()

Arrows

The following code generates 100 random sized arrows, with begins and direction represents their begin points and incrementals:

python
import numpy as np
begins = np.random.random((100, 2))
directions = np.random.uniform(low=-0.05, high=0.05, size=(100, 2))
gui = ti.GUI('arrows', res=(400, 400))
while gui.running:
    gui.arrows(orig=begins, direction=directions, radius=1)
    gui.show()

Notice that we used low and high in the call to np.random.uniform() to limit the range of generated random numbers.

Event handling

Taichi's GUI system also provides a set of methods for mouse and keyboard control. Input events are classified into three types:

python
ti.GUI.RELEASE  # key up or mouse button up
ti.GUI.PRESS    # key down or mouse button down
ti.GUI.MOTION   # mouse motion or mouse wheel

Event key is the key that you press from your keyboard or mouse. It can be one of:

python
# for ti.GUI.PRESS and ti.GUI.RELEASE event:
ti.GUI.ESCAPE  # Esc
ti.GUI.SHIFT   # Shift
ti.GUI.LEFT    # Left Arrow
'a'            # we use lowercase for alphabet
'b'
...
ti.GUI.LMB     # Left Mouse Button
ti.GUI.RMB     # Right Mouse Button

# for ti.GUI.MOTION event:
ti.GUI.MOVE    # Mouse Moved
ti.GUI.WHEEL   # Mouse Wheel Scrolling

An event filter is a combined list of key, type, and (type, key) tuple. For example:

python preludes:gui
# if ESC pressed or released:
gui.get_event(ti.GUI.ESCAPE)

# if any key is pressed:
gui.get_event(ti.GUI.PRESS)

# if ESC is pressed or SPACE is released:
gui.get_event((ti.GUI.PRESS, ti.GUI.ESCAPE), (ti.GUI.RELEASE, ti.GUI.SPACE))

gui.get_event() pops an event from the queue and saves it to gui.event. For example:

python preludes:gui
if gui.get_event():
    print('Got event, key =', gui.event.key)

The following code defines that the while loop goes on until ESC is pressed:

python preludes:gui
while gui.running:
    if gui.get_event(ti.GUI.ESCAPE):
        break
    gui.show()

gui.is_pressed() detects the pressed keys. As the following code snippet shows, you must use it together with gui.get_event(). Otherwise, it is not updated.

For example:

python preludes:gui
while gui.running:
    gui.get_event()  # must be called before is_pressed
    if gui.is_pressed('a', ti.GUI.LEFT):
        print('Go left!')
    elif gui.is_pressed('d', ti.GUI.RIGHT):
        print('Go right!')
    gui.show()

:::caution

Call gui.get_event() before calling gui.is_pressed(). Otherwise, gui.is_pressed() does not take effect.

:::

Retrieve cursor position

gui.get_cursor_pos() returns the cursor's current position in the window. The return value is a pair of floats in the range [0.0, 1.0]. For example:

python preludes:gui
mouse_x, mouse_y = gui.get_cursor_pos()

GUI Widgets

Taichi's GUI system also provides widgets, including slider(), label(), and button(), for you to customize your control interface. Take a look at the following code snippet:

python
import taichi as ti
gui = ti.GUI('GUI widgets')

radius = gui.slider('Radius', 1, 50, step=1)
xcoor = gui.label('X-coordinate')
okay = gui.button('OK')

xcoor.value = 0.5
radius.value = 10

while gui.running:
    for e in gui.get_events(gui.PRESS):
        if e.key == gui.ESCAPE:
            gui.running = False
        elif e.key == 'a':
            xcoor.value -= 0.05
        elif e.key == 'd':
            xcoor.value += 0.05
        elif e.key == 's':
            radius.value -= 1
        elif e.key == 'w':
            radius.value += 1
        elif e.key == okay:
            print('OK clicked')

    gui.circle((xcoor.value, 0.5), radius=radius.value)
    gui.show()

Explore Related Technologies & Topics

View All 9,111+ Repositories โ†’

Discover similar high-velocity repositories, agent skills, and OpenAPI specifications across the ecosystem.

#computer-graphics Browse computer-graphics catalog โ†’ #differentiable-programming Browse differentiable-programming catalog โ†’ #gpu Browse gpu catalog โ†’ #gpu-programming Browse gpu-programming catalog โ†’ #sparse-computation Browse sparse-computation catalog โ†’ #taichi Browse taichi catalog โ†’ โšก MCP Servers Model Context Protocol โ†’ ๐Ÿง  Agent Skills Prompt & rule files โ†’ ๐Ÿ”Œ OpenAPI Specs REST & Swagger schemas โ†’