[SYCL][CUDA] Add bfloat16 and marray support to ldg builtin#22732
Open
zjin-lcf wants to merge 2 commits into
Open
[SYCL][CUDA] Add bfloat16 and marray support to ldg builtin#22732zjin-lcf wants to merge 2 commits into
zjin-lcf wants to merge 2 commits into
Conversation
The CUDA read-only-cache `ldg` builtin previously rejected `bfloat16` and `marray` arguments, producing "no matching function for call to 'ldg'" (intel#15418). Add support for scalar `bfloat16`, `vec<bfloat16, 2/3/4>`, and generic `marray<E, 2/3/4>` for all supported element types. `marray` shares the `E[N]` storage layout, so it delegates to the existing `vec<E, N>` loads (which read exactly N elements, keeping the padded N=3 case safe) and thus gets the same `ld.global.nc` read-only-cache / vectorized codegen. Update the extension specification and check_device_code test accordingly. Fixes intel#15418 Co-authored-by: Cursor <cursoragent@cursor.com>
Address code_formatter CI failure on PR intel#22732 by reflowing the bfloat16 vec branches and the bfloat16 ldg template instantiation to match clang-format. Co-authored-by: Cursor <cursoragent@cursor.com>
bratpiorka
approved these changes
Jul 24, 2026
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
Fixes #15418.
The CUDA read-only-cache
ldgbuiltin previously rejectedbfloat16andmarrayarguments, producing "no matching function for call to 'ldg'".bfloat16andvec<bfloat16, 2/3/4>support.marray<E, 2/3/4>support for all supported element types.marrayshares theE[N]storage layout, so it delegates to the existingvec<E, N>loads (which read exactly N elements, keeping the padded N=3 case safe) and gets the sameld.global.ncread-only-cache / vectorized codegen.sycl_ext_oneapi_cuda_tex_cache_readextension spec (documentbfloat16,marray, and the alignment requirement).sycl/test/check_device_code/cuda/ldg.cpp.Test plan
check_device_code/cuda/ldg.cppFileCheck passes (verified with an NVPTX-enabled DPC++ build).ldgcompiles formarrayof bfloat16/half/float/double/int/char.ld.global.ncfor bf16; thebit_castworkaround adds no extra instructions.addBiasResidualLayerNorm-syclcompiles and runs correctly on an H100 (checksums unchanged); bf16ldgis up to ~7-8% faster than plain loads.