Skip to content

[BugFix] Skip source-compilation options when exporting LLVM module#2467

Open
penguin-wwy wants to merge 1 commit into
tile-ai:mainfrom
penguin-wwy:fix_cache
Open

[BugFix] Skip source-compilation options when exporting LLVM module#2467
penguin-wwy wants to merge 1 commit into
tile-ai:mainfrom
penguin-wwy:fix_cache

Conversation

@penguin-wwy

@penguin-wwy penguin-wwy commented Jun 27, 2026

Copy link
Copy Markdown
Contributor

Running the LLVM backend tests on macOS produces a cache-save compilation error, even though the kernel itself compiles and runs correctly:

PYTHONPATH=$(pwd) python -m pytest testing/python/llvm/test_tilelang_llvm_gemm.py -s
The test passes (kernel executes correctly), but the disk-cache save fails with:
[TileLang:tilelang.cache.kernel_cache:ERROR] Error during atomic cache save Compilation error:
/var/folders/.../lib0.o:1:1: error: source file is not valid UTF-8
/var/folders/.../lib1.o:1:1: error: source file is not valid UTF-8
/var/folders/.../devc.o:1:1: error: source file is not valid UTF-8
...

The failing command (from the log):

/usr/bin/g++ -shared -fPIC -undefined dynamic_lookup -o ...so \
  -x objective-c++ -g -std=gnu++17 -I...torch/include \
  lib0.o lib1.o devc.o

The binary .o object files are being parsed as source text.

The error is triggered by the interaction between tilelang's cache-save path and TVM's export_library:

  • target=llvm: TVM's export_library emits binary .o object files and constructs a link command (-shared -fPIC -undefined dynamic_lookup) to combine them into a .so. No source compilation is involved.
  • tilelang on macOS: _get_compile_args() unconditionally returns -x objective-c++ (originally intended for the Metal codegen path, which emits C source).

Summary

  • Fixed macOS kernel cache export failures for the LLVM backend by ensuring LLVM exports don’t receive Metal/source-compilation flags (e.g., -x objective-c++) when exporting to a shared library.
  • Updated KernelCache._safe_write_executable to accept an optional target and to strip the "options" compile-arg entry for darwin + target.kind.name == "llvm" before calling TVM export_library.
  • Passed the kernel’s target through the JIT adapter (_save_so_cubin_to_disk_safe_write_executable) so the LLVM-specific export adjustment can be applied during cache save.

@github-actions

Copy link
Copy Markdown

👋 Hi! Thank you for contributing to the TileLang project.

Please remember to run pre-commit run --all-files in the root directory of the project to ensure your changes are properly linted and formatted. This will help ensure your contribution passes the format check.

We appreciate you taking this step! Our team will review your contribution, and we look forward to your awesome work! 🚀

@coderabbitai

coderabbitai Bot commented Jun 27, 2026

Copy link
Copy Markdown
Contributor

Review Change Stack

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Pro

Run ID: fa03252c-9392-4738-b6ed-a3a75a830fb2

📥 Commits

Reviewing files that changed from the base of the PR and between ce7871d and 46775ae.

📒 Files selected for processing (2)
  • tilelang/cache/kernel_cache.py
  • tilelang/jit/adapter/kernel_cache.py
🚧 Files skipped from review as they are similar to previous changes (1)
  • tilelang/jit/adapter/kernel_cache.py

📝 Walkthrough

Walkthrough

The PR updates executable export handling to skip compile arguments for LLVM targets and passes the kernel target through the adapter when saving the executable.

Changes

Target-aware kernel export

Layer / File(s) Summary
Target-aware export path
tilelang/cache/kernel_cache.py, tilelang/jit/adapter/kernel_cache.py
KernelCache._safe_write_executable now accepts an optional target and conditionally omits compile arguments for LLVM exports; the JIT adapter passes kernel.target into that call.

🎯 2 (Simple) | ⏱️ ~10 minutes

A rabbit hopped by with a kernel and grin,
"LLVM gets a gentler export from within!"
Hop-hop went the cache, target in sight,
No extra args for LLVM tonight 🐇

🚥 Pre-merge checks | ✅ 4 | ❌ 1

❌ Failed checks (1 warning)

Check name Status Explanation Resolution
Docstring Coverage ⚠️ Warning Docstring coverage is 25.00% which is insufficient. The required threshold is 80.00%. Write docstrings for the functions missing them to satisfy the coverage threshold.
✅ Passed checks (4 passed)
Check name Status Explanation
Description Check ✅ Passed Check skipped - CodeRabbit’s high-level summary is enabled.
Title check ✅ Passed The title clearly matches the main fix: skipping source-compilation options during LLVM module export.
Linked Issues check ✅ Passed Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check ✅ Passed Check skipped because no linked issues were found for this pull request.
✨ Finishing Touches
🧪 Generate unit tests (beta)
  • Create PR with unit tests

Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out.

❤️ Share

Comment @coderabbitai help to get the list of available commands.

@penguin-wwy penguin-wwy changed the title [BugFix] Skip source-compilation options when exporting LLVM module c… [BugFix] Skip source-compilation options when exporting LLVM module Jun 27, 2026

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actionable comments posted: 1

🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Inline comments:
In `@tilelang/cache/kernel_cache.py`:
- Around line 462-469: The LLVM export gate in kernel_cache.py is too broad: it
skips all compile args for LLVM targets and accidentally drops the Windows
fcompile path from _get_compile_args(). Update the export logic around
cls._get_compile_args() and executable.export_library so only the Darwin
source-language options are suppressed for LLVM, while preserving the Windows
msvc_create_shared/fcompile hook and any other non-Darwin compile settings.
🪄 Autofix (Beta)

Fix all unresolved CodeRabbit comments on this PR:

  • Push a commit to this branch (recommended)
  • Create a new PR with the fixes

ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Pro

Run ID: 8aa43433-21e4-4334-b350-10483c601444

📥 Commits

Reviewing files that changed from the base of the PR and between 46f3d1c and ce7871d.

📒 Files selected for processing (2)
  • tilelang/cache/kernel_cache.py
  • tilelang/jit/adapter/kernel_cache.py

Comment thread tilelang/cache/kernel_cache.py Outdated
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant