[clr] Allow 32-bit warp-sync masks on wave32 architectures

[iassiour]

Motivation

Allow 32-bit warp-sync masks on wave32 AMD targets. Code such as:
__shfl_sync(0xffffffff, value, 0, width)
currently fails to compile because HIP requires every explicit warp-sync mask to be a 64-bit integer, even on wave32 architectures where a 32-bit mask is the natural width.

Technical Details

Replaces the duplicated sizeof(MaskT) == 8 checks in amd_warp_sync_functions.h with a shared mask-type predicate.
Allows 32-bit integer masks for wave32-only targets covered by this change while keeping the existing 64-bit requirement for other device targets.
Leaves the existing runtime behavior intact: wave32 masks still ignore the upper half via __hip_adjust_mask_for_wave32.
Adds Unit_hipShflSync_32BitMask to validate that __shfl_sync accepts a 32-bit mask and broadcasts lane 0 correctly.

JIRA ID

Resolves ROCM-27352.

Test Plan

Compile raw 32-bit-mask repro for gfx1250/gfx942/gfx950.
Compile the new guarded unit test for gfx1250/gfx942/gfx950.
Run existing CI.

Test Result

Raw 32-bit mask compiles for wave32 targets covered by this change.
Raw 32-bit mask remains rejected on gfx942/gfx950 wave64 targets.
New Unit_hipShflSync_32BitMask syntax-checks cleanly for gfx1250/gfx942/gfx950.

Submission Checklist

• Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

[therock-pr-bot]

❌ PR Check — Action Required

Check	Status	Details
🌿 Branch Name	✅ Pass	—
📝 PR Title/Description	❌ Fail	Error: Title does not follow Conventional Commits style. Expected: start with a valid type (feat, fix, docs, …). Desired format: type(optional-scope): short description ─── Error: PR description must reference a JIRA ID, ISSUE ID, or a GitHub closing keyword. Expected: include a JIRA ID / ISSUE ID line (separator : or -, or omitted; value may be a JIRA key, a number with/without #, or a link), OR a closing keyword + issue reference. Accepted examples: • JIRA ID : TESTAUTO-6039 • JIRA ID - #330 • JIRA ID #330 • ISSUE ID : TESTUTO-3334 • ISSUE ID #3334 • ISSUE ID - TESTAUTO-3433 • ISSUE ID : https://github.com/<org_name>/<repo_name>/issues/1234 • Closes #10 • Fixes octo-org/octo-repo#100 • Resolves: #123 • #123 • https://github.com/<org_name>/<repo_name>/issues/123 Current: no valid JIRA/ISSUE/closing-keyword reference found
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	❌ Fail	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/clr/hipamd/include/hip/amd_detail/amd_warp_sync_functions.h, projects/hip-tests/catch/unit/warp/hipShflSyncTests.cc; no test file found
🔎 pre-commit	⏳ Pending	⏳ Still running…
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—

⚠️ 2 policy check(s) failed. Please address the issues above before this PR can be Reviewed.

🚫 Please fix the failed policies

• ❌ PR Title/Description
• ❌ Unit Test

The Not ready to Review label was added to this PR. Once all policies pass, the label is removed automatically.

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

[therock-pr-bot]

🚫 Please fix the failed policies before requesting reviews.

The following policy checks failed:

• ❌ PR Title/Description
• ❌ Unit Test

The Not ready to Review label has been added to this PR.
Once all policies pass, the label will be removed automatically.

[Copilot]

Pull request overview

This PR updates HIP’s warp-sync builtins to permit a 32-bit lane mask on wave32-class targets (instead of requiring a 64-bit integer mask everywhere), and adds a unit test intended to validate __shfl_sync behavior with a 32-bit mask.

Changes:

• Relax MaskT static-assert validation in amd_warp_sync_functions.h to allow 32-bit masks on specific targets.
• Add a new HIP test kernel + Catch2 test case to validate __shfl_sync with a 32-bit mask.
• Register the new test in the warp unit test YAML config.

Reviewed changes

Copilot reviewed 3 out of 3 changed files in this pull request and generated 3 comments.

File	Description
projects/hip-tests/catch/unit/warp/hipShflSyncTests.cc	Adds a new kernel and test case for __shfl_sync mask-type acceptance and broadcast correctness.
projects/hip-tests/catch/config/configs/unit/warp.yaml	Registers the new unit test in the warp test suite.
projects/clr/hipamd/include/hip/amd_detail/amd_warp_sync_functions.h	Adjusts mask type validation to permit 32-bit masks under a target-specific compile-time guard.

---

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

[g-h-c]

If we do this, for consistency, all the __reduce_sync operations should also allow a 32-bit mask

[ssahasra]

Keeping the mask argument independent of wavesize was a deliberate decision. The mask is always 64-bit so that a single HIP source will compile correctly to any target. Perhaps we need a deeper understanding of the need for a 32-bit mask. Is it just convenience? Or a real issue? Even if it is just convenience, how important is it?

[iassiour]

If we do this, for consistency, all the __reduce_sync operations should also allow a 32-bit mask

I think this is covered by the changes in reduce_op_sync https://github.com/ROCm/rocm-systems/pull/8385/changes#diff-c885872acc0fe466252f4a87941bfc669032ba494f1107cbd020e94e2847ac4bR327

But modified the test to cover the case as well

[iassiour]

closing based on the comments on https://amd-hub.atlassian.net/browse/ROCM-27352