[GH-ISSUE #27] why your guess? "st.global.L1::no_allocate" that doesn't mean L2 won't be allocated cache. #27

Closed
opened 2026-10-26 09:52:43 +00:00 by bsakaguchi · 2 comments
Owner

Originally created by @zzJingHong on GitHub (Feb 26, 2025).
Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/27

// st.global.L1::no_allocate will be translated into ST.E.NA.[width] in SASS,
// which does not have cache allocation (obviously in L1, I guess not in L2 too)
#ifndef DISABLE_AGGRESSIVE_PTX_INSTRS
#define ST_NA_FUNC "st.global.L1::no_allocate"
#else
#define ST_NA_FUNC "st.global"
#endif

Originally created by @zzJingHong on GitHub (Feb 26, 2025). Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/27 // `st.global.L1::no_allocate` will be translated into `ST.E.NA.[width]` in SASS, // which does not have cache allocation (obviously in L1, I guess not in L2 too) #ifndef DISABLE_AGGRESSIVE_PTX_INSTRS #define ST_NA_FUNC "st.global.L1::no_allocate" #else #define ST_NA_FUNC "st.global" #endif
Author
Owner

@LyricZhao commented on GitHub (Feb 27, 2025):

Sorry, the comment is somehow misleading: here, I'm guessing the "NA" in ST.E.NA not the PTX. Will update related docs later.

<!-- gh-comment-id:2686647898 --> @LyricZhao commented on GitHub (Feb 27, 2025): Sorry, the comment is somehow misleading: here, I'm guessing the "NA" in `ST.E.NA` not the PTX. Will update related docs later.
Author
Owner

@zsnoob commented on GitHub (Mar 29, 2025):

Yeah, though a bit misleading here, it's actually great to see these kinds of speculations and discussions. Even if unverified, they're valuable—after all, the PTX's undefined behavior is exactly the kind of thing we are speculating about in the first place.

<!-- gh-comment-id:2764289274 --> @zsnoob commented on GitHub (Mar 29, 2025): Yeah, though a bit misleading here, it's actually great to see these kinds of speculations and discussions. Even if unverified, they're valuable—after all, the PTX's undefined behavior is exactly the kind of thing we are speculating about in the first place.
Sign in to join this conversation.
No labels
pull-request
No milestone
No project
No assignees
1 participant
Notifications
Due date
The due date is invalid or out of range. Please use the format "yyyy-mm-dd".

No due date set.

Dependencies

No dependencies set.

Reference
deepseek-ai/DeepEP#27
No description provided.