Skip to content

[TLE][MTHREADS] Support transposed operands in tle.wgmma #837

[TLE][MTHREADS] Support transposed operands in tle.wgmma

[TLE][MTHREADS] Support transposed operands in tle.wgmma #837

# Copyright 2025- FlagOS Contributors
#
# Permission is hereby granted, free of charge, to any person obtaining a copy
# of this software and associated documentation files (the "Software"), to deal
# in the Software without restriction, including without limitation the rights
# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
# copies of the Software, and to permit persons to whom the Software is
# furnished to do so, subject to the following conditions:
#
# The above copyright notice and this permission notice shall be included in all
# copies or substantial portions of the Software.
#
# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
# SOFTWARE.
name: Iluvatar3.6-Build-And-Test
on:
push:
branches: [ "main", "triton_v3.6.x" ]
pull_request:
branches: [ "main", "triton_v3.6.x" ]
types: [opened, synchronize, reopened, ready_for_review]
permissions:
contents: read
pull-requests: read
concurrency:
group: ${{ github.workflow }}-${{ github.event.pull_request.number || github.ref }}
cancel-in-progress: true
jobs:
iluvatar36x-unit-test:
runs-on: iluvatar3.6
if: ${{ github.repository == 'flagos-ai/flagtree' && github.event.pull_request.draft == false }}
steps:
- name: Setup environment
shell: bash
run: |
source ~/env.sh
env | grep -E '^(http_proxy|https_proxy|all_proxy|no_proxy)=' >> $GITHUB_ENV || true
- name: Smart Checkout
uses: flagos-ai/FlagTree/.github/actions/smart-checkout@main
with:
checkout_version: 'v6'
- name: Check if backend-relevant files changed
id: check_backend
uses: flagos-ai/FlagTree/.github/actions/check-backend-changed@main
with:
backend: iluvatar
- name: Cleanup FlagTree Environment
if: steps.check_backend.outputs.should_skip != 'true'
uses: flagos-ai/FlagTree/.github/actions/cleanup-flagtree-env@main
- name: Build FlagTree
if: steps.check_backend.outputs.should_skip != 'true'
shell: bash
run: |
set -x
export FLAGTREE_BACKEND=iluvatar
export TRITON_OFFLINE_BUILD=1
export TRITON_BUILD_PROTON=OFF
export TRITON_ILU_BUILD_GLUON=OFF
MAX_JOBS=32 python3 -m pip install . --no-build-isolation -v
- name: Clear cache if ClearCache label present
if: steps.check_backend.outputs.should_skip != 'true'
id: clear_cache
uses: flagos-ai/FlagTree/.github/actions/clear-cache-if-labeled@main
with:
github-token: ${{ secrets.GITHUB_TOKEN }}
backend: iluvatar
- name: Unit Test
if: steps.check_backend.outputs.should_skip != 'true'
shell: bash
run: |
set -x
source ~/env.sh
pip3 install "pytest==9.1.1"
for pkg in hypothesis absl-py scipy pytest-forked expecttest; do
pip3 list "$pkg" | grep "$pkg" || pip3 install "$pkg"
done
if [[ "$(uname -m)" == "aarch64" ]]; then
libgomp_path=$(find /usr/lib /usr/lib64 /lib /lib64 -type f -name 'libgomp.so*' 2>/dev/null | head -n 1)
if [[ -n "$libgomp_path" ]]; then
export LD_PRELOAD="$libgomp_path${LD_PRELOAD:+:$LD_PRELOAD}"
fi
fi
# triton language unit test
UMD_CUDAMODULELOADING=0 pytest -v third_party/iluvatar/python/test/unit/language/test_core.py
pytest -v third_party/iluvatar/python/test/unit/language/test_annotations.py
pytest -v third_party/iluvatar/python/test/unit/language/test_block_pointer.py
pytest -v third_party/iluvatar/python/test/unit/language/test_compile_errors.py
pytest -v third_party/iluvatar/python/test/unit/language/test_conversions.py
pytest -v third_party/iluvatar/python/test/unit/language/test_decorator.py
pytest -v third_party/iluvatar/python/test/unit/language/test_frontend.py
pytest -v third_party/iluvatar/python/test/unit/language/test_libdevice.py
pytest -v third_party/iluvatar/python/test/unit/language/test_iluvatar_math_fp16_bf16.py
pytest -v third_party/iluvatar/python/test/unit/language/test_iluvatar_int8_upcast_dot_pipeline.py
pytest -v third_party/iluvatar/python/test/unit/language/test_matmul.py
pytest -v third_party/iluvatar/python/test/unit/language/test_mxfp.py
pytest -v third_party/iluvatar/python/test/unit/language/test_pipeliner.py
pytest -v third_party/iluvatar/python/test/unit/language/test_random.py
pytest -v third_party/iluvatar/python/test/unit/language/test_reproducer.py
pytest -v third_party/iluvatar/python/test/unit/language/test_standard.py
pytest -v third_party/iluvatar/python/test/unit/language/test_subprocess.py
UMD_CUDAMODULELOADING=0 pytest -v third_party/iluvatar/python/test/unit/language/test_tensor_descriptor.py
pytest -v third_party/iluvatar/python/test/unit/language/test_tuple.py
pytest -v third_party/iluvatar/python/test/unit/language/test_warp_specialization.py
# triton operators unit test
pytest -v third_party/iluvatar/python/test/unit/operators/test_blocksparse.py
pytest -v third_party/iluvatar/python/test/unit/operators/test_cross_entropy.py
pytest -v third_party/iluvatar/python/test/unit/operators/test_dot_trans.py
pytest -v third_party/iluvatar/python/test/unit/operators/test_flash_attention.py
pytest -v third_party/iluvatar/python/test/unit/operators/test_inductor.py
pytest -v third_party/iluvatar/python/test/unit/operators/test_matmul.py
pytest -v third_party/iluvatar/python/test/unit/operators/test_sme.py
# triton runtime unit test
pytest -v third_party/iluvatar/python/test/unit/runtime/test_autotuner.py
pytest -v third_party/iluvatar/python/test/unit/runtime/test_bindings.py
pytest -v third_party/iluvatar/python/test/unit/runtime/test_build.py
pytest -v third_party/iluvatar/python/test/unit/runtime/test_cache.py
pytest -v third_party/iluvatar/python/test/unit/runtime/test_compilation_listener.py
pytest -v third_party/iluvatar/python/test/unit/runtime/test_driver.py
pytest -v third_party/iluvatar/python/test/unit/runtime/test_launch.py
pytest -v third_party/iluvatar/python/test/unit/runtime/test_specialize.py
pytest -v third_party/iluvatar/python/test/unit/runtime/test_subproc.py
pytest -v third_party/iluvatar/python/test/unit/runtime/test_iluvatar_loop_unroll_warning.py
# triton instrumentation unit test
pytest -v third_party/iluvatar/python/test/unit/instrumentation/test_gpuhello.py
# triton tools unit test
pytest -v third_party/iluvatar/python/test/unit/tools/test_irsource.py
pytest -v third_party/iluvatar/python/test/unit/tools/test_linear_layout.py
# triton misc unit test
pytest -v third_party/iluvatar/python/test/unit/test_debug_dump.py
pytest -v third_party/iluvatar/python/test/unit/test_debug.py
pytest -v third_party/iluvatar/python/test/unit/test_debuginfo.py
pytest -v third_party/iluvatar/python/test/unit/test_filecheck.py
pytest -v third_party/iluvatar/python/test/unit/test_knobs.py
UMD_CUDAMODULELOADING=0 pytest -v third_party/iluvatar/python/test/unit/test_link.py
pytest -v third_party/iluvatar/python/test/unit/test_perf_warning.py
# triton regression test
pytest -v third_party/iluvatar/python/test/regression/test_cast_matmul.py
pytest -v third_party/iluvatar/python/test/regression/test_functional_regressions.py
# triton integrations unit test
pytest -v third_party/iluvatar/python/test/unit/integrations/vllm/chunk_o/test_chunk_fwd_kernel_o.py
pytest -v third_party/iluvatar/python/test/unit/integrations/vllm/wy_fast/test_recompute_w_u.py
pytest -v third_party/iluvatar/python/test/unit/integrations/fbgemm/test_jagged_flash_attention_bwd_basic_min.py
PUNICA_TEST_LEVEL=quick pytest -v third_party/iluvatar/python/test/unit/integrations/vllm/punica_lora/test_punica_ops.py
pytest -v third_party/iluvatar/python/test/unit/integrations/sglang/flash_mla/test_flash_mla_ut.py
pytest -v third_party/iluvatar/python/test/unit/integrations/inductor/test_bucketize_matmul.py
pytest -v third_party/iluvatar/python/test/unit/integrations/inductor/test_swfw3103_flex_attention_precision.py
# TLE unit test
pytest -v python/test/tle/integration/test_tle_local_store.py
pytest -v python/test/tle/unit/test_tle_gpu_local_ptr.py
pytest -v python/test/tle/unit/test_extract_tile_static_index.py
pytest -v python/test/tle/unit/test_extract_tile_dynamic_index.py
pytest -v python/test/tle/unit/test_insert_tile_static_index.py
pytest -v python/test/tle/unit/test_insert_tile_dynamic_index.py
pytest -v python/test/tle/unit/test_tle.py
pytest -v third_party/iluvatar/python/test/unit/tle/test_tle_copy.py
pytest -v third_party/iluvatar/python/test/unit/tle/test_tle_async_load.py
pytest -v third_party/iluvatar/python/test/unit/tle/test_tle_memory_space.py
pytest -v third_party/iluvatar/python/test/unit/tle/test_tle_cumsum.py
pytest -v third_party/iluvatar/python/test/unit/tle/test_tle_pipeline.py
pytest -v third_party/iluvatar/python/test/unit/tle/test_tle_pipeline_e2e.py
pytest -v third_party/iluvatar/python/test/unit/tle/test_tle_warp_specialize.py
pytest -v third_party/iluvatar/python/test/unit/tle/test_tle_pipe.py
# TLE tutorials unit tests
python3 third_party/iluvatar/python/tutorials/tle/01-fft.py --only_unit_test
python3 third_party/iluvatar/python/tutorials/tle/03-topk.py --only_unit_test