[Bugfix] Add kernel_global_source property to TVMFFIKernelAdapter

[haok1402]

fix #1576

JITKernel.kernel_source falls back to self.adapter.kernel_global_source when artifact is None:

tilelang/tilelang/jit/kernel.py

Lines 616 to 618 in dcacc5a

	@property
	def kernel_source(self) -> str:
	return self.artifact.kernel_source if self.artifact else self.adapter.kernel_global_source

which causes the following error,

Error saving kernel source code to disk: 'TVMFFIKernelAdapter' object has no attribute 'kernel_global_source'

TVMFFIKernelAdapter was missing this attribute.

Add kernel_global_source assignment in init, consistent with CuTeDSLKernelAdapter:

tilelang/tilelang/jit/adapter/cutedsl/adapter.py

Line 95 in dcacc5a

self.kernel_global_source = self.device_kernel_source

Summary by CodeRabbit

• Chores
  • Internal improvements to kernel source code handling for enhanced consistency.

_{✏️ Tip: You can customize this high-level summary in your review settings.}

[github-actions]

👋 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]

📝 Walkthrough

Walkthrough

This change fixes a missing attribute error in the TVMFFIKernelAdapter class by adding a single attribute assignment during initialization. The kernel_global_source attribute is now set to mirror the device_kernel_source value, enabling the autotuner cache dump functionality.

Changes

Cohort / File(s)	Summary
TVMFFIKernelAdapter attribute initialization tilelang/jit/adapter/tvm_ffi.py	Added self.kernel_global_source = self.device_kernel_source assignment in __init__ to provide the missing attribute referenced by cache serialization logic

Estimated code review effort

🎯 1 (Trivial) | ⏱️ ~2 minutes

Poem

🐰 A kernel attribute was lost in the night,
But a mirror reflects it back bright,
Now caches are saved with a hop and a bound,
The autotuner's treasure will never be drowned! ✨

Pre-merge checks and finishing touches

✅ Passed checks (5 passed)

Check name	Status	Explanation
Description Check	✅ Passed	Check skipped - CodeRabbit’s high-level summary is enabled.
Linked Issues check	✅ Passed	The PR directly addresses issue #1576 by adding the missing kernel_global_source attribute to TVMFFIKernelAdapter, which resolves the AttributeError preventing cache dumps.
Out of Scope Changes check	✅ Passed	All changes are directly related to fixing the linked issue by adding the kernel_global_source attribute; no out-of-scope modifications detected.
Docstring Coverage	✅ Passed	Docstring coverage is 100.00% which is sufficient. The required threshold is 80.00%.
Title check	✅ Passed	The title accurately describes the main change: adding a kernel_global_source property to TVMFFIKernelAdapter to fix the missing attribute issue in JITKernel.kernel_source fallback logic.

✨ Finishing touches

• 📝 Generate docstrings

---

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

• X
• Mastodon
• Reddit
• LinkedIn

_{Comment @coderabbitai help to get the list of available commands and usage tips.}

[coderabbitai]

Actionable comments posted: 0

Caution

Some comments are outside the diff and can’t be posted inline due to platform limitations.

⚠️ Outside diff range comments (2)

tilelang/jit/adapter/tvm_ffi.py (2)

37-55: Add kernel_global_source as a class attribute for consistency.

The new attribute kernel_global_source is set in __init__ at line 100 but is not declared as a class attribute like its counterparts host_kernel_source and device_kernel_source (lines 42–43). This inconsistency can affect type checking and readability.

🔎 Proposed fix to add class attribute declaration

  host_kernel_source: str | None = None
  device_kernel_source: str | None = None
+ kernel_global_source: str | None = None
  executable: tvm.runtime.Executable | None = None

---

264-297: Missing kernel_global_source assignment in from_database classmethod affects multiple adapters.

The from_database alternative constructor in TVMFFIKernelAdapter, CythonKernelAdapter, and NVRTCKernelAdapter does not set kernel_global_source, unlike their __init__ methods. When loaded from the autotuner cache via from_database, accessing the kernel_source property will raise an AttributeError. This affects the autotuner cache loading path confirmed in tilelang/autotuner/param.py and tilelang/cache/kernel_cache.py.

🔎 Proposed fix

Add the assignment in all three adapter classes' from_database methods:

  adapter.host_kernel_source = host_kernel_source
  adapter.device_kernel_source = device_kernel_source
+ adapter.kernel_global_source = device_kernel_source

(CuTeDSLKernelAdapter already has this assignment correctly at line 151.)

📜 Review details

Configuration used: defaults

Review profile: CHILL

Plan: Pro

📥 Commits

Reviewing files that changed from the base of the PR and between dcacc5a and 1b9b4a0.

📒 Files selected for processing (1)

• tilelang/jit/adapter/tvm_ffi.py

🧰 Additional context used

🧠 Learnings (2)

📓 Common learnings

Learnt from: lucifer1004
Repo: tile-ai/tilelang PR: 1483
File: tilelang/jit/adapter/cutedsl/adapter.py:93-95
Timestamp: 2025-12-26T06:45:51.789Z
Learning: For the CuTeDSL backend in tilelang/jit/adapter/cutedsl/adapter.py, the host_kernel_source and device_kernel_source have the same value.

📚 Learning: 2025-12-26T06:45:51.789Z

Learnt from: lucifer1004
Repo: tile-ai/tilelang PR: 1483
File: tilelang/jit/adapter/cutedsl/adapter.py:93-95
Timestamp: 2025-12-26T06:45:51.789Z
Learning: For the CuTeDSL backend in tilelang/jit/adapter/cutedsl/adapter.py, the host_kernel_source and device_kernel_source have the same value.

Applied to files:

• tilelang/jit/adapter/tvm_ffi.py

[LeiWang1999]

LGTM, Thanks @haok1402 !