From dd0f87261a80cf0ce8aa66e4ab2041843851d810 Mon Sep 17 00:00:00 2001 From: Keita Watanabe Date: Fri, 21 Aug 2026 07:34:25 +0000 Subject: [PATCH] fix: use a 64-bit operand for st.bulk's size in st_bulk() PTX ISA 9.0 added support for a 32-bit st.bulk size operand, but ptxas 13.0.88 rejects the register form even when nvcc emits PTX .version 9.0. This breaks the affected JIT kernels on sm_100-family GPUs with CUDA 13.0. Pass kNumBytes through a 64-bit "l" constraint instead. The value is unchanged, and the 64-bit form works with both CUDA 13.0 and CUDA 13.1. A CUDA 13.1.80 compile probe emits PTX .version 9.1 and accepts both the 32-bit and 64-bit register forms. Observed on 2 x p6-b300.48xlarge instances (B300, sm_103): tests/elastic/test_ep.py fails at the first dispatch on every rank with CUDA 13.0. With this fix, the 2-node internode run of tests/elastic/test_ep.py --test-first-only passes end-to-end over EFA-GDA, with RUN_EXIT=0 on both nodes and all correctness checks green. Signed-off-by: Keita Watanabe --- deep_ep/include/deep_ep/common/ptx.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/deep_ep/include/deep_ep/common/ptx.cuh b/deep_ep/include/deep_ep/common/ptx.cuh index ac8b1489..dc29546b 100644 --- a/deep_ep/include/deep_ep/common/ptx.cuh +++ b/deep_ep/include/deep_ep/common/ptx.cuh @@ -107,7 +107,7 @@ __forceinline__ __device__ void st_bulk(void* smem_ptr) { if (elect_one_sync()) { asm volatile("st.bulk.weak.shared::cta [%0], %1, 0;\n" :: "r"(static_cast(__cvta_generic_to_shared(smem_ptr))), - "r"(kNumBytes) + "l"(static_cast(kNumBytes)) : "memory"); } #else