GitLab

Signed-off-by: lijian <lijian6@sugon.com>

fix rocshmem compile err and update rocshmem_dir.

See merge request dcutoolkit/deeplearing/DeepEP!4

Signed-off-by: lijian <lijian6@sugon.com>

支持int8类型的kernel接口

See merge request dcutoolkit/deeplearing/DeepEP!3

Signed-off-by: lijian <lijian6@sugon.com>

Signed-off-by: lijian <lijian6@sugon.com>

Rocshmem dushmem

See merge request dcutoolkit/deeplearing/DeepEP!2

2. Add internode ll mode.
3. Add test internode ll mode.
Signed-off-by: lijian <lijian6@sugon.com>

2. 修改测试脚本，降低显存占用。使用量从17G -> 8G.
Signed-off-by: lijian <lijian6@sugon.com>

Signed-off-by: lijian <lijian6@sugon.com>

     backup pass way.
Signed-off-by: lishen <lishen@sugon.com>

Signed-off-by: lijian <lijian6@sugon.com>

Signed-off-by: lijian <lijian6@sugon.com>

Signed-off-by: lijian <lijian6@sugon.com>

Co-authored-by: Yifei Zhang <219273404+yifeizhang-c@users.noreply.github.com>

* Fix hidden_size % 128 != 0

* Add `align_down()` function

* Use the full warp to wait TMA store

* Support arbitrary hidden sizes in fp8 cast

* lint

* Remove redundant TMA flushes

* Less barrier initialization overhead

* Simplify `elect_one_sync`

* Use `elect_one_sync` instead of lanes

* Minor fix

* Polish testing prints

* Refactor for internode kernels

* Better performance

* Suppress kineto output

* Add pressure test mode

* Add `x_pure_rand_e4m3` test

* Add more results into hash value

* Fix mbarrier

* Remove redundant store

Each thread is responsible for one target rank

* nits

* hack unrolled warp copy

* Revert "nits"

This reverts commit 3e1b28d9b17f2c1cc46403d432ca576dbf15bd45.