[WIP] refactor: unifying return status of different backend implementation

[yzh119]

📌 Description

We currently support kernel implementations from multiple backends (Cutlass, TRT-LLM-Gen, etc.), each of which uses its own return codes or error types. Handling these inconsistently with separate macros (CUTLASS_CHECK, TLLM_CHECK, etc.) has become chaotic.

In This PR, we

• Introduces a single Status class that can wrap and implicitly convert from various backend error types (cudaError_t, cutlass::Status, etc.).
• Replaces all backend-specific check macros with one unified call macro: FLASHINFER_CALL.

Future API conventions

1. Return codes, not exceptions
   • APIs under include/ will return a Status instead of using C++ exceptions.
2. Unified call-site error handling
   • Every call (to FlashInfer or backend functions) must be wrapped in FLASHINFER_CALL.
3. Custom errors via Status
   • Replace inline exceptions with Status error codes using FLASHINFER_ERROR("message").
4. Exceptions in bindings
   • In binding code (csrc/), parse the returned Status and rethrow as C++ exceptions using TORCH_CHECK.

This centralizes error handling, makes it easy to add new backends, and aligns FlashInfer with HPC-library practices while still supporting exception-based checks at the Python/C++ interface.

🔍 Related Issues

🚀 Pull Request Checklist

Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete.

✅ Pre-commit Checks

• I have installed pre-commit by running pip install pre-commit (or used your preferred method).
• I have installed the hooks with pre-commit install.
• I have run the hooks manually with pre-commit run --all-files and fixed any reported issues.

If you are unsure about how to set up pre-commit, see the pre-commit documentation.

🧪 Tests

• Tests have been added or updated as needed.
• All tests are passing (unittest, etc.).

Reviewer Notes

[gemini-code-assist]

Summary of Changes

Hello @yzh119, I'm Gemini Code Assist¹! I'm currently reviewing this pull request and will post my feedback shortly. In the meantime, here's a summary to help you and other reviewers quickly get up to speed!

This pull request implements a significant refactor of the error handling mechanism within FlashInfer. The goal is to unify how errors from different hardware backends are reported and checked, moving from disparate backend-specific methods to a single, consistent Status class and a universal checking macro. This change centralizes error logic, simplifies the addition of new backends, and aligns the library's internal error reporting with common HPC practices, while preserving exception-based error handling at the user-facing binding layer.

Highlights

• Unified Error Handling: Introduced a new Status class to provide a consistent way to represent operation results across different backends (like CUDA and Cutlass).
• Single Check Macro: Replaced backend-specific error checking macros (CUTLASS_CHECK, FLASHINFER_CUDA_CALL) with a single unified macro, FLASHINFER_CALL, which works with the new Status class.
• API Convention Changes: Outlined future API conventions where core library functions will return Status objects, and error-to-exception conversion will be handled in the Python/C++ bindings layer.
• Codebase Refactor: Updated numerous function signatures and call sites across the codebase (csrc/ and include/) to use the new Status return type and the FLASHINFER_CALL macro.

Using Gemini Code Assist

The full guide for Gemini Code Assist can be found on our documentation page, here are some quick tips.

Invoking Gemini

You can request assistance from Gemini at any point in your pull request via creating an issue comment (i.e. comment on the pull request page) using either /gemini <command> or @gemini-code-assist <command>. Below is a summary of the supported commands.

Feature	Command	Description
Code Review	/gemini review	Performs a code review for the current pull request in its current state.
Pull Request Summary	/gemini summary	Provides a summary of the current pull request in its current state.
Comment	@gemini-code-assist	Responds in comments when explicitly tagged, both in issue comments and review comments.
Help	/gemini help	Displays a list of available commands.

Customization

To customize Gemini Code Assist for GitHub experience, repository maintainers can create a configuration file and/or provide a custom code review style guide (such as PEP-8 for Python) by creating and adding files to a .gemini/ folder in the base of the repository. Detailed instructions can be found here.

Limitations & Feedback

Gemini Code Assist is currently in preview and may make mistakes. Please leave feedback on any instances where its feedback is incorrect or counter productive. You can react with 👍 and 👎 on @gemini-code-assist comments to provide feedback.

You can also get AI-powered code generation, chat, as well as code reviews directly in the IDE at no cost with the Gemini Code Assist IDE Extension.

Footnotes

1. Review the Privacy Notices, Generative AI Prohibited Use Policy, Terms of Service, and learn how to configureGemini Code Assist in GitHub here. Gemini can make mistakes, so double check it and use code with caution. ↩

[gemini-code-assist]

Code Review

This pull request introduces a unified Status class and FLASHINFER_CALL macro to standardize error handling across different backends like CUDA and CUTLASS. This is a positive step towards improving code clarity and maintainability.

The changes largely involve replacing cudaError_t with Status and backend-specific check macros with FLASHINFER_CALL. The new Status class and its integration into the core library functions (.cuh files) seem correct.

A key area for review is the update of TORCH_CHECK macros in the csrc/*.cu files. Many of these still use status == cudaSuccess and cudaGetErrorString(status) after status has been changed to the Status type. These need to be updated to status.success() and status.error_message() to ensure correct error handling and prevent compilation issues. One file (csrc/group_gemm_sm90.cu) correctly demonstrates this updated usage.

Overall, the refactoring is well-structured. Addressing the TORCH_CHECK inconsistencies will complete the integration of the new error handling system.