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

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
__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(__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(smem_ptr)[i] = 0;
}
#endif
}
```

Start at the copy_epilogue code that uses the st.bulk instruction and inspect the shown architecture guard and fallback path. Verify that zero padding works on pre-sm100 architectures without using st.bulk, while retaining the instruction for sm100 and newer builds.

Written by the indexing model from the issue text.