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] bitmasked support #737

Closed
wants to merge 10 commits into from
Closed

Conversation

archibate
Copy link
Collaborator

@archibate archibate commented Apr 9, 2020

Related issue = #711

First problem is bitmasked cannot be assigned correctly:

import taichi as ti

ti.init(arch=ti.opengl)
ti.misc.util.set_gdb_trigger(False)

x = ti.var(ti.i32)
y = ti.var(ti.i32, shape=())
z = ti.var(ti.i32, shape=())
ti.root.bitmasked(ti.i, 4).place(x)

@ti.kernel
def func():
    y[None] = ti.is_active(x.parent(), 0)
    x[0] = 233
    z[None] = ti.is_active(x.parent(), 0)

func()
print(x[0])
print(y[None])
print(z[None])

It shows 0 0 0, it should be at least 233 0 0.
Good night!

[Click here for the format server]

@archibate archibate requested review from yuanming-hu and k-ye April 10, 2020 06:15
@archibate archibate changed the title [OpenGL] Sparse data structure support [OpenGL] bitmasked support Apr 10, 2020
@archibate archibate marked this pull request as ready for review April 10, 2020 06:29
@archibate
Copy link
Collaborator Author

Metal root buffer layout:

SNodeA: i32 | i32 | ...(12)... | i32 | i32 | 00...(12)...00
SNodeB: i32 | i32 | ...(233)... | i32 | i32 | 00...(233)...00

OpenGL root buffer layout (for now):

Unified bitmask array:  00...(12)...00 | 00...(233)...00
SNodeA: i32 | i32 | ...(12)... | i32 | i32
SNodeB: i32 | i32 | ...(233)... | i32 | i32

@archibate archibate self-assigned this Apr 10, 2020
TI_ASSERT(stmt->width() == 1);
auto dt = stmt->data->element_type();
auto ch_addr = stmt->ptr->short_name();
emit("atomicOr(_rt_.bitmask[{} >> 8], 1 << ({} & 31));", ch_addr, ch_addr);
Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is this the correct usage to activate when store?

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actually, activation should happen on SNodeLookUp statements. When doing GlobalStore you should assume the element is already activated.

@archibate
Copy link
Collaborator Author

How to struct_for a bitmasked snode? if (!is_active) return;? Or something better?

@yuanming-hu
Copy link
Member

How to struct_for a bitmasked snode? if (!is_active) return;? Or something better?

That depends - for intermediate SNode we need a list generation (see x64/CUDA/Metal backends as examples), for leaf SNode if (!is_active) return; is fine.

@archibate
Copy link
Collaborator Author

What is a list generation?
Why it's fine for leaf SNode like dense().bitmasked()?
A non-trival example might be bitmasked().bitmasked() or bitmasked().dense()?

@yuanming-hu
Copy link
Member

Good questions - I should probably do my job to improve the documentation #736 .... I plan to do this on Sunday. Sorry for blocking you. If you have extra cycles maybe we can have #653 in.

Note that it's fine if OpenGL backend in v0.6.0 doesn't support sparse data structures. Metal doesn't support sparse either on its initial release. Most demos so far doesn't require sparsity.

We should also include glfw and glew as submodules, and statically link their binary into libtaichi_sore.so so that we can really ship the OpenGL backend.

@archibate
Copy link
Collaborator Author

Right, we'll first ship it, then improve it!
Btw I found #653 too hard thanks to the 30m signal delay from the CI virtual machines...

@yuanming-hu
Copy link
Member

Right, we'll first ship it, then improve it!
Btw I found #653 too hard thanks to the 30m signal delay from the CI virtual machines...

Understood... Setting up CI can be boring due to the long build delay - feel free to do that while watching Bilibili...

@archibate archibate marked this pull request as draft April 16, 2020 16:32
@archibate
Copy link
Collaborator Author

Won't proceed before doc complete, closing for now.

@archibate archibate closed this Apr 22, 2020
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

3 participants