bugfix: hotfix of PR 2366 (mamba kernel)

[yzh119]

📌 Description

The commit 1938c5c we added to #2366 cause hanging issue when running the unittests because the FLASHINFER_CUDA_CALL macro expects the caller environment returns a cudaError_t which contradicts with the structure of the lambda functions inside invokeSelectiveStateUpdate of mamba kernel. This PR fixes the issue by adding another macro FLASHINFER_CUDA_CHECK.

🔍 Related Issues

#2366

🚀 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

Summary by CodeRabbit

• Refactor
  • Improved CUDA error handling mechanisms to provide clearer error messages and enhanced diagnostics during kernel configuration and execution.

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

[yzh119]

/bot run

[coderabbitai]

📝 Walkthrough

Walkthrough

A new CUDA error-checking macro FLASHINFER_CUDA_CHECK is introduced in the utilities header, which validates CUDA function results and formats error messages with context. Two existing error-check invocations in the selective state update kernel are updated to use this new macro.

Changes

Cohort / File(s)	Summary
CUDA Error Handling Macro Definition include/flashinfer/utils.cuh	Adds FLASHINFER_CUDA_CHECK(func) macro that executes a CUDA function and validates the result equals cudaSuccess, emitting formatted error messages with CUDA error string, code, file, line, and function text.
Macro Usage in Selective State Update include/flashinfer/mamba/selective_state_update.cuh	Updates two instances of cudaFuncSetAttribute calls to use the new FLASHINFER_CUDA_CHECK macro instead of FLASHINFER_CUDA_CALL for consistent error handling.

Estimated code review effort

🎯 2 (Simple) | ⏱️ ~10 minutes

Poem

🐰 A macro hops in to check CUDA's way,
Validating success throughout the day,
Error messages clear in the morning light,
Two kernels now check in a unified might!
Better error handling, what a delight!

🚥 Pre-merge checks | ✅ 3

✅ Passed checks (3 passed)

Check name	Status	Explanation
Title check	✅ Passed	The title clearly and specifically describes the main change: fixing a hotfix related to PR 2366 affecting the mamba kernel.
Description check	✅ Passed	The description follows the template structure with a clear explanation of the issue, related issues linked, and pre-commit/test checklist items addressed.
Docstring Coverage	✅ Passed	No functions found in the changed files to evaluate docstring coverage. Skipping docstring coverage check.

_{✏️ Tip: You can configure your own custom pre-merge checks in the settings.}

---

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.}

[flashinfer-bot]

GitLab MR !249 has been created, and the CI pipeline #42094754 is currently running. I'll report back once the pipeline job completes.

[flashinfer-bot]

[FAILED] Pipeline #42094754: 10/20 passed

[yzh119]

Failed UTs are not relevant, should be ready to merge.