关于 iterator_sm70 文件中 copy 函数调用的作用

[vicety]

// src/turbomind/kernels/attention/iterator_sm70.h
template<bool is_residue, class TileIter>
    __device__ void Load(const TileIter& tile_iter, Fragment& rmem, int max_s)
    {
        auto src_data = tile_iter.OffsetPtr<Idx>(src_offset_);
        int  offset_s = Map::get_offset(threadIdx.x / WARP_SIZE, threadIdx.x % WARP_SIZE).y;
        PRAGMA_UNROLL
        for (int s = 0; s < Map::kIterS; ++s) {
            PRAGMA_UNROLL
            for (int c = 0; c < Map::kIterC; ++c) {
                copy(Array<T, Map::kAccessC>{}, rmem[s][c]);
                // ...

我理解整个函数的目的是从 global memory 中加载 kvcache 到 rmem，rmem 的类型是 Array<T, Map::kAccessC>[Map::kIterS][Map::kIterC]。关于这一行 copy(Array<T, Map::kAccessC>{}, rmem[s][c]);，为什么需要从一个临时的 Array 拷贝内容去覆盖 rmem 的内容呢？

如果移除这一行，程序会 crash；如果改为用任意值初始化 rmem[s][c]，似乎也能得到正确结果。

[lzhangzz]

这里是提前清零越界值，否则可能会有 NAN 或者 INF。随便设个值也行估计是因为会被 attention score mask 掉。