Morph MORPH
SIGN IN SIGN UP
taichi-dev / taichi UNCLAIMED

Productive, portable, and performant GPU programming in Python.

28123 0 0 C++
Normal View History Raw
[Doc] Document lazy_grad() function (#2456) * [doc] Document lazy_grad() function * Auto Format Co-authored-by: Taichi Gardener <taichigardener@gmail.com>
2021-06-23 17:12:25 +08:00
import taichi as ti
[ci] Move _testing.py into tests folder (#4247)
2022-02-10 12:37:36 +08:00
from tests import test_utils
[Doc] Document lazy_grad() function (#2456) * [doc] Document lazy_grad() function * Auto Format Co-authored-by: Taichi Gardener <taichigardener@gmail.com>
2021-06-23 17:12:25 +08:00
[type] [Bug] Fix bit array vectorization on GPU (#2120)
2020-12-26 12:31:21 +08:00
[ci] Move _testing.py into tests folder (#4247)
2022-02-10 12:37:36 +08:00
@test_utils.test(require=ti.extension.quant, debug=True, cfg_optimization=False)
[type] [Bug] Fix bit array vectorization on GPU (#2120)
2020-12-26 12:31:21 +08:00
def test_vectorized_struct_for():
[type] [refactor] Consistently use quant_xxx in quant-related names (#5166) * Rename CustomXXType to QuantXXType * Further rename custom to quant * Further rename * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Rename test_quant_time_integration Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
2022-06-15 09:43:49 +08:00
qu1 = ti.types.quant.int(1, False)
[type] [Bug] Fix bit array vectorization on GPU (#2120)
2020-12-26 12:31:21 +08:00
[type] [refactor] Consistently use quant_xxx in quant-related names (#5166) * Rename CustomXXType to QuantXXType * Further rename custom to quant * Further rename * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Rename test_quant_time_integration Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
2022-06-15 09:43:49 +08:00
x = ti.field(dtype=qu1)
y = ti.field(dtype=qu1)
[type] [Bug] Fix bit array vectorization on GPU (#2120)
2020-12-26 12:31:21 +08:00
N = 4096
n_blocks = 4
bits = 32
boundary_offset = 1024
block = ti.root.pointer(ti.ij, (n_blocks, n_blocks))
[Lang] [type] Refine SNode with quant 5/n: Rename bit_array to quant_array (#5344) * [Lang] [type] Rename bit_array to quant_array * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Fix Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
2022-07-06 16:45:40 +08:00
block.dense(ti.ij, (N // n_blocks, N // (bits * n_blocks))).quant_array(ti.j, bits, max_num_bits=bits).place(x)
block.dense(ti.ij, (N // n_blocks, N // (bits * n_blocks))).quant_array(ti.j, bits, max_num_bits=bits).place(y)
[type] [Bug] Fix bit array vectorization on GPU (#2120)
2020-12-26 12:31:21 +08:00
@ti.kernel
def init():
for i, j in ti.ndrange(
(boundary_offset, N - boundary_offset),
[async] [IR] More accurate same_value analysis (#2118) * [async] [IR] More precise same_value analysis * Remove debug outputs * Fix std::unordered_set::merge()... (with debug outputs) * Remove debug outputs * Remove do ... while (false); * [skip ci] minor change * [skip ci] minor change * [skip ci] enforce code format * Apply review * Apply review Co-authored-by: Taichi Gardener <taichigardener@gmail.com>
2020-12-27 15:28:16 +08:00
(boundary_offset, N - boundary_offset),
):
[type] [Bug] Fix bit array vectorization on GPU (#2120)
2020-12-26 12:31:21 +08:00
x[i, j] = ti.random(dtype=ti.i32) % 2
@ti.kernel
def assign_vectorized():
[Lang] [type] Make bit_vectorize a parameter of ti.loop_config (#5334)
2022-07-05 21:35:06 +08:00
ti.loop_config(bit_vectorize=True)
[type] [Bug] Fix bit array vectorization on GPU (#2120)
2020-12-26 12:31:21 +08:00
for i, j in x:
y[i, j] = x[i, j]
@ti.kernel
def verify():
for i, j in ti.ndrange(
(boundary_offset, N - boundary_offset),
[async] [IR] More accurate same_value analysis (#2118) * [async] [IR] More precise same_value analysis * Remove debug outputs * Fix std::unordered_set::merge()... (with debug outputs) * Remove debug outputs * Remove do ... while (false); * [skip ci] minor change * [skip ci] minor change * [skip ci] enforce code format * Apply review * Apply review Co-authored-by: Taichi Gardener <taichigardener@gmail.com>
2020-12-27 15:28:16 +08:00
(boundary_offset, N - boundary_offset),
):
[type] [Bug] Fix bit array vectorization on GPU (#2120)
2020-12-26 12:31:21 +08:00
assert y[i, j] == x[i, j]
init()
assign_vectorized()
verify()
[type] Support offset load in bit vectorized loop (#2127)
2021-01-01 06:17:14 +08:00
[type] [bug] Refine SNode with quant 1/n: Fix (atomic_)set_mask_b##N (#5238) * [type] [bug] Refine SNode with quant 1/n: Fix (atomic_)set_mask_b##N * Update metal
2022-06-24 08:55:49 +08:00
@test_utils.test(require=ti.extension.quant, debug=True)
[type] Support offset load in bit vectorized loop (#2127)
2021-01-01 06:17:14 +08:00
def test_offset_load():
[type] [bug] Refine SNode with quant 1/n: Fix (atomic_)set_mask_b##N (#5238) * [type] [bug] Refine SNode with quant 1/n: Fix (atomic_)set_mask_b##N * Update metal
2022-06-24 08:55:49 +08:00
qu1 = ti.types.quant.int(1, False)
[type] Support offset load in bit vectorized loop (#2127)
2021-01-01 06:17:14 +08:00
[type] [bug] Refine SNode with quant 1/n: Fix (atomic_)set_mask_b##N (#5238) * [type] [bug] Refine SNode with quant 1/n: Fix (atomic_)set_mask_b##N * Update metal
2022-06-24 08:55:49 +08:00
x = ti.field(dtype=qu1)
y = ti.field(dtype=qu1)
z = ti.field(dtype=qu1)
[type] Support offset load in bit vectorized loop (#2127)
2021-01-01 06:17:14 +08:00
N = 4096
n_blocks = 4
bits = 32
boundary_offset = 1024
assert boundary_offset >= N // n_blocks
block = ti.root.pointer(ti.ij, (n_blocks, n_blocks))
[Lang] [type] Refine SNode with quant 5/n: Rename bit_array to quant_array (#5344) * [Lang] [type] Rename bit_array to quant_array * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Fix Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
2022-07-06 16:45:40 +08:00
block.dense(ti.ij, (N // n_blocks, N // (bits * n_blocks))).quant_array(ti.j, bits, max_num_bits=bits).place(x)
block.dense(ti.ij, (N // n_blocks, N // (bits * n_blocks))).quant_array(ti.j, bits, max_num_bits=bits).place(y)
block.dense(ti.ij, (N // n_blocks, N // (bits * n_blocks))).quant_array(ti.j, bits, max_num_bits=bits).place(z)
[type] Support offset load in bit vectorized loop (#2127)
2021-01-01 06:17:14 +08:00
@ti.kernel
def init():
for i, j in ti.ndrange(
(boundary_offset, N - boundary_offset),
(boundary_offset, N - boundary_offset),
):
x[i, j] = ti.random(dtype=ti.i32) % 2
@ti.kernel
def assign_vectorized(dx: ti.template(), dy: ti.template()):
[Lang] [type] Make bit_vectorize a parameter of ti.loop_config (#5334)
2022-07-05 21:35:06 +08:00
ti.loop_config(bit_vectorize=True)
[type] Support offset load in bit vectorized loop (#2127)
2021-01-01 06:17:14 +08:00
for i, j in x:
y[i, j] = x[i + dx, j + dy]
z[i, j] = x[i + dx, j + dy]
@ti.kernel
def verify(dx: ti.template(), dy: ti.template()):
for i, j in ti.ndrange(
(boundary_offset, N - boundary_offset),
(boundary_offset, N - boundary_offset),
):
assert y[i, j] == x[i + dx, j + dy]
init()
assign_vectorized(0, 1)
verify(0, 1)
assign_vectorized(1, 0)
verify(1, 0)
assign_vectorized(0, -1)
verify(0, -1)
assign_vectorized(-1, 0)
verify(-1, 0)
[type] Local adder structure (#2136) * test all common 8 cases * transform atomic add and clean up * format * transform boolean expr and make test pass * refine test case * inline function
2021-01-06 22:37:15 +08:00
assign_vectorized(1, 1)
verify(1, 1)
assign_vectorized(1, -1)
verify(1, -1)
assign_vectorized(-1, -1)
verify(-1, -1)
assign_vectorized(-1, 1)
verify(-1, 1)
[lang] Added u1 as boolean type to taichi lang, replacing i32 ghstack-source-id: 242a65540f7c8ddd1501d2b2ba0a2f3d68629639 Pull Request resolved: https://github.com/taichi-dev/taichi/pull/8021
2023-05-17 15:07:09 +08:00
# FIXME:
# this test fails after we introduced type u1. Actually before we introduced u1 to taichi, this test has already
# appeared to be problematic. All problems are related to this code:
# `y[i, j] = (num_active_neighbors == 3) | ((num_active_neighbors == 2) & (x[i, j] == 1))`
# Before we introduce new type u1, problems arise when:
# 1. Replace | and & with `or` and `and`
# 2. Wrap this expression with `1 if ... else 0
# After we introduced new type u1, we can't pass this test with or without those modifications.
# Some experiments had been carried out on this problem. The results are as follows.
# +--------+-------------------------------+--------------------------------+---------------+
# | (y, z) | Replace `|``&` with `or``and` | Wrap expr with `1 if ... 0` | Do nothing |
# +--------+-------------------------------+--------------------------------+---------------+
# | Before | always (0, 1) | often (0, 1), sometimes (1, 0) | OK |
# | After | always (0, 1) | always(0, 1) | always (0, 1) |
# +--------+-------------------------------+--------------------------------+---------------+
# @test_utils.test(require=ti.extension.quant, debug=True)
# def test_evolve():
# qu1 = ti.types.quant.int(1, False)
#
# x = ti.field(dtype=qu1)
# y = ti.field(dtype=qu1)
# z = ti.field(dtype=qu1)
#
# N = 4096
# n_blocks = 4
# bits = 32
# boundary_offset = 1024
# assert boundary_offset >= N // n_blocks
#
# block = ti.root.pointer(ti.ij, (n_blocks, n_blocks))
# block.dense(ti.ij, (N // n_blocks, N // (bits * n_blocks))).quant_array(ti.j, bits, max_num_bits=bits).place(x)
# block.dense(ti.ij, (N // n_blocks, N // (bits * n_blocks))).quant_array(ti.j, bits, max_num_bits=bits).place(y)
# block.dense(ti.ij, (N // n_blocks, N // (bits * n_blocks))).quant_array(ti.j, bits, max_num_bits=bits).place(z)
#
# @ti.kernel
# def init():
# for i, j in ti.ndrange(
# (boundary_offset, N - boundary_offset),
# (boundary_offset, N - boundary_offset),
# ):
# x[i, j] = ti.random(dtype=ti.i32) % 2
#
# @ti.kernel
# def evolve_vectorized(x: ti.template(), y: ti.template()):
# ti.loop_config(bit_vectorize=True)
# for i, j in x:
# num_active_neighbors = 0
# num_active_neighbors += ti.cast(x[i - 1, j - 1], ti.u32)
# num_active_neighbors += ti.cast(x[i - 1, j], ti.u32)
# num_active_neighbors += ti.cast(x[i - 1, j + 1], ti.u32)
# num_active_neighbors += ti.cast(x[i, j - 1], ti.u32)
# num_active_neighbors += ti.cast(x[i, j + 1], ti.u32)
# num_active_neighbors += ti.cast(x[i + 1, j - 1], ti.u32)
# num_active_neighbors += ti.cast(x[i + 1, j], ti.u32)
# num_active_neighbors += ti.cast(x[i + 1, j + 1], ti.u32)
# y[i, j] = (num_active_neighbors == 3) | ((num_active_neighbors == 2) & (x[i, j] == 1))
#
# @ti.kernel
# def evolve_naive(x: ti.template(), y: ti.template()):
# for i, j in ti.ndrange(
# (boundary_offset, N - boundary_offset),
# (boundary_offset, N - boundary_offset),
# ):
# num_active_neighbors = 0
# num_active_neighbors += ti.cast(x[i - 1, j - 1], ti.u32)
# num_active_neighbors += ti.cast(x[i - 1, j], ti.u32)
# num_active_neighbors += ti.cast(x[i - 1, j + 1], ti.u32)
# num_active_neighbors += ti.cast(x[i, j - 1], ti.u32)
# num_active_neighbors += ti.cast(x[i, j + 1], ti.u32)
# num_active_neighbors += ti.cast(x[i + 1, j - 1], ti.u32)
# num_active_neighbors += ti.cast(x[i + 1, j], ti.u32)
# num_active_neighbors += ti.cast(x[i + 1, j + 1], ti.u32)
# y[i, j] = (num_active_neighbors == 3) or (num_active_neighbors == 2 and x[i, j] == 1)
#
# @ti.kernel
# def verify():
# for i, j in ti.ndrange(
# (boundary_offset, N - boundary_offset),
# (boundary_offset, N - boundary_offset),
# ):
# assert y[i, j] == z[i, j]
#
# init()
# evolve_naive(x, z)
# evolve_vectorized(x, y)
# verify()
Reference in New Issue Copy Permalink