[shantipriya-amd]

Proposed changes

ISSUE ID: ROCm/aiter#6142

Upstreams two commits originally developed locally alongside ROCm/aiter#5708 (relu2/Squared-ReLU MoE support for CK-Tile fused MoE GEMM). They were applied directly to a local CK submodule checkout rather than submitted upstream, which caused aiter's CK submodule pin to reference an unfetchable commit — see ROCm/aiter#6142 (revert) for full context.

1. Add MoeRelu2 activation functor for gate-only MoE flatmm kernel — adds the MoeRelu2 functor to projects/composablekernel/include/ck_tile/ops/flatmm/kernel/moe_flatmm_kernel.hpp, needed for non-gated (Squared ReLU) MoE activations used by NemotronH.
2. Fix half_t intrinsics (neg/sin/exp/log) via float round-trip for gfx950 — fixes projects/composablekernel/include/ck/utility/math_v2.hpp.

Once merged, aiter's CK submodule pin will be updated to reference this commit, and the relu2 MoE GEMM support PR will be resubmitted against aiter (see ROCm/aiter#5708 for the original implementation and test results, which remain valid).

Note on test coverage: these two commits were already exercised as part of ROCm/aiter#5708's test plan (op-level and kernel-level tests on real gfx950 hardware, including production-scale E=512/topk=22 sweeps — see that PR for full results). No new CK-level unit test is included in this PR since the functor is only exercised through aiter's op_tests, not CK's own test suite; happy to add a CK-side test if the review team wants one here specifically.

Checklist

• I have added tests relevant to the introduced functionality, and the unit tests are passing locally (carried over from original local testing in [HIP] [CK] Add relu2 (Squared ReLU) activation support to CK-Tile fused MoE GEMM aiter#5708 — op-level and kernel-level tests passed on real gfx950 hardware)
• I have added inline documentation
• I have run clang-format on all changed files

[therock-pr-bot]

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	⚠️ Warning	Error: Source/code files changed without an accompanying unit test. Expected: add at least one test file named like test_<name>.py / test_<name>.cpp (or <name>_test.*). Current: code file(s) changed: projects/composablekernel/include/ck/utility/math_v2.hpp, projects/composablekernel/include/ck_tile/ops/flatmm/kernel/moe_flatmm_kernel.hpp; no test file found
🔎 pre-commit	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—
🤖 therock-pr-bot	✅ Pass	—

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

[therock-pr-bot]

🎉 All checks passed! This PR is ready for review.

[JiaLuo-CAN]

Thanks for the PR. Could you be more specific about what issue requires the fix "Fix half_t intrinsics (neg/sin/exp/log) via float round-trip for gfx950 — fixes projects/composablekernel/include/ck/utility/math_v2.hpp."?

[shantipriya-amd]

@JiaLuo-CAN Thanks for your review. These were added as part of the changes to include the new ReLU2 activation function in AITER. These half_t specializations (neg, sin, ceil, floor, exp, log) call HIP's native half-precision intrinsics (__hneg, hsin, hceil, ...) directly on __half. On gfx950, these intrinsics don't compile/behave correctly, producing incorrect results. The fix upcasts to float, applies the float32 intrinsic, and downcasts back to half_t, which is portable across targets including gfx950.