fix: negative zero by type trait --> binary value

[yyihuang]

📌 Description

We are going to fix the has_neg_zero.

🔍 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 @yyihuang, 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 addresses a potential issue with detecting negative zero in floating-point values within the CUDA all-reduce implementation by switching from a type-trait based comparison to a direct binary representation check. It also includes some adjustments to the associated test suite to refine the correctness checks.

Highlights

• Negative Zero Detection: Implemented a new device function is_negative_zero with specializations for floating-point types (float, double, __half, __nv_bfloat16) that checks for negative zero by inspecting the binary representation rather than using standard equality comparison, which can be unreliable for floating-point zeros.
• Refactor has_neg_zero and remove_neg_zero: Updated the existing has_neg_zero and remove_neg_zero functions to utilize the new is_negative_zero helper function for accurate detection and handling of negative zero within vector operations.
• CUDA Synchronization: Enabled cudaDeviceSynchronize() in the lamportInitializeAll function, which was previously commented out.
• Test Suite Adjustments: Modified the correctness test for trtllm_allreduce by reducing the set of tested hidden sizes, increasing the number of test loops, changing test failure handling from printing messages to using assert, and slightly increasing the tolerance for non-float16 dtypes.

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 refactors the detection of negative zero in CUDA kernels by switching to a more robust bitwise comparison method, which is a good improvement. The core C++ changes implement this well. Key areas for attention include a critical issue with an unreachable cudaDeviceSynchronize call that needs to be repositioned, and a potentially unused variable that could be cleaned up.
In the Python tests, several parameters have been adjusted (test loop count, hidden sizes, and numerical tolerance). The rationale for the changes to hidden sizes and tolerance would be helpful to understand their impact on test coverage and reliability. The simplification of assertion logic in tests using assert directly is a positive change.