Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[opengl] Add ti.gles arch and enable tests #6988

Merged
merged 1 commit into from
Dec 27, 2022
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions docs/lang/articles/deployment/ndarray_android.md
Original file line number Diff line number Diff line change
Expand Up @@ -66,7 +66,7 @@ The following Python script defines a Taichi AOT module for generating and savin
```python
import taichi as ti

ti.init(arch=ti.opengl, use_gles=True, allow_nv_shader_extension=False)
ti.init(arch=ti.gles, allow_nv_shader_extension=False)

# Define constants for computation
G = 1
Expand Down Expand Up @@ -126,7 +126,7 @@ aot()

**In line 3, you initialize Taichi:**

1. Set `use_gles` to `True` to generate GLES compute shaders for Android.
1. Set `arch=ti.gles` to generate GLES compute shaders for Android.
2. Set `allow_nv_shader_extension` to `False` to prevent the generated GLES compute shaders from using Nvidia GL extensions on Android.

> This setting is because Android supports GLES APIs but GLES does not support `NV_SHADER_EXTENSION`.
Expand Down
3 changes: 2 additions & 1 deletion python/taichi/lang/impl.py
Original file line number Diff line number Diff line change
Expand Up @@ -627,7 +627,8 @@ def create_field_member(dtype, name, needs_grad, needs_dual):
# adjoint checkbit
x_grad_checkbit = Expr(get_runtime().prog.make_id_expr(""))
dtype = u8
if prog.config().arch in (_ti_core.opengl, _ti_core.vulkan):
if prog.config().arch in (_ti_core.opengl, _ti_core.vulkan,
_ti_core.gles):
dtype = i32
x_grad_checkbit.ptr = _ti_core.expr_field(x_grad_checkbit.ptr,
cook_dtype(dtype))
Expand Down
25 changes: 14 additions & 11 deletions python/taichi/lang/misc.py
Original file line number Diff line number Diff line change
Expand Up @@ -129,6 +129,11 @@
"""
# ----------------------

gles = _ti_core.gles
"""The OpenGL ES backend. OpenGL ES 3.1 required.
"""
# ----------------------

# Skip annotating this one because it is barely maintained.
cc = _ti_core.cc

Expand All @@ -154,9 +159,9 @@
"""
# ----------------------

gpu = [cuda, metal, vulkan, opengl, dx11, dx12]
gpu = [cuda, metal, vulkan, opengl, dx11, dx12, gles]
"""A list of GPU backends supported on the current system.
Currently contains 'cuda', 'metal', 'opengl', 'vulkan', 'dx11', 'dx12'.
Currently contains 'cuda', 'metal', 'opengl', 'vulkan', 'dx11', 'dx12', 'gles'.
When this is used, Taichi automatically picks the matching GPU backend. If no
GPU is detected, Taichi falls back to the CPU backend.
Expand Down Expand Up @@ -454,7 +459,7 @@ def init(arch=None,
if env_arch is not None:
_logging.info(f'Following TI_ARCH setting up for arch={env_arch}')
arch = _ti_core.arch_from_name(env_arch)
cfg.arch = adaptive_arch_select(arch, enable_fallback, cfg.use_gles)
cfg.arch = adaptive_arch_select(arch, enable_fallback)
if cfg.arch == cc:
_ti_core.set_tmp_dir(locale_encode(prepare_sandbox()))
print(f'[Taichi] Starting on arch={_ti_core.arch_name(cfg.arch)}')
Expand Down Expand Up @@ -717,14 +722,11 @@ def mesh_patch_idx():
)


def is_arch_supported(arch, use_gles=False):
def is_arch_supported(arch):
"""Checks whether an arch is supported on the machine.
Args:
arch (taichi_python.Arch): Specified arch.
use_gles (bool): If True, check is GLES is available otherwise
check if GLSL is available. Only effective when `arch` is `ti.opengl`.
Default is `False`.
Returns:
bool: Whether `arch` is supported on the machine.
Expand All @@ -733,7 +735,8 @@ def is_arch_supported(arch, use_gles=False):
arch_table = {
cuda: _ti_core.with_cuda,
metal: _ti_core.with_metal,
opengl: functools.partial(_ti_core.with_opengl, use_gles),
opengl: functools.partial(_ti_core.with_opengl, False),
gles: functools.partial(_ti_core.with_opengl, True),
cc: _ti_core.with_cc,
vulkan: _ti_core.with_vulkan,
dx11: _ti_core.with_dx11,
Expand All @@ -753,13 +756,13 @@ def is_arch_supported(arch, use_gles=False):
return False


def adaptive_arch_select(arch, enable_fallback, use_gles):
def adaptive_arch_select(arch, enable_fallback):
if arch is None:
return cpu
if not isinstance(arch, (list, tuple)):
arch = [arch]
for a in arch:
if is_arch_supported(a, use_gles):
if is_arch_supported(a):
return a
if not enable_fallback:
raise RuntimeError(f'Arch={arch} is not supported')
Expand All @@ -778,7 +781,7 @@ def get_compute_stream_device_time_elapsed_us() -> float:
__all__ = [
'i', 'ij', 'ijk', 'ijkl', 'ijl', 'ik', 'ikl', 'il', 'j', 'jk', 'jkl', 'jl',
'k', 'kl', 'l', 'x86_64', 'x64', 'dx11', 'dx12', 'wasm', 'arm64', 'cc',
'cpu', 'cuda', 'gpu', 'metal', 'opengl', 'vulkan', 'extension',
'cpu', 'cuda', 'gles', 'gpu', 'metal', 'opengl', 'vulkan', 'extension',
'loop_config', 'global_thread_idx', 'assume_in_range', 'block_local',
'cache_read_only', 'init', 'mesh_local', 'no_activate', 'reset',
'mesh_patch_idx', 'get_compute_stream_device_time_elapsed_us'
Expand Down
3 changes: 1 addition & 2 deletions taichi/analysis/offline_cache_util.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -52,9 +52,8 @@ static std::vector<std::uint8_t> get_offline_cache_key_of_compile_config(
if (config->arch == Arch::cc) {
serializer(config->cc_compile_cmd);
serializer(config->cc_link_cmd);
} else if (config->arch == Arch::opengl) {
} else if (config->arch == Arch::opengl || config->arch == Arch::gles) {
serializer(config->allow_nv_shader_extension);
serializer(config->use_gles);
}
serializer(config->make_mesh_block_local);
serializer(config->optimize_mesh_reordered_mapping);
Expand Down
4 changes: 4 additions & 0 deletions taichi/aot/module_loader.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -37,6 +37,10 @@ std::unique_ptr<Module> Module::load(Arch arch, std::any mod_params) {
} else if (arch == Arch::opengl) {
#ifdef TI_WITH_OPENGL
return gfx::make_aot_module(mod_params, arch);
#endif
} else if (arch == Arch::gles) {
#ifdef TI_WITH_OPENGL
return gfx::make_aot_module(mod_params, arch);
#endif
} else if (arch == Arch::dx11) {
#ifdef TI_WITH_DX11
Expand Down
1 change: 1 addition & 0 deletions taichi/inc/archs.inc.h
Original file line number Diff line number Diff line change
Expand Up @@ -16,3 +16,4 @@ PER_ARCH(dx12) // Microsoft DirectX 12, WIP
PER_ARCH(opencl) // OpenCL, N/A
PER_ARCH(amdgpu) // AMD GPU, N/A
PER_ARCH(vulkan) // Vulkan
PER_ARCH(gles) // OpenGL ES
1 change: 0 additions & 1 deletion taichi/program/compile_config.h
Original file line number Diff line number Diff line change
Expand Up @@ -77,7 +77,6 @@ struct CompileConfig {

// Opengl backend options:
bool allow_nv_shader_extension{true};
bool use_gles{false};

bool quant_opt_store_fusion{true};
bool quant_opt_atomic_demotion{true};
Expand Down
1 change: 1 addition & 0 deletions taichi/program/extension.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -23,6 +23,7 @@ bool is_extension_supported(Arch arch, Extension ext) {
{Arch::metal,
{Extension::adstack, Extension::assertion, Extension::sparse}},
{Arch::opengl, {Extension::extfunc}},
{Arch::gles, {}},
{Arch::cc, {Extension::data64, Extension::extfunc, Extension::adstack}},
};
// if (with_opengl_extension_data64())
Expand Down
10 changes: 9 additions & 1 deletion taichi/program/program.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -123,7 +123,14 @@ Program::Program(Arch desired_arch) : snode_rw_accessors_bank_(this) {
#endif
} else if (config.arch == Arch::opengl) {
#ifdef TI_WITH_OPENGL
TI_ASSERT(opengl::initialize_opengl(config.use_gles));
TI_ASSERT(opengl::initialize_opengl(false));
program_impl_ = std::make_unique<OpenglProgramImpl>(config);
#else
TI_ERROR("This taichi is not compiled with OpenGL");
#endif
} else if (config.arch == Arch::gles) {
#ifdef TI_WITH_OPENGL
TI_ASSERT(opengl::initialize_opengl(true));
program_impl_ = std::make_unique<OpenglProgramImpl>(config);
#else
TI_ERROR("This taichi is not compiled with OpenGL");
Expand Down Expand Up @@ -552,6 +559,7 @@ std::unique_ptr<AotModuleBuilder> Program::make_aot_module_builder(
this_thread_config().arch == Arch::metal ||
this_thread_config().arch == Arch::vulkan ||
this_thread_config().arch == Arch::opengl ||
this_thread_config().arch == Arch::gles ||
this_thread_config().arch == Arch::dx12) {
return program_impl_->make_aot_module_builder(cfg);
}
Expand Down
1 change: 0 additions & 1 deletion taichi/python/export_lang.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -209,7 +209,6 @@ void export_lang(py::module &m) {
&CompileConfig::quant_opt_atomic_demotion)
.def_readwrite("allow_nv_shader_extension",
&CompileConfig::allow_nv_shader_extension)
.def_readwrite("use_gles", &CompileConfig::use_gles)
.def_readwrite("make_mesh_block_local",
&CompileConfig::make_mesh_block_local)
.def_readwrite("mesh_localize_to_end_mapping",
Expand Down
6 changes: 4 additions & 2 deletions taichi/transforms/offload.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -112,7 +112,8 @@ class Offloader {
offloaded->const_end = true;
offloaded->end_value = val->val.val_int32();
} else {
if ((arch == Arch::opengl || arch == Arch::vulkan) &&
if ((arch == Arch::opengl || arch == Arch::vulkan ||
arch == Arch::gles) &&
demotable_axis_load(s->end)) {
// TODO: We need to update codegen for each backend gradually so
// let's limit it to opengl backend for now.
Expand Down Expand Up @@ -385,7 +386,8 @@ class IdentifyValuesUsedInOtherOffloads : public BasicStmtVisitor {
if (top_level_ptr->is<GlobalPtrStmt>() || stmt->is<ExternalPtrStmt>() ||
(stmt->is<ArgLoadStmt>() && stmt->as<ArgLoadStmt>()->is_ptr))
return;
if ((config_.arch == Arch::opengl || config_.arch == Arch::vulkan) &&
if ((config_.arch == Arch::opengl || config_.arch == Arch::vulkan ||
config_.arch == Arch::gles) &&
demotable_axis_load(stmt))
return;
// Not yet allocated
Expand Down
3 changes: 2 additions & 1 deletion taichi/util/offline_cache.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -22,7 +22,8 @@ std::string get_cache_path_by_arch(const std::string &base_path, Arch arch) {
std::string subdir;
if (arch_uses_llvm(arch)) {
subdir = kLlvmCachSubPath;
} else if (arch == Arch::vulkan || arch == Arch::opengl) {
} else if (arch == Arch::vulkan || arch == Arch::opengl ||
arch == Arch::gles) {
subdir = kSpirvCacheSubPath;
} else if (arch == Arch::metal) {
subdir = kMetalCacheSubPath;
Expand Down
4 changes: 2 additions & 2 deletions tests/python/test_ad_grad_check.py
Original file line number Diff line number Diff line change
Expand Up @@ -6,7 +6,7 @@


@test_utils.test(default_fp=ti.f64,
exclude=[ti.cc, ti.vulkan, ti.opengl, ti.metal])
exclude=[ti.cc, ti.vulkan, ti.opengl, ti.gles, ti.metal])
def test_general():
x1 = ti.field(dtype=float, shape=(2, 2), needs_grad=True)
y1 = ti.field(dtype=float, shape=(), needs_grad=True)
Expand Down Expand Up @@ -71,6 +71,6 @@ def func():
lambda x: ti.atan2(x, 0.4), lambda x: 0.4**x, lambda x: x**0.4
])
@test_utils.test(default_fp=ti.f64,
exclude=[ti.cc, ti.vulkan, ti.opengl, ti.metal])
exclude=[ti.cc, ti.vulkan, ti.opengl, ti.gles, ti.metal])
def test_basics(tifunc):
grad_test(tifunc)
6 changes: 2 additions & 4 deletions tests/python/test_aot.py
Original file line number Diff line number Diff line change
Expand Up @@ -138,10 +138,8 @@ def test_non_dense_snode():
m.add_field('y', y)


@pytest.mark.parametrize('use_gles', [True, False])
@test_utils.test(arch=[ti.opengl, ti.vulkan])
def test_mpm88_aot(use_gles):
ti.init(ti.lang.impl.current_cfg().arch, use_gles=use_gles)
@test_utils.test(arch=[ti.opengl, ti.gles, ti.vulkan])
def test_mpm88_aot():
n_particles = 8192
n_grid = 128
dx = 1 / n_grid
Expand Down
28 changes: 14 additions & 14 deletions tests/python/test_api.py
Original file line number Diff line number Diff line change
Expand Up @@ -72,20 +72,20 @@ def _get_expected_matrix_apis():
'deactivate_all_snodes', 'dx11', 'dx12', 'eig', 'exp', 'experimental',
'extension', 'f16', 'f32', 'f64', 'field', 'float16', 'float32', 'float64',
'floor', 'func', 'get_addr', 'get_compute_stream_device_time_elapsed_us',
'global_thread_idx', 'gpu', 'graph', 'grouped', 'hex_to_rgb', 'i', 'i16',
'i32', 'i64', 'i8', 'ij', 'ijk', 'ijkl', 'ijl', 'ik', 'ikl', 'il', 'init',
'int16', 'int32', 'int64', 'int8', 'is_active', 'is_logging_effective',
'j', 'jk', 'jkl', 'jl', 'k', 'kernel', 'kl', 'l', 'lang', 'length',
'linalg', 'log', 'loop_config', 'math', 'max', 'mesh_local',
'mesh_patch_idx', 'metal', 'min', 'ndarray', 'ndrange', 'no_activate',
'one', 'opengl', 'polar_decompose', 'pow', 'profiler', 'randn', 'random',
'raw_div', 'raw_mod', 'ref', 'rescale_index', 'reset', 'rgb_to_hex',
'root', 'round', 'rsqrt', 'select', 'set_logging_level', 'simt', 'sin',
'solve', 'sparse_matrix_builder', 'sqrt', 'static', 'static_assert',
'static_print', 'stop_grad', 'svd', 'swizzle_generator', 'sym_eig', 'sync',
'tan', 'tanh', 'template', 'tools', 'types', 'u16', 'u32', 'u64', 'u8',
'ui', 'uint16', 'uint32', 'uint64', 'uint8', 'vulkan', 'wasm', 'x64',
'x86_64', 'zero'
'gles', 'global_thread_idx', 'gpu', 'graph', 'grouped', 'hex_to_rgb', 'i',
'i16', 'i32', 'i64', 'i8', 'ij', 'ijk', 'ijkl', 'ijl', 'ik', 'ikl', 'il',
'init', 'int16', 'int32', 'int64', 'int8', 'is_active',
'is_logging_effective', 'j', 'jk', 'jkl', 'jl', 'k', 'kernel', 'kl', 'l',
'lang', 'length', 'linalg', 'log', 'loop_config', 'math', 'max',
'mesh_local', 'mesh_patch_idx', 'metal', 'min', 'ndarray', 'ndrange',
'no_activate', 'one', 'opengl', 'polar_decompose', 'pow', 'profiler',
'randn', 'random', 'raw_div', 'raw_mod', 'ref', 'rescale_index', 'reset',
'rgb_to_hex', 'root', 'round', 'rsqrt', 'select', 'set_logging_level',
'simt', 'sin', 'solve', 'sparse_matrix_builder', 'sqrt', 'static',
'static_assert', 'static_print', 'stop_grad', 'svd', 'swizzle_generator',
'sym_eig', 'sync', 'tan', 'tanh', 'template', 'tools', 'types', 'u16',
'u32', 'u64', 'u8', 'ui', 'uint16', 'uint32', 'uint64', 'uint8', 'vulkan',
'wasm', 'x64', 'x86_64', 'zero'
]
user_api[ti.ad] = [
'FwdMode', 'Tape', 'clear_all_gradients', 'grad_for', 'grad_replaced',
Expand Down
4 changes: 2 additions & 2 deletions tests/python/test_arg_alignment.py
Original file line number Diff line number Diff line change
Expand Up @@ -2,7 +2,7 @@
from tests import test_utils


@test_utils.test(exclude=[ti.opengl])
@test_utils.test(exclude=[ti.opengl, ti.gles])
def test_ret_write():
@ti.kernel
def func(a: ti.i16) -> ti.f32:
Expand All @@ -11,7 +11,7 @@ def func(a: ti.i16) -> ti.f32:
assert func(255) == 3.0


@test_utils.test(exclude=[ti.opengl])
@test_utils.test(exclude=[ti.opengl, ti.gles])
def test_arg_read():
x = ti.field(ti.i32, shape=())

Expand Down
2 changes: 1 addition & 1 deletion tests/python/test_argument.py
Original file line number Diff line number Diff line change
Expand Up @@ -26,7 +26,7 @@ def foo2(a: ti.i32, b: ti.i32, c: ti.i32, d: ti.i32, e: ti.i32, f: ti.i32,
foo2(1, 2, 3, 4, 5, 6, 7, 8, 9)


@test_utils.test(exclude=[ti.opengl, ti.cc])
@test_utils.test(exclude=[ti.opengl, ti.gles, ti.cc])
def test_exceed_max_64():
N = 64

Expand Down
4 changes: 2 additions & 2 deletions tests/python/test_cast.py
Original file line number Diff line number Diff line change
Expand Up @@ -5,7 +5,7 @@


@pytest.mark.parametrize('dtype', [ti.u8, ti.u16, ti.u32])
@test_utils.test(exclude=ti.opengl)
@test_utils.test(exclude=[ti.opengl, ti.gles])
def test_cast_uint_to_float(dtype):
@ti.kernel
def func(a: dtype) -> ti.f32:
Expand All @@ -19,7 +19,7 @@ def func_sugar(a: dtype) -> ti.f32:


@pytest.mark.parametrize('dtype', [ti.u8, ti.u16, ti.u32])
@test_utils.test(exclude=ti.opengl)
@test_utils.test(exclude=[ti.opengl, ti.gles])
def test_cast_float_to_uint(dtype):
@ti.kernel
def func(a: ti.f32) -> dtype:
Expand Down
2 changes: 1 addition & 1 deletion tests/python/test_cli.py
Original file line number Diff line number Diff line change
Expand Up @@ -213,7 +213,7 @@ def test_cli_run():


def test_cli_cache():
archs = {ti.cpu, ti.cuda, ti.opengl, ti.vulkan, ti.metal}
archs = {ti.cpu, ti.cuda, ti.opengl, ti.vulkan, ti.metal, ti.gles}
archs = {v for v in archs if v in test_utils.expected_archs()}
exts = ('ll', 'bc', 'spv', 'metal', 'tcb', 'lock')
tmp_path = tempfile.mkdtemp()
Expand Down
2 changes: 1 addition & 1 deletion tests/python/test_fields_builder.py
Original file line number Diff line number Diff line change
Expand Up @@ -196,7 +196,7 @@ def test_field_initialize_zero():
assert b[0] == 0


@test_utils.test(exclude=[ti.opengl, ti.cc])
@test_utils.test(exclude=[ti.opengl, ti.gles, ti.cc])
def test_field_builder_place_grad():
@ti.kernel
def mul(arr: ti.template(), out: ti.template()):
Expand Down
15 changes: 10 additions & 5 deletions tests/python/test_internal_func.py
Original file line number Diff line number Diff line change
Expand Up @@ -6,7 +6,8 @@
from tests import test_utils


@test_utils.test(exclude=[ti.metal, ti.opengl, ti.cuda, ti.vulkan, ti.cc])
@test_utils.test(
exclude=[ti.metal, ti.opengl, ti.gles, ti.cuda, ti.vulkan, ti.cc])
def test_basic():
@ti.kernel
def test():
Expand All @@ -16,7 +17,8 @@ def test():
test()


@test_utils.test(exclude=[ti.metal, ti.opengl, ti.cuda, ti.vulkan, ti.cc])
@test_utils.test(
exclude=[ti.metal, ti.opengl, ti.gles, ti.cuda, ti.vulkan, ti.cc])
def test_host_polling():
return

Expand All @@ -30,7 +32,8 @@ def test():
time.sleep(0.1)


@test_utils.test(exclude=[ti.metal, ti.opengl, ti.cuda, ti.vulkan, ti.cc])
@test_utils.test(
exclude=[ti.metal, ti.opengl, ti.gles, ti.cuda, ti.vulkan, ti.cc])
def test_list_manager():
@ti.kernel
def test():
Expand All @@ -40,7 +43,8 @@ def test():
test()


@test_utils.test(exclude=[ti.metal, ti.opengl, ti.cuda, ti.vulkan, ti.cc])
@test_utils.test(
exclude=[ti.metal, ti.opengl, ti.gles, ti.cuda, ti.vulkan, ti.cc])
def test_node_manager():
@ti.kernel
def test():
Expand All @@ -50,7 +54,8 @@ def test():
test()


@test_utils.test(exclude=[ti.metal, ti.opengl, ti.cuda, ti.vulkan, ti.cc])
@test_utils.test(
exclude=[ti.metal, ti.opengl, ti.gles, ti.cuda, ti.vulkan, ti.cc])
def test_node_manager_gc():
@ti.kernel
def test_cpu():
Expand Down
Loading