[GH-ISSUE #675] [issue] st.bulk instruction requires sm_100, don't support pre-sm100 arch #397

New issue

Open

opened 2026-10-26 10:00:03 +00:00 by bsakaguchi · 0 comments

bsakaguchi commented

2026-10-26 10:00:03 +00:00

Owner

Originally created by @huakainomore on GitHub (Jul 3, 2026).
Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/675

in latest code, use st.bulk to do Zero padding in copy_epilogue, but it requires sm_100 or higher. pre-sm100 arch support may like this.

template <int kNumBytes>
__forceinline__ __device__ void st_bulk(void* smem_ptr) {
    EP_STATIC_ASSERT(kNumBytes % 8 == 0, "`st.bulk` requires size to be a multiple of 8");
#if defined(__CUDA_ARCH__) and (__CUDA_ARCH__ >= 1000)
    if (ptx::elect_one_sync())
        asm volatile("st.bulk.weak.shared::cta [%0], %1, 0;\n" ::
                    "r"(static_cast<uint32_t>(__cvta_generic_to_shared(smem_ptr))),
                    "r"(kNumBytes)
                    : "memory");
#else
    #pragma unroll
    for (int i = get_lane_idx(); i < kNumBytes/8; i += 32) {
        reinterpret_cast<uint64_t*>(smem_ptr)[i] = 0;
    }
#endif
}

Originally created by @huakainomore on GitHub (Jul 3, 2026). Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/675 in latest code, use `st.bulk` to do Zero padding in copy_epilogue, but it requires sm_100 or higher. pre-sm100 arch support may like this. ``` template <int kNumBytes> __forceinline__ __device__ void st_bulk(void* smem_ptr) { EP_STATIC_ASSERT(kNumBytes % 8 == 0, "`st.bulk` requires size to be a multiple of 8"); #if defined(__CUDA_ARCH__) and (__CUDA_ARCH__ >= 1000) if (ptx::elect_one_sync()) asm volatile("st.bulk.weak.shared::cta [%0], %1, 0;\n" :: "r"(static_cast<uint32_t>(__cvta_generic_to_shared(smem_ptr))), "r"(kNumBytes) : "memory"); #else #pragma unroll for (int i = get_lane_idx(); i < kNumBytes/8; i += 32) { reinterpret_cast<uint64_t*>(smem_ptr)[i] = 0; } #endif } ```

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#397

No description provided.

Rows
Columns