[Metal] Changes to enable bitmasked on Metal!

[k-ye]

Related issue id = #593

[Click here for the format server]

Feel like the tests are pretty weak...

[archibate]

Interesting! Is there any pre-requirements for OpenGL to also support this? Not familiar to what bitmasked mean yet...

[archibate]

So, the difference bitmasked introduced is, only those assigned slots are marked as is_active?

[archibate]

If I'm thinking right, this can be used for performance optimization, like Game of Life - only the non-static cells are calculated?

[k-ye]

Exactly. For sparse data structures, struct_for will only loop over the active indices (it could still iterate over more slots due to POT paddings, but not all of them).

[yuanming-hu]

If I'm thinking right, this can be used for performance optimization, like Game of Life - only the non-static cells are calculated?

Great idea! If you are interested, having a game of life demo with sparse grids would be super cool!!

[k-ye]

Not familiar to what bitmasked mean yet...

It's the simplest sparse data structure in taichi. Under the hood, it is dense, with each element annotated by 1 bit to indicate its activeness.

taichi/taichi/runtime/llvm/node_dense.h

Lines 19 to 25 in 068af03

	if (DenseMeta_get_bitmasked(dmeta)) {
	auto element_size = StructMeta_get_element_size(smeta);
	auto num_elements = Dense_get_num_elements(meta, node);
	auto data_section_size = element_size * num_elements;
	auto mask_begin = (uint64 *)(node + data_section_size);
	atomic_or_u64(&mask_begin[i / 64], 1UL << (i % 64));
	}

Is there any pre-requirements for OpenGL to also support this?

For each bitmasked (or any other sparse) SNode, it is associated with a ListManager stored on the device side. The active indices, along with some other properties, are appended to this list. So if you preallocate the list memory, and use an atomic to simulate the append behavior, then i guess OpenGL can also support it.

• Metal

taichi/taichi/platform/metal/shaders/runtime_utils.metal.h

Lines 39 to 53 in 068af03

	int append(device ListManager *list,
	thread const T &elem,
	device byte *data_addr) {
	thread char *elem_ptr = (thread char *)(&elem);
	int me = atomic_fetch_add_explicit(
	reinterpret_cast<device atomic_int *>(&(list->next)), 1,
	metal::memory_order_relaxed);
	device byte *ptr =
	data_addr + list->mem_begin + (me * list->element_stride);
	for (int i = 0; i < list->element_stride; ++i) {
	*ptr = *elem_ptr;
	++ptr;
	++elem_ptr;
	}
	return me;

• LLVM

taichi/taichi/runtime/llvm/runtime.cpp

Lines 1105 to 1108 in 068af03

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

If you want to support sparse SNodes on OpenGL, i think https://github.com/taichi-dev/taichi/blob/master/taichi/runtime/llvm/runtime.cpp is a really good start.

[yuanming-hu]

Looks good to me! Thanks so much for contributing this!!

Looking at the recurring pattern snode->is_dense && snode->is_bitmasked, I'm wondering if we should promote bitmasked to be a standalone SNodeType, instead of a decorator to dense. We used to have dense.pointer, yet later pointers are made a standalone data structure (thanks to @KLozes). Maybe we should do the same for bitmasked since

• bitmasked is always used with dense, so maybe we should save the verbosity;
• dense.bitmasked is too different from dense so bitmasked worths a standalone SNodeType.

[k-ye]

Maybe we should do the same for bitmasked since

SGTM! Shall I make the change before merging this?

[yuanming-hu]

That doesn't sound like a small change to me. Let's merge this in first :-)