[SYCL] Add CUDA/HIP hardware acceleration for dot_acc (dp4a)#22756
Open
zjin-lcf wants to merge 2 commits into
Open
[SYCL] Add CUDA/HIP hardware acceleration for dot_acc (dp4a)#22756zjin-lcf wants to merge 2 commits into
zjin-lcf wants to merge 2 commits into
Conversation
The compare_exchange definition guard grouped __NVPTX__ with __SPIR32__ and only emitted atomics for 32-bit types, so 64-bit integer __clc_atomic_compare_exchange was declared (via atomic_decl.inc) but never defined for CUDA, causing "Unresolved extern function '_Z29__clc_atomic_compare_exchange...'" (ptxas fatal) for 64-bit CAS. Match the guard used by the sibling clc_atomic_def.inc (restrict only __SPIR32__ to 32-bit) so both CUDA (NVPTX) and HIP (AMDGPU) get 64-bit compare_exchange definitions. Co-authored-by: Cursor <cursoragent@cursor.com>
The `sycl_ext_oneapi_dot_accumulate` `dot_acc` helpers previously always
used a scalar implementation. Lower them to the native dp4a instruction
on hardware that supports it, keeping the scalar path as a portable
fallback (the numerical result is identical):
- NVPTX (CUDA): emit the `dp4a.{s32,u32}.{s32,u32}` instruction on
sm_61+ (Pascal and later), covering all four operand-signedness
combinations.
- AMDGCN (HIP): use `__builtin_amdgcn_sdot4` / `udot4` (v_dot4, gfx906+/
CDNA/RDNA2) and `__builtin_amdgcn_sudot4` (v_dot4_i32_iu8, gfx11/RDNA3+
and gfx12) for the mixed- and, where sdot4 is unavailable, same-sign
cases. Each path is guarded by `__builtin_amdgcn_is_invocable` so
architectures without a given instruction fall back to scalar.
The scalar fallbacks are moved into `detail` (dropping the `Us`/`Uu`
unions that leaked into the public namespace) and now read the packed
words through `int8_t`/`uint8_t` views, which fixes signedness on targets
where plain `char` is unsigned (e.g. AArch64) and avoids type-punning UB.
The `vec` overloads bit-cast directly to the packed 32-bit word.
Add check_device_code tests for the CUDA (sm_90) and HIP (gfx90a, gfx942,
gfx1030, gfx1100, gfx1200) codegen, and update the existing E2E tests'
comments to reflect that they now validate the hardware paths too.
Co-authored-by: Cursor <cursoragent@cursor.com>
wenju-he
reviewed
Jul 25, 2026
Contributor
There was a problem hiding this comment.
could you also upstream this change to llvm-project?
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Lower the
sycl_ext_oneapi_dot_accumulatedot_acchelpers to the native dp4a instruction on hardware that supports it, keeping the existing scalar implementation as a portable fallback. The numerical result is identical on every path.dp4a.{s32,u32}.{s32,u32}on sm_61+ (Pascal and later) for all four operand-signedness combinations.__builtin_amdgcn_sdot4/udot4(v_dot4, gfx906+/CDNA/RDNA2) and__builtin_amdgcn_sudot4(v_dot4_i32_iu8, gfx11/RDNA3+ and gfx12/RDNA4) for the mixed-sign cases and, wheresdot4is unavailable (RDNA3+ droppedv_dot4_i32_i8), the same-sign signed case. Every path is guarded by__builtin_amdgcn_is_invocable, so architectures lacking a given instruction fall back to scalar.Cleanups
detailand drop theUs/Uuunions that leaked into the publicsycl::ext::oneapinamespace.int8_t/uint8_tviews instead of achar-based union: fixes signedness on targets where plaincharis unsigned (e.g. AArch64) and removes type-punning UB.vec<*8_t, 4>overloads nowsycl::bit_castdirectly to the packed 32-bit word.Test plan
check_device_codecodegen tests:sm_90:dp4a.s32.s32,dp4a.u32.u32,dp4a.s32.u32,dp4a.u32.s32.gfx90a,gfx942,gfx1030(sdot4/udot4) andgfx1100,gfx1200(sudot4).DotProductint + vec tests (all four sign combinations) pass on: