Skip to content

Add Loads from fixed inputs#162031

Closed
drisspg wants to merge 22 commits intogh/drisspg/197/basefrom
gh/drisspg/197/head
Closed

Add Loads from fixed inputs#162031
drisspg wants to merge 22 commits intogh/drisspg/197/basefrom
gh/drisspg/197/head

Conversation

@drisspg
Copy link
Contributor

@drisspg drisspg commented Sep 3, 2025 •
edited by pytorch-bot bot
Loading

Stack from ghstack (oldest at bottom):

TODO

Check on multi indices

    @cute.jit
    def score_mod(tSrS_ssa, b_idx, h_idx, q_idx, kv_idx, buffers):
        in_ptr4 = buffers[0]
        tmp0 = tSrS_ssa
        tmp1 = b_idx
        tmp2 = h_idx
        tmp3 = cute.make_fragment(1, cutlass.Int32)
        tmp4 = tmp3.store(32*tmp1 + tmp2)
        tmp5 = cute.make_fragment(1, cutlass.BFloat16)
        tmp6 = tmp3[0]
        tmp7 = tmp5[0] = (in_ptr4[tmp6])
        tmp8 = (tmp5.load()).to(cutlass.Float32)
        tmp9 = (tmp0 + tmp8)
        tSrS_ssa = tmp9

        return tSrS_ssa

I dont think that

        tmp4 = tmp3.store(32*tmp1 + tmp2)
        tmp5 = cute.make_fragment(1, cutlass.BFloat16)
        tmp6 = tmp3[0]
        tmp7 = tmp5[0] = (in_ptr4[tmp6]

is right since this tmp6 value will be larger than the actual index dim int his case its B -> see if its possible to 1d index

cc @voznesenskym @penguinwu @EikanWang @jgong5 @Guobing-Chen @XiaobingSuper @zhuhaozhe @blzheng @wenzhe-nrv @jiayisunx @ipiszy @chenyang78 @kadeng @muchulee8 @amjames @chauhang @aakhundov @coconutruben

@drisspg
Update
eba6001 
[ghstack-poisoned]
@pytorch-bot
Copy link

pytorch-bot bot commented Sep 3, 2025 •
edited
Loading

🔗 Helpful Links

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

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

⏳ 1 Pending, 1 Unrelated Failure

As of commit 954c5f1 with merge base 086dec3 (image):

BROKEN TRUNK - The following job failed but were present on the merge base:

👉 Rebase onto the `viable/strict` branch to avoid these failures

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

drisspg added a commit that referenced this pull request Sep 3, 2025
@drisspg
Add Loads from fixe inputs
32a76b4 
ghstack-source-id: 03e809f
Pull-Request: #162031
@drisspg drisspg mentioned this pull request Sep 3, 2025
Closed
@drisspg drisspg marked this pull request as draft September 3, 2025 22:24
@drisspg
Update
481e4c6 
[ghstack-poisoned]
drisspg added a commit that referenced this pull request Sep 3, 2025
@drisspg
Add Loads from fixe inputs
8bd82aa 
ghstack-source-id: 4ae4670
Pull-Request: #162031
@drisspg
Update
823ebf2 
[ghstack-poisoned]
drisspg added a commit that referenced this pull request Sep 4, 2025
@drisspg
Add Loads from fixed inputs
b0a8261 
ghstack-source-id: 5f81aeb
Pull-Request: #162031
@drisspg
Update
81c1a03 
[ghstack-poisoned]
drisspg added a commit that referenced this pull request Sep 4, 2025
@drisspg
Add Loads from fixed inputs
0843c00 
ghstack-source-id: 2b3f5c7
Pull-Request: #162031
@drisspg drisspg changed the title Add Loads from fixe inputs Add Loads from fixed inputs Sep 4, 2025
@drisspg drisspg added the topic: not user facing topic category label Sep 4, 2025
@drisspg drisspg mentioned this pull request Sep 4, 2025
Merged
@drisspg
Update
1705f35 
[ghstack-poisoned]
drisspg added a commit that referenced this pull request Sep 5, 2025
@drisspg
Add Loads from fixed inputs
77730f2 
ghstack-source-id: 3a4328f
Pull-Request: #162031
@drisspg
Update
512960d 
[ghstack-poisoned]
drisspg added a commit that referenced this pull request Sep 10, 2025
@drisspg
Add Loads from fixed inputs
2157284 
ghstack-source-id: b6aba51
Pull-Request: #162031
@drisspg
Update
fc7de50 
[ghstack-poisoned]
drisspg added a commit that referenced this pull request Sep 11, 2025
@drisspg
Add Loads from fixed inputs
6a90c62 
ghstack-source-id: 8709217
Pull-Request: #162031
@drisspg
Update
3e00d34 
[ghstack-poisoned]
drisspg added a commit that referenced this pull request Sep 11, 2025
@drisspg
Add Loads from fixed inputs
2dcede6 
ghstack-source-id: 2dfedca
Pull-Request: #162031
@drisspg
Update
01bd655 
[ghstack-poisoned]
drisspg added a commit that referenced this pull request Sep 11, 2025
@drisspg
Add Loads from fixed inputs
f17e4c3 
ghstack-source-id: b07003d
Pull-Request: #162031
@drisspg
Update
0a2073e 
[ghstack-poisoned]
drisspg added a commit that referenced this pull request Sep 17, 2025
@drisspg
Add Loads from fixed inputs
98a9f2d 
ghstack-source-id: d11e9cb
Pull-Request: #162031
@drisspg
Update
691f8b0 
[ghstack-poisoned]
drisspg added a commit that referenced this pull request Sep 18, 2025
@drisspg
Add Loads from fixed inputs
d47b422 
ghstack-source-id: 5471f7e
Pull-Request: #162031
@drisspg
Update
7da7472 
[ghstack-poisoned]
drisspg added a commit that referenced this pull request Oct 5, 2025
@drisspg
Add Loads from fixed inputs
dc6484a 
ghstack-source-id: 66877d2
Pull-Request: #162031
drisspg added a commit that referenced this pull request Oct 5, 2025
@drisspg
Add Loads from fixed inputs
8108523 
ghstack-source-id: 66877d2
Pull-Request: #162031
@drisspg drisspg marked this pull request as ready for review October 7, 2025 03:28
@drisspg
Update
37c6db3 
[ghstack-poisoned]
drisspg added a commit that referenced this pull request Oct 8, 2025
@drisspg
Add Loads from fixed inputs
f3933b9 
ghstack-source-id: d8f172c
Pull-Request: #162031
@drisspg
Update
bca3ad4 
[ghstack-poisoned]
drisspg added a commit that referenced this pull request Oct 8, 2025
@drisspg
Add Loads from fixed inputs
ca85a49 
ghstack-source-id: 5d43690
Pull-Request: #162031
@drisspg
Update
2ba05c2 
[ghstack-poisoned]
drisspg added a commit that referenced this pull request Oct 9, 2025
@drisspg
Add Loads from fixed inputs
5d0bac0 
ghstack-source-id: 3da0b78
Pull-Request: #162031
@drisspg
Copy link
Contributor Author

drisspg commented Oct 9, 2025

@pytorchbot merge

@pytorch-bot pytorch-bot bot added the ciflow/trunk Trigger trunk jobs on your pull request label Oct 9, 2025
@pytorchmergebot
Copy link
Collaborator

pytorchmergebot commented Oct 9, 2025

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

@pytorchmergebot
Copy link
Collaborator

pytorchmergebot commented Oct 9, 2025

Merge failed

Reason: 1 jobs have failed, first few of them are: trunk / linux-jammy-cuda12.8-py3.10-gcc11 / test (default, 1, 5, lf.linux.g6.4xlarge.experimental.nvidia.gpu)

Details for Dev Infra team Raised by workflow job

@drisspg
Update
b254144 
[ghstack-poisoned]
drisspg added a commit that referenced this pull request Oct 9, 2025
@drisspg
Add Loads from fixed inputs
df22f4c 
ghstack-source-id: 5cb3d5c
Pull-Request: #162031
@drisspg
Update
954c5f1 
[ghstack-poisoned]
drisspg added a commit that referenced this pull request Oct 9, 2025
@drisspg
Add Loads from fixed inputs
d2752fa 
ghstack-source-id: dd3cd80
Pull-Request: #162031
@drisspg
Copy link
Contributor Author

drisspg commented Oct 10, 2025

@pytorchbot merge

@pytorchmergebot
Copy link
Collaborator

pytorchmergebot commented Oct 10, 2025

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

Chao1Han pushed a commit to Chao1Han/pytorch that referenced this pull request Oct 21, 2025
@drisspg @Chao1Han
Add Loads from fixed inputs (pytorch#162031)
3792fed 
## TODO
Check on multi indices
```Python

    @cute.jit
    def score_mod(tSrS_ssa, b_idx, h_idx, q_idx, kv_idx, buffers):
        in_ptr4 = buffers[0]
        tmp0 = tSrS_ssa
        tmp1 = b_idx
        tmp2 = h_idx
        tmp3 = cute.make_fragment(1, cutlass.Int32)
        tmp4 = tmp3.store(32*tmp1 + tmp2)
        tmp5 = cute.make_fragment(1, cutlass.BFloat16)
        tmp6 = tmp3[0]
        tmp7 = tmp5[0] = (in_ptr4[tmp6])
        tmp8 = (tmp5.load()).to(cutlass.Float32)
        tmp9 = (tmp0 + tmp8)
        tSrS_ssa = tmp9

        return tSrS_ssa

 ```

I dont think that
```
        tmp4 = tmp3.store(32*tmp1 + tmp2)
        tmp5 = cute.make_fragment(1, cutlass.BFloat16)
        tmp6 = tmp3[0]
        tmp7 = tmp5[0] = (in_ptr4[tmp6]

```

 is right since this tmp6 value will be larger than the actual index dim int his case its B -> see if its possible to 1d index

Pull Request resolved: pytorch#162031
Approved by: https://github.com/v0i0
ghstack dependencies: pytorch#161118
@github-actions github-actions bot deleted the gh/drisspg/197/head branch November 9, 2025 02:19
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

@v0i0 v0i0 v0i0 approved these changes

Assignees

No one assigned

Labels

ciflow/inductor ciflow/trunk Trigger trunk jobs on your pull request Merged module: inductor topic: not user facing topic category

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

3 participants

@drisspg @pytorchmergebot @v0i0