[SYCL] Add CUDA/HIP hardware acceleration for dot_acc (dp4a)

[zjin-lcf]

Summary

Lower the sycl_ext_oneapi_dot_accumulate dot_acc helpers 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.

• NVPTX (CUDA): emit dp4a.{s32,u32}.{s32,u32} on sm_61+ (Pascal and later) for 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/RDNA4) for the mixed-sign cases and, where sdot4 is unavailable (RDNA3+ dropped v_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

• Move the scalar fallbacks into detail and drop the Us/Uu unions that leaked into the public sycl::ext::oneapi namespace.
• Read the packed words through int8_t/uint8_t views instead of a char-based union: fixes signedness on targets where plain char is unsigned (e.g. AArch64) and removes type-punning UB.
• vec<*8_t, 4> overloads now sycl::bit_cast directly to the packed 32-bit word.

Test plan

• New check_device_code codegen tests:
  • CUDA sm_90: dp4a.s32.s32, dp4a.u32.u32, dp4a.s32.u32, dp4a.u32.s32.
  • HIP gfx90a, gfx942, gfx1030 (sdot4/udot4) and gfx1100, gfx1200 (sudot4).
• E2E DotProduct int + vec tests (all four sign combinations) pass on:
  • NVIDIA H100 (sm_90)
  • AMD Instinct MI210 (gfx90a) and MI300A (gfx942)
• Host equivalence of the refactored fallbacks vs. reference: 0 mismatches over millions of random inputs.

[wenju-he]

could you also upstream this change to llvm-project?

[zjin-lcf]

Done - upstreamed as llvm/llvm-project#213691.

I have also dropped that commit from this PR: it was only here because the branch was accidentally created off my other topic branch, and it is unrelated to dp4a. It is still tracked separately in #22750, so this PR is now purely the dp4a change.

While preparing the upstream patch I found the __NVPTX__ exclusion came from llvm/llvm-project#146814, where it was added to silence the large atomic operation may incur significant performance penalty warning on 32-bit targets. Since NVPTX sets MaxAtomicInlineWidth to the target pointer width, nvptx64 never emits that warning, and libclc no longer builds a 32-bit nvptx target, so the guard is obsolete. The upstream PR therefore drops __NVPTX__ from clc_atomic_def.inc as well, which is the same change intel/llvm carries locally as 5f67a6c (marked "This change will be upstreamed").

[bratpiorka]

@zjin-lcf please fix the code formatting

[zjin-lcf]

Thanks for the review, both comments should be addressed now:

• @bratpiorka formatting is fixed. git clang-format is now clean on the whole change, and no line exceeds 80 columns. The main fixes were the line continuations in the #if defined(...) guards, the wrapping of the return expressions in the scalar fallbacks, and an over-long comment block in dot_product.hpp which I rewrapped by hand so clang-format would not leave a dangling word. I also rewrapped the // RUN: lines in the two new codegen tests, which were 81 columns for the longer gfx* names.
• @wenju-he the libclc change is upstreamed as [libclc] Enable 64-bit atomics for nvptx64 llvm/llvm-project#213691 and is no longer part of this PR (details in the thread above).

The branch was recreated so the PR is now a single dp4a commit. Apart from the formatting, the content is unchanged: ignoring whitespace, dot_product.hpp is identical to the previous revision, and the RUN commands lit assembles from the rewrapped lines are byte-identical. I re-ran the HIP codegen test for all five architectures (gfx90a, gfx942, gfx1030, gfx1100, gfx1200) and they still pass; the CUDA one I could not re-run locally since my build has no NVPTX backend registered, but that file only changed by moving %s onto the next RUN line.