[Lang] Simplify dense.bitmasked to bitmasked

[k-ye]

Related issue = #663

Also fixed a bug where the SNode type is not set correctly on Metal:

taichi/taichi/codegen/codegen_metal.cpp

Line 719 in 1371a07

la.append("{}.type = {};", var_name, meta.num_slots);

[Click here for the format server]

[k-ye]

I have a further question regarding this test

taichi/tests/python/test_sparse_basics.py

Lines 4 to 27 in 1371a07

	@ti.archs_support_sparse
	def test_bitmasked():
	x = ti.var(ti.f32)
	s = ti.var(ti.i32)
	n = 128
	@ti.layout
	def place():
	ti.root.dense(ti.i, n).bitmasked().dense(ti.i, n).place(x)
	ti.root.place(s)
	@ti.kernel
	def func():
	for i in x:
	s[None] += 1
	x[0] = 1
	x[127] = 1
	x[256] = 1
	x[257] = 1
	func()
	assert s[None] == 256

This test checks the number of iterations over the leaf SNode x. If we modify it to ti.root.bitmasked(ti.i, n).bitmasked(ti.i, n).place(x), then on CPU, it is 256. On Metal, however, it is 4 (only those being activated are iterated).

Should this be made as part of Taichi's semantics eventually? Or should we just treat this as an implementation detail? Personally I prefer the former, as it won't surprise the users. That said, this might restrict each backend's implementation.

[archibate]

Maybe two dense(128) caused this? The actual size of x is 128x128, however the second dim is not bitmasked, so the result is 128x2 (0 and 127 is one group, 256 and 257 is another group)

[k-ye]

Maybe two dense(128) caused this? The actual size of x is 128x128, however the second dim is not bitmasked, so the result is 128x2 (0 and 127 is one group, 256 and 257 is another group)

Yep, the behavior of bitmasked.dense is the same. But bitmasked.bitmasked are different.

[yuanming-hu]

Your Metal implementation is correct. On CPUs, bitmasked that directly includes place is not correctly implemented. I assumed the parent of place must be dense. I'll fix this: #671

[yuanming-hu]

LGTM!

[yuanming-hu]

Should this be made as part of Taichi's semantics eventually?

Yes! The ultimate goal of Taichi sparse programming is to make algorithms fully independent of data structures. Assuming the finest level must be a dense node actually harms, since the data mask now has a granularity of the dense node dim (e.g. 4x4x8 if the dense node has this dim, which is data-structure-dependent) instead of 1x1x1.

1x1x1 sparsity mask is the only option that is independent of the data structure.