Skip to content

[rocPRIM] Guard __builtin_amdgcn_mov_dpp with __builtin_amdgcn_is_invocable - #8418

Closed
MrSidims wants to merge 1 commit into
developfrom
guard-dpp-with-builtin
Closed

MrSidims wants to merge 1 commit into
developfrom
guard-dpp-with-builtin

Conversation

@MrSidims

Copy link
Copy Markdown

For the generic 'amdgcnspirv' target the 'dpp' feature is unknown until the SPIR-V is JIT-compiled for a concrete GPU by comgr, so clang rejects the bare builtin with "cannot be invoked in the current context, as it requires the 'dpp' feature(s)". Guard it with a resolved at JIT 'if(__builtin_amdgcn_is_invocable(...))'.

More technical details:
llvm/llvm-project#134016

Other similar changes in ROCm stack:
ROCm/rocm-systems#5581

Test Result

Without the patch all related to rocPRIM rocm-examples compiled for amdgcnspirv fail with the
error mentioned above.
With the patch all rocm-examples can be compiled and pass. Same story for HeCBench samples.

…ocable

For the generic 'amdgcnspirv' target the 'dpp' feature is unknown until the
SPIR-V is JIT-compiled for a concrete GPU by comgr, so clang rejects the
bare builtin with "cannot be invoked in the current context, as it requires the
'dpp' feature(s)". Guard it with a resolved at JIT
'if(__builtin_amdgcn_is_invocable(...))'.
@MrSidims

Copy link
Copy Markdown
Author

Closing in favor of #8102

@MrSidims MrSidims closed this Jun 16, 2026
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.

1 participant