[issue] st.bulk instruction requires sm_100, don't support pre-sm100 arch

[huakainomore]

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
}