From 287bf3a9f2c844dfb209d37a590e6152d26090f1 Mon Sep 17 00:00:00 2001 From: LeiWang1999 Date: Tue, 2 Sep 2025 16:08:35 +0800 Subject: [PATCH 1/6] Fix type hint for target_host parameter in compile function to allow None value --- tilelang/jit/__init__.py | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/tilelang/jit/__init__.py b/tilelang/jit/__init__.py index 8f9a4a3819..b64ff134a1 100644 --- a/tilelang/jit/__init__.py +++ b/tilelang/jit/__init__.py @@ -34,7 +34,7 @@ def compile( out_idx: Union[List[int], int, None] = None, execution_backend: Literal["dlpack", "ctypes", "cython", "nvrtc"] = "cython", target: Union[str, Target] = "auto", - target_host: Union[str, Target] = None, + target_host: Union[str, Target, None] = None, verbose: bool = False, pass_configs: Optional[Dict[str, Any]] = None, compile_flags: Optional[Union[List[str], str]] = None, @@ -69,6 +69,10 @@ def compile( assert isinstance(func, PrimFunc), f"target function must be a PrimFunc but got {type(func)}" if isinstance(compile_flags, str): compile_flags = [compile_flags] + + # This path is not a performance critical path, so we can afford to convert the target. + target, target_host = Target(target), Target(target_host) if target_host else None + return cached( func=func, out_idx=out_idx, From 5f44aef7383952d694b3c1afe758c05b22f3e186 Mon Sep 17 00:00:00 2001 From: LeiWang1999 Date: Tue, 2 Sep 2025 16:21:03 +0800 Subject: [PATCH 2/6] Refactor target handling in compile function to utilize determine_target for improved clarity and consistency --- tilelang/jit/__init__.py | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/tilelang/jit/__init__.py b/tilelang/jit/__init__.py index b64ff134a1..4d9edd54c1 100644 --- a/tilelang/jit/__init__.py +++ b/tilelang/jit/__init__.py @@ -20,6 +20,7 @@ from tvm.target import Target from tilelang.jit.kernel import JITKernel +from tilelang.utils.target import determine_target from tilelang.cache import cached from os import path, makedirs from logging import getLogger @@ -71,7 +72,7 @@ def compile( compile_flags = [compile_flags] # This path is not a performance critical path, so we can afford to convert the target. - target, target_host = Target(target), Target(target_host) if target_host else None + target = Target(determine_target(target)) return cached( func=func, From b2e85c091c3f0317ee37ee07aa6d9e556d9b87e8 Mon Sep 17 00:00:00 2001 From: LeiWang1999 Date: Tue, 2 Sep 2025 16:47:57 +0800 Subject: [PATCH 3/6] Update PrintConst function in codegen_cuda.cc to use hexfloat format for bfloat16 and float8/float4 types, while adding scientific notation comments for clarity. This change enhances the representation of floating-point constants in the generated code. --- src/target/codegen_cuda.cc | 11 ++++++++--- 1 file changed, 8 insertions(+), 3 deletions(-) diff --git a/src/target/codegen_cuda.cc b/src/target/codegen_cuda.cc index d2826f6ef9..12bf0dbd14 100644 --- a/src/target/codegen_cuda.cc +++ b/src/target/codegen_cuda.cc @@ -1960,13 +1960,17 @@ inline void PrintConst(const FloatImmNode *op, std::ostream &os, // Type code is kBFloat if (op->dtype.is_bfloat16()) { os << "bfloat16_t"; - os << '(' << std::scientific << op->value << 'f' << ')'; + os << '(' << std::hexfloat << op->value << 'f'; + os << "/*" << std::scientific << op->value << "*/"; + os << ')'; return; } // Type code is kFloat8_e5m2 or kE4M4Float if (op->dtype.is_float8() || op->dtype.is_float4()) { p->PrintType(op->dtype, os); - os << '(' << std::scientific << op->value << 'f' << ')'; + os << '(' << std::hexfloat << op->value << 'f'; + os << "/*" << std::scientific << op->value << "*/"; + os << ')'; return; } // Type code is kFloat @@ -1984,9 +1988,10 @@ inline void PrintConst(const FloatImmNode *op, std::ostream &os, temp << ((op->dtype.bits() == 32) ? "CUDART_NAN_F" : "CUDART_NAN"); p->need_math_constants_h_ = true; } else { - temp << std::scientific << op->value; + temp << std::hexfloat << op->value; if (op->dtype.bits() == 32) temp << 'f'; + temp << "/*" << std::scientific << op->value << "*/"; } p->MarkConst(temp.str()); os << temp.str(); From aeff94d5b1aff5de91af1b029cdb3bd4f76cbc10 Mon Sep 17 00:00:00 2001 From: LeiWang1999 Date: Tue, 2 Sep 2025 16:50:45 +0800 Subject: [PATCH 4/6] Refactor PrintType function in codegen_cuda.cc to remove unnecessary failure conditions for floating-point types with lane counts greater than 4. This change simplifies the logic and improves code clarity. --- src/target/codegen_cuda.cc | 4 ---- 1 file changed, 4 deletions(-) diff --git a/src/target/codegen_cuda.cc b/src/target/codegen_cuda.cc index 12bf0dbd14..2a4bb9c17e 100644 --- a/src/target/codegen_cuda.cc +++ b/src/target/codegen_cuda.cc @@ -325,16 +325,12 @@ void CodeGenTileLangCUDA::PrintType(DataType t, std::ostream &os) { // NOLINT(*) enable_fp6_ = true; if (t.lanes() <= 4) { os << GetFP6Type(t); - } else { - fail = true; } return; } else if (t.is_float4()) { enable_fp4_ = true; if (t.lanes() <= 4) { os << GetFP4Type(t); - } else { - fail = true; } return; } else if (t == DataType::Bool()) { From 200e366b57aa7c6c8ec240fdfbc7c41840b4d42d Mon Sep 17 00:00:00 2001 From: LeiWang1999 Date: Tue, 2 Sep 2025 19:12:01 +0800 Subject: [PATCH 5/6] Enhance benchmark_matmul.py to conditionally print Reference TFlops only if ref_latency is not None. Update param.py to ensure target is converted to string for consistency. Refactor tuner.py to utilize determine_target for improved clarity in target handling. --- benchmark/matmul/benchmark_matmul.py | 3 ++- tilelang/autotuner/param.py | 2 +- tilelang/autotuner/tuner.py | 3 ++- 3 files changed, 5 insertions(+), 3 deletions(-) diff --git a/benchmark/matmul/benchmark_matmul.py b/benchmark/matmul/benchmark_matmul.py index 39063b6f26..981f0225f7 100644 --- a/benchmark/matmul/benchmark_matmul.py +++ b/benchmark/matmul/benchmark_matmul.py @@ -243,4 +243,5 @@ def main( print(f"Best TFlops: {total_flops / best_latency * 1e-9:.3f}") print(f"Best config: {best_config}") - print(f"Reference TFlops: {total_flops / ref_latency * 1e-9:.3f}") + if ref_latency is not None: + print(f"Reference TFlops: {total_flops / ref_latency * 1e-9:.3f}") diff --git a/tilelang/autotuner/param.py b/tilelang/autotuner/param.py index fcf9eb7ff2..5807b8c77b 100644 --- a/tilelang/autotuner/param.py +++ b/tilelang/autotuner/param.py @@ -68,7 +68,7 @@ def __hash__(self): "execution_backend": self.execution_backend, "target": - self.target, + str(self.target), "target_host": str(self.target_host) if self.target_host else None, "verbose": diff --git a/tilelang/autotuner/tuner.py b/tilelang/autotuner/tuner.py index 2ed38c58c6..9078884a5f 100644 --- a/tilelang/autotuner/tuner.py +++ b/tilelang/autotuner/tuner.py @@ -28,6 +28,7 @@ from tilelang import env from tilelang.autotuner.param import CompileArgs, ProfileArgs, AutotuneResult from tilelang.autotuner.capture import get_autotune_inputs +from tilelang.utils.target import determine_target from tilelang.jit.param import _P, _RProg from tilelang.version import __version__ @@ -150,7 +151,7 @@ def set_compile_args(self, """ self.compile_args = CompileArgs( out_idx=out_idx, - target=target, + target=Target(determine_target(target)), execution_backend=execution_backend, target_host=target_host, verbose=verbose, From 0b4dc0751961d66dba46ab83613aded273609063 Mon Sep 17 00:00:00 2001 From: LeiWang1999 Date: Tue, 2 Sep 2025 19:17:46 +0800 Subject: [PATCH 6/6] Remove automatic commit and push step from AMD and NVIDIA CI workflows to streamline the process and avoid unnecessary commits. --- .github/workflows/amd_ci.yml | 5 ----- .github/workflows/ci.yml | 5 ----- 2 files changed, 10 deletions(-) diff --git a/.github/workflows/amd_ci.yml b/.github/workflows/amd_ci.yml index 784f342084..49e7037984 100644 --- a/.github/workflows/amd_ci.yml +++ b/.github/workflows/amd_ci.yml @@ -60,11 +60,6 @@ jobs: exit 1 fi rm -rf build - - - name: Commit and Push Changes - uses: stefanzweifel/git-auto-commit-action@v5 - with: - commit_message: "lint" build-test-amd: runs-on: [self-hosted, amd, gpu] diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 0826e5d3a2..541931cedf 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -60,11 +60,6 @@ jobs: exit 1 fi rm -rf build - - - name: Commit and Push Changes - uses: stefanzweifel/git-auto-commit-action@v5 - with: - commit_message: "lint" build-test-nvidia: runs-on: [self-hosted, nvidia]