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

Prefer Array class over register arrays. #3737

Open
wants to merge 7 commits into
base: main
Choose a base branch
from
Open

Prefer Array class over register arrays. #3737

wants to merge 7 commits into from

Conversation

csarofeen
Copy link
Collaborator

Prefer Array class over register arrays.
For example in code generation use Array<float, 2> not float[2].

@csarofeen
Copy link
Collaborator Author

!test

Copy link

github-actions bot commented Jan 20, 2025

PR Reviewer Guide 🔍

(Review updated until commit 1e7014a)

Here are some key observations to aid the review process:

⏱️ Estimated effort to review: 3 🔵🔵🔵⚪⚪
🧪 No relevant tests
⚡ Recommended focus areas for review

Logic Change

The aligned_array_of_regs_ set has been removed, and the genVariableNameConvertAlignedArray function has been replaced with genVariableName. This change may affect the handling of aligned arrays in the code generation process.

void initStringStreamFormat(std::stringstream& ss) {
Type Change

The philox function now returns an Array<uint32_t, 4> instead of a uint4. This change may affect the usage of the philox function in other parts of the code.

__device__ Array<uint32_t, 4> philox(

@csarofeen
Copy link
Collaborator Author

!test --pybench

@csarofeen
Copy link
Collaborator Author

csarofeen commented Jan 21, 2025

All tests passing.
Benchmarks do not appear to regress. https://nv/ei9/22766159

@naoyam
Copy link
Collaborator

naoyam commented Jan 21, 2025

Prefer Array class over register arrays. For example in code generation use Array<float, 2> not float[2].

Any specific reason?

@zasdfgbnm
Copy link
Collaborator

Prefer Array class over register arrays. For example in code generation use Array<float, 2> not float[2].

Any specific reason?

Likely because currently we are using float[2] in some cases, but Array<float, 2> in other cases, and sometimes there are compatibility issue for these two types? I remembered Liqiang fixed one such issue on vectorization.

@zasdfgbnm
Copy link
Collaborator

zasdfgbnm commented Jan 21, 2025

Possibly related: #1475

csrc/codegen.cpp Outdated
@@ -265,7 +265,9 @@ class CudaKernelGenerator : private kir::ConstIrVisitor {
} else if (v->isA<TensorView>()) {
tv = v->as<TensorView>();
}
if (tv && aligned_array_of_regs_.count(tv)) {
if (tv &&
(aligned_array_of_regs_.count(tv) ||
Copy link
Collaborator

Choose a reason for hiding this comment

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

Seems that aligned_array_of_regs_ can be removed now.

Copy link
Collaborator Author

Choose a reason for hiding this comment

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

Makes sense, will remove.

@csarofeen
Copy link
Collaborator Author

@naoyam I thought it would help to make things consistent. I also wanted to convert RNG to all use the class instead of literal arrays. This would help do something like generate bit array with Philox and reinterpret it as more lower precision integer types.

Copy link
Collaborator

@naoyam naoyam left a comment

Choose a reason for hiding this comment

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

LGTM

@csarofeen
Copy link
Collaborator Author

!test

@csarofeen
Copy link
Collaborator Author

!test

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