deepseek-ai/DeepEP

Fork 0

mirror of https://github.com/deepseek-ai/DeepEP.git synced 2026-11-05 00:47:32 +00:00

[GH-ISSUE #21] Question about GPU Driver #21

New issue

Closed

opened 2026-10-26 09:52:37 +00:00 by bsakaguchi · 5 comments

bsakaguchi commented

2026-10-26 09:52:37 +00:00

Owner

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

I would like to ask which version of the GPU driver you have installed on the H800 hopper architecture to support the smooth operation of PTX?

Originally created by @Qizhi697 on GitHub (Feb 26, 2025). Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/21 I would like to ask which version of the `GPU driver` you have installed on the `H800 hopper architecture` to support the smooth operation of `PTX`?

bsakaguchi closed this issue

2026-10-26 09:52:37 +00:00

bsakaguchi commented

2026-10-26 09:52:38 +00:00

Author

Owner

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

I guess the PTX is not related to OS driver, but an undoced Hopper hardware feature: L1::no_allocate is not a hint but a strong option to disable the L1 cache (so that ensuring the volatile data read correctness).

@LyricZhao commented on GitHub (Feb 26, 2025): I guess the PTX is not related to OS driver, but an undoced Hopper hardware feature: `L1::no_allocate` is not a hint but a strong option to disable the L1 cache (so that ensuring the volatile data read correctness).

bsakaguchi commented

2026-10-26 09:52:38 +00:00

Author

Owner

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

I guess the PTX is not related to OS driver, but an undoced Hopper hardware feature: L1::no_allocate is not a hint but a strong option to disable the L1 cache (so that ensuring the volatile data read correctness).

Brief discussion:

In this context, considering both the repository and the documentation, could we interpret this "disable" as follows: for volatile data, the use of no_allocate makes it impossible for the data to be written into L1 (unified with nc to make sure correctness). As a result, accesses to violatile data always miss in L1/nc. Since reading from nc doesn't require engaging the cache coherence protocol, it effectively bypasses "L1 overhead" as well. In that sense, the L1 cache is "disabled" for such accesses—not just by avoiding writing into it, but also by excluding L1 from the read overhead. The access will always miss in nc.

Or there is a better way to explain "disable" rather than "Do not allocate data to cache. (.level::eviction_priority)" for no_allocate in the PTX documentation? Thank you!

@zsnoob commented on GitHub (Mar 30, 2025): > I guess the PTX is not related to OS driver, but an undoced Hopper hardware feature: `L1::no_allocate` is not a hint but a strong option to disable the L1 cache (so that ensuring the volatile data read correctness). Brief discussion: In this context, considering both the repository and the documentation, could we interpret this "disable" as follows: for volatile data, the use of no_allocate makes it impossible for the data to be written into L1 (unified with nc to make sure correctness). As a result, accesses to violatile data always miss in L1/nc. Since reading from nc doesn't require engaging the cache coherence protocol, it effectively bypasses "L1 overhead" as well. In that sense, the L1 cache is "disabled" for such accesses—not just by avoiding writing into it, but also by excluding L1 from the read overhead. The access will always miss in nc. Or there is a better way to explain "disable" rather than "Do not allocate data to cache. (.level::eviction_priority)" for `no_allocate` in the PTX documentation? Thank you!

bsakaguchi commented

2026-10-26 09:52:38 +00:00

Author

Owner

@LyricZhao commented on GitHub (Mar 31, 2025):

You can read the README:

Undefined-behavior PTX usage

For extreme performance, we discover and use an undefined-behavior PTX usage: using read-only PTX ld.global.nc.L1::no_allocate.L2::256B to read volatile data. The PTX modifier .nc indicates that a non-coherent cache is used. But the correctness is tested to be guaranteed with .L1::no_allocate on Hopper architectures, and performance will be much better. The reason we guess may be: the non-coherent cache is unified with L1, and the L1 modifier is not just a hint but a strong option, so that the correctness can be guaranteed by no dirty data in L1.
Initially, because NVCC could not automatically unroll volatile read PTX, we tried using __ldg (i.e., ld.nc). Even compared to manually unrolled volatile reads, it was significantly faster (likely due to additional compiler optimizations). However, the results could be incorrect or dirty. After consulting the PTX documentation, we discovered that L1 and non-coherent cache are unified on Hopper architectures. We speculated that .L1::no_allocate might resolve the issue, leading to this discovery.
If you find kernels not working on some other platforms, you may add DISABLE_AGGRESSIVE_PTX_INSTRS=1 to setup.py and disable this, or file an issue.

@LyricZhao commented on GitHub (Mar 31, 2025): You can read the README: #### Undefined-behavior PTX usage - For extreme performance, we discover and use an undefined-behavior PTX usage: using read-only PTX `ld.global.nc.L1::no_allocate.L2::256B` to **read volatile data**. The PTX modifier `.nc` indicates that a non-coherent cache is used. But the correctness is tested to be guaranteed with `.L1::no_allocate` on Hopper architectures, and performance will be much better. The reason we guess may be: the non-coherent cache is unified with L1, and the L1 modifier is not just a hint but a strong option, so that the correctness can be guaranteed by no dirty data in L1. - Initially, because NVCC could not automatically unroll volatile read PTX, we tried using `__ldg` (i.e., `ld.nc`). Even compared to manually unrolled volatile reads, it was significantly faster (likely due to additional compiler optimizations). However, the results could be incorrect or dirty. After consulting the PTX documentation, we discovered that L1 and non-coherent cache are unified on Hopper architectures. We speculated that `.L1::no_allocate` might resolve the issue, leading to this discovery. - If you find kernels not working on some other platforms, you may add `DISABLE_AGGRESSIVE_PTX_INSTRS=1` to `setup.py` and disable this, or file an issue.

bsakaguchi commented

2026-10-26 09:52:38 +00:00

Author

Owner

@LyricZhao commented on GitHub (Mar 31, 2025):

We observe that ld.nc is much faster than ld, but it is wrong. As L1 and non-coherent cache are unified on Hopper architectures, so we tried ld.nc.L1::no_allocate and it worked (always not storing into L1 as you said).

But I guess ld.L1::no_allocate (no nc) also works, you can have a try. If so, .L1::no_allocate is used to ensure correctness, .nc encourages NVCC to do more optimizations. They are separate two options.

@LyricZhao commented on GitHub (Mar 31, 2025): We observe that `ld.nc` is much faster than `ld`, but it is wrong. As L1 and non-coherent cache are unified on Hopper architectures, so we tried `ld.nc.L1::no_allocate` and it worked (always not storing into L1 as you said). But I guess `ld.L1::no_allocate` (no `nc`) also works, you can have a try. If so, `.L1::no_allocate` is used to ensure correctness, `.nc` encourages NVCC to do more optimizations. They are separate two options.

bsakaguchi commented

2026-10-26 09:52:39 +00:00

Author

Owner

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

Thanks, very insightful for me!

@zsnoob commented on GitHub (Mar 31, 2025): Thanks, very insightful for me!

bsakaguchi referenced this issue

2026-10-26 09:53:39 +00:00

[GH-ISSUE #74] test_internode.py hangs at internode.cu:dispatch(...) #68

bsakaguchi referenced this issue

2026-10-26 09:56:29 +00:00