Skip to content

AIRUNTIME-2821 - fix two misuses of __builtin_amdgcn_is_invocable() - #13281

Merged
g-h-c merged 2 commits into
developfrom
users/g-h-c/AIRUNTIME-2821_fix_invocable_guards
Oct 14, 2026
Merged

g-h-c merged 2 commits into
developfrom
users/g-h-c/AIRUNTIME-2821_fix_invocable_guards

Conversation

@g-h-c

@g-h-c g-h-c commented Oct 12, 2026

Copy link
Copy Markdown
Contributor

Motivation

The test Unit_hiprtc_spirv_compilation fails in some architectures with:

/tmp/comgr-3248377-0-f4c112/include/hiprtc_runtime.h:9028:7: error: '__builtin_amdgcn_query_shared_rank' cannot be invoked in the current context, as it requires the 'gfx13-insts,wavefrontsize32' feature(s)
 9028 |       __builtin_amdgcn_query_shared_rank((__attribute__((address_space(11))) const void*)in));
      |       ^
1 error generated when compiling for spirv64-amd-amdhsa.
comgr:  ReturnStatus: AMD_COMGR_STATUS_ERROR

Technical Details

The intrinsics must be guarded using __builtin_is_invocable() but instead there are two instances where the code tries to exit early if an intrinsic is invocable via __builtin_trap(), as opposed to be guarded with if()

Issue Tracking

AIRUNTIME-2821

Test Plan

Re-run with:

Unit_hiprtc_spirv_compilation
Unit_hiprtc_spirv_linker

Test Result

Both test passed OK now. Before Unit_hiprtc_spirv_compilation was failing for me.

Submission Checklist

@g-h-c
g-h-c requested review from a team as code owners October 12, 2026 15:59
Copilot AI balanced review requested due to automatic review settings October 12, 2026 15:59
@g-h-c
g-h-c requested a review from cjatin October 12, 2026 16:00
@therock-pr-bot

Copy link
Copy Markdown

❌ PR Check — Action Required

Check Status Details
📝 PR Description ❌ Fail 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
• JIRA ID (on separate line)
ROCM-25757
• ISSUE ID : TESTUTO-3334
• ISSUE ID #3334
• ISSUE ID - TESTAUTO-3433
• ISSUE ID (on separate line)
AIRUNTIME-2352
• 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 ⚠️ 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/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h; no test file found
🔎 pre-commit ⏳ Pending ⏳ Still running…
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —

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

🚫 Please fix the failed policies

  • ❌ PR Description

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.

🙋 Wish to Override Policy?

@therock-pr-bot

Copy link
Copy Markdown

🚫 Please fix the failed policies before requesting reviews.

The following policy checks failed:

  • ❌ PR Description

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

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟢 Approval recommended

The functional fix is sound; only minor formatting corrections remain.

1 open finding
What changed in this PR

Fixes architecture-specific HIPRTC compilation failures by correctly guarding AMDGCN intrinsics.

Changes:

  • Guards mbcnt and shared-rank intrinsics with invocability checks.
  • Preserves trap-based fallbacks when intrinsics are unavailable.
File Description
projects/​clr/​hipamd/​include/​hip/​amd_detail/​hip_cooperative_groups_helper.h Corrects intrinsic guards and fallback control flow.

🧠 Review effort: Balanced


💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

@mangupta

Copy link
Copy Markdown
Contributor

🤖 AI-assisted pre-review (for @mangupta)

Re-reviewed the updated head (1309a4344413) — no new findings at or above the posting threshold.

@g-h-c
g-h-c requested a review from iassiour October 13, 2026 09:40
@g-h-c
g-h-c force-pushed the users/g-h-c/AIRUNTIME-2821_fix_invocable_guards branch from 1309a43 to cfced37 Compare October 13, 2026 15:10
@g-h-c
g-h-c merged commit bfffd9e into develop Oct 14, 2026
168 of 169 checks passed
@g-h-c
g-h-c deleted the users/g-h-c/AIRUNTIME-2821_fix_invocable_guards branch October 14, 2026 08:24
@systems-assistant

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

5 participants