-
Notifications
You must be signed in to change notification settings - Fork 2.3k
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
[doc] Document list generation #990
Conversation
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Thank for adding this, this is really helpful for backend maintainers, lgtm in general except for a few questions.
docs/internal.rst
Outdated
- (Tasks ``$0`` and ``$1``) from the list of ``root`` node (``S0``) to list of the ``dense`` nodes (``S1``); | ||
- (Tasks ``$2`` and ``$3``) from the list of ``dense`` nodes (``S1``) to the list of ``bitmasked`` nodes (``S2``). |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
What does from ... to
mean? Can we clarify this?
@ti.kernel | ||
def func(): | ||
for i in x: | ||
print(i) |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I guess this will always get not printed since bitmasked is inactivated by default.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
That's true. We are just illustrating list generation here.
docs/internal.rst
Outdated
Our strategy is to generate lists of SNode elements layer by layer. | ||
This flattens the data structure leaf elements into a 1D dense array, circumventing the irregularity of | ||
incomplete trees. Then we can simply invoke a regular **parallel for** over the list. |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Our strategy is to generate lists of SNode elements layer by layer. | |
This flattens the data structure leaf elements into a 1D dense array, circumventing the irregularity of | |
incomplete trees. Then we can simply invoke a regular **parallel for** over the list. | |
Our strategy is to generate a list of SNode element indices for each layer. | |
This flattens the data structure leaf elements into a 1D dense array, circumventing the irregularity of | |
incomplete trees. Then we can simply invoke a regular **parallel for** over the list. |
Also, is the generated list stored in CPU or GPU memory? If you find the process hard to explain, please use LLVM implementation for example.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
It's on GPU, and these lists are updated in a kernel. There are multiple lists so I think plural form is the correct one. I'd suggest to say lists of active SNode elements layer by layer
.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Our strategy is to generate lists of SNode elements layer by layer. | |
This flattens the data structure leaf elements into a 1D dense array, circumventing the irregularity of | |
incomplete trees. Then we can simply invoke a regular **parallel for** over the list. | |
Our strategy is to generate several lists of active SNode elements on GPU, layer by layer. | |
This flattens the data structure leaf elements into a 1D dense array, circumventing the irregularity of | |
incomplete trees. Then we can simply invoke a regular **parallel for** over the list. |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I'll be more specific and say generation will happen on the device.
docs/internal.rst
Outdated
Our strategy is to generate lists of SNode elements layer by layer. | ||
This flattens the data structure leaf elements into a 1D dense array, circumventing the irregularity of | ||
incomplete trees. Then we can simply invoke a regular **parallel for** over the list. |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
It's on GPU, and these lists are updated in a kernel. There are multiple lists so I think plural form is the correct one. I'd suggest to say lists of active SNode elements layer by layer
.
Thanks for the reviews! I'm merging this, so that I can add more details on the internal designs (for #986) following this PR. |
Related issue = #709 #941 #986
[Click here for the format server]