Skip to content

[CUDA][64-bit indexing] Fix some existing problematic int64_t _ = blockIdx.* * blockDim.* code #142010

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.

Sign up for GitHub

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

Jump to bottom
Closed
eqy wants to merge 1 commit into from
Closed

[CUDA][64-bit indexing] Fix some existing problematic int64_t _ = blockIdx.* * blockDim.* code #142010

eqy wants to merge 1 commit into from

Conversation

@eqy
Copy link
Collaborator

@eqy eqy commented Dec 3, 2024 •
edited by pytorch-bot bot
Loading

grep didn't surface any blockIdx.z * blockDim.z cases

git grep -l "int64_t.*=.*blockIdx.x \* blockDim.x.*" | xargs sed -i 's/int64_t \(.*\) = blockIdx.x \* blockDim.x + threadIdx.x;.*/int64_t \1 = ((int64_t) blockIdx.x) * blockDim.x + threadIdx.x;/g'
git grep -l "int64_t.*=.*blockIdx.x \* blockDim.x.*" | xargs sed -i 's/int64_t \(.*\) = threadIdx.x + blockIdx.x \* blockDim.x;.*/int64_t \1 = threadIdx.x + ((int64_t) blockIdx.x) * blockDim.x;/g'
git grep -l "int64_t.*=.*blockIdx.y \* blockDim.y.*" | xargs sed -i 's/int64_t \(.*\) = blockIdx.y \* blockDim.y + threadIdx.y;.*/int64_t \1 = ((int64_t) blockIdx.y) * blockDim.y + threadIdx.y;/g'
git grep -l "int64_t.*=.*blockIdx.y \* blockDim.y.*" | xargs sed -i 's/int64_t \(.*\) = threadIdx.y + blockIdx.y \* blockDim.y;.*/int64_t \1 = threadIdx.y + ((int64_t) blockIdx.y) * blockDim.y;/g'
git grep -l "int64_t.*=.*blockDim.x \* blockIdx.x.*" | xargs sed -i 's/int64_t \(.*\) = blockDim.x \* blockIdx.x + threadIdx.x;.*/int64_t \1 = ((int64_t) blockIdx.x) * blockDim.x + threadIdx.x;/g'

See also https://github.com/pytorch/pytorch/pull/141922/files#r1868262823 in #141999 141922

cc @ptrblck @msaroufim

@eqy eqy added module: cuda Related to torch.cuda, and CUDA support in general module: 64-bit Problems related to incorrectly using 32-bit integers when 64-bit is needed (e.g., 8G tensors) open source better-engineering Relatively self-contained tasks for better engineering contributors topic: not user facing topic category labels Dec 3, 2024
@eqy eqy requested a review from syed-ahmed as a code owner December 3, 2024 22:35
@pytorch-bot
Copy link

pytorch-bot bot commented Dec 3, 2024 •
edited
Loading

🔗 Helpful Links

🧪 See artifacts and rendered test results at hud.pytorch.org/pr/142010

Note: Links to docs will display an error until the docs builds have been completed.

✅ No Failures

As of commit 6a9f155 with merge base c06b504 (image):
💚 Looks good so far! There are no failures yet. 💚

This comment was automatically generated by Dr. CI and updates every 15 minutes.

@cpuhrsch cpuhrsch requested a review from ngimel December 4, 2024 00:01
@cpuhrsch cpuhrsch added the triaged This issue has been looked at a team member, and triaged and prioritized into an appropriate module label Dec 4, 2024
@eqy eqy added ciflow/trunk Trigger trunk jobs on your pull request ciflow/rocm Trigger "default" config CI on ROCm labels Dec 18, 2024
@eqy
Copy link
Collaborator Author

eqy commented Dec 18, 2024

@pytorchmergebot merge

@pytorchmergebot
Copy link
Collaborator

pytorchmergebot commented Dec 18, 2024

Merge started

Your change will be merged once all checks pass (ETA 0-4 Hours).

Learn more about merging in the wiki.

Questions? Feedback? Please reach out to the PyTorch DevX Team

Advanced Debugging
Check the merge workflow status
here

@eqy eqy mentioned this pull request Oct 9, 2025
Closed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

@ngimel ngimel ngimel approved these changes

@syed-ahmed syed-ahmed Awaiting requested review from syed-ahmed syed-ahmed is a code owner

Assignees

No one assigned

Labels

better-engineering Relatively self-contained tasks for better engineering contributors ciflow/rocm Trigger "default" config CI on ROCm ciflow/trunk Trigger trunk jobs on your pull request Merged module: cuda Related to torch.cuda, and CUDA support in general module: 64-bit Problems related to incorrectly using 32-bit integers when 64-bit is needed (e.g., 8G tensors) open source topic: not user facing topic category triaged This issue has been looked at a team member, and triaged and prioritized into an appropriate module

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

4 participants

@eqy @pytorchmergebot @ngimel @cpuhrsch