Skip to content

sycl: fix compiler warnings - missing prototypes, dead code, extra semicolons - #30271

Merged
Titaniumtown merged 7 commits into
ggml-org:masterfrom
mctylr-gh:sycl-cmake
Oct 20, 2026
Merged

Titaniumtown merged 7 commits into
ggml-org:masterfrom
mctylr-gh:sycl-cmake

Conversation

@mctylr-gh

@mctylr-gh mctylr-gh commented Oct 10, 2026 •

Copy link
Copy Markdown
Contributor

Overview

Primary: The current Intel oneAPI DPC++ and I believe intel/llvm (sycl branch) identify as IntelLLVM not as Clang in CMake's CMAKE_COMPILER_ID, which means that the default warnings are not configured for some/many SYCL builds.

Secondary: Small, behaviour-preserving cleanups in the SYCL backend, found while building with the Intel oneAPI DPC++ compiler (IntelLLVM), these are pre-existing warnings, once the cmake build includes the usual warning flags for IntelLLVM:

  • mark 8 file-local functions static (only referenced in their own file,
    no prior declaration; -Wmissing-prototypes)
  • include the existing declaration header (gla.hpp) in gla.cpp so the
    cross-TU prototype of ggml_sycl_op_gated_linear_attn is visible at the definition
  • remove unreachable break/return after noreturn GGML_ABORT / GGML_ASSERT(false)
  • remove extra ';' after function / member definitions No behaviour change.

Verified: the SYCL backend recompiles clean under -Wall -Wextra -Wpedantic -Wmissing-prototypes -Wextra-semi. No new warnings introduced.

Issue: #30269: Feature Request: SYCL build: file-local functions missing static, unreachable code after abort, extra semicolons

Additional information

  • The dead break/return removals are the most likely to draw a preference
    either way; they are separate commit if maintainers want to keep the defensive
    break/return after the abort.
  • the fattn-common.hpp used a deprecated parallel_for(properties) overload:
warning: 'parallel_for(sycl::nd_range<3, sycl::default_accessor>,
 sycl::ext::oneapi::experimental::properties<...grf_size...>,
 <lambda>)' is deprecated: Use a launch_config object or a kernel functor
 with a get(properties_tag) member function to specify kernel properties

Requirements

  • I have read and agree with the contributing guidelines: Yes
  • AI usage disclosure: Yes an AI coding assistant identified the
    warnings and drafted the minimal edits from the compiler output; I
    reviewed each change, understand it, and take responsibility for all of
    it.

@mctylr-gh

mctylr-gh commented Oct 10, 2026 •

Copy link
Copy Markdown
Contributor Author

It looks like a based this branch on my yet unmerged (PR 29889) rather than off of the master branch.

This branch should be 3549e9b .. d07cc1f

Rebased to correct this. Hopefully keeping it easier to merge, and the other PR needs reviewing.

@github-actions github-actions Bot added build Compilation issues ggml changes relating to the ggml tensor library for machine learning SYCL https://en.wikipedia.org/wiki/SYCL - GPU programming language labels Oct 10, 2026
@mctylr-gh
mctylr-gh force-pushed the sycl-cmake branch 2 times, most recently from a9153b2 to 62fd9a0 Compare October 10, 2026 09:14
Comment thread ggml/src/ggml-sycl/fattn-common.hpp
The Intel fork of LLVM for SYCL support, and Intel oneAPI/DPC++ identifies
itself as IntelLLVM, which cmake recongises. So treat it much
like LLVM.

Tell IntelLLVM to not disable optimizations when debugging
information is includes, mostly of use for build type RelWithDebInfo,
and Debug. I don't believe it is supported by mainstream versions of
LLVM/Clang/Clang++.

Treat oneAPI and LevelZero headers as system headers (suppress warnings).
- mark file-local functions `static` (only referenced in their own TU,
  fixes `-Wmissing-prototypes`): `ggml_sycl_op_swiglu_oai`,
  `ggml_sycl_op_upscale`, `get_zes_init_res`, `initialize_sycl_begining`,
  `is_bmg_g31_arch`, `sycl_get_mem_info`, `query_free_memory_by_ze`,
  `get_memory_size_by_sycl_api`
- include the declaration header `gla.hpp` in `gla.cpp` so the prototype of
  `ggml_sycl_op_gated_linear_attn` is visible at its definition (called
  from `ggml-sycl.cpp`)
- remove extra ';' after function / member definitions: `upscale` (x3),
  `vecdotq`, `dpct/helper.hpp`
- remove break / return unreachable after the `noreturn GGML_ABORT` /
  `GGML_ASSERT(false)` (`ggml_abort` is `GGML_NORETURN`): `concat`,
  `count-equal`, `fattn-tile`, `getrows`, `topk-moe`, `ggml-sycl` (x2),
  `convert` (x3)

If the `GGML_ASSERT` definition was modified or redefined to be empty,
similar to how C `assert()` acts when `NDEBUG` is defined, these changes
could generate errors.
- remove extra ';' after function / member definitions:
  `dpct/helper.hpp`
…load

`parallel_for(range, properties, kernel)`, used to advertise the Intel
`grf_size_automatic` property for the flash-attention kernels, is
deprecated in SYCL 2020. The recommended replacements are a
`launch_config` object or a kernel *functor* with a `get(properties_tag)`
member. `launch_config` cannot carry `grf_size` (it static_asserts against
properties that have a compile-time kernel effect), so the kernel is
wrapped in a functor that advertises the property through the
properties_tag getter.

- add `grf_size_kernel<KernelFunc, warp_size>`: a kernel functor that
  forwards to the kernel lambda and returns `grf_size_automatic` from a
  `const get(properties_tag)` member
- `lauch_kernel`: use `grf_size_kernel` in the `use_large_grf` branch instead
  of the deprecated overload; the `use_large_grf = false` branch is unchanged
- the `reqd_sub_group_size` attribute is also placed on the functor operator()
  it is written at the start of the declaration because icpx rejects it
  after the const qualifier

Assisted-by: pi (unsloth/Qwen3.8-27B-GGUF:Q4_K_M via llama.cpp)
The kernel carried both the `reqd_sub_group_size` attribute and a kernel
property (`grf_size_automatic`), which icpx flags with
`-Wignored-attributes`. The oneAPI kernel-properties framework already
provides `sub_group_size`, so move the sub-group size from the attribute
into the same get(properties_tag) object that carries
`grf_size_automatic`.

- rename `grf_size_kernel` to `props_kernel<KernelFunc, Size, UseGrf>;` the
  sub-group size is now a property (`sub_group_size<Size>`) returned from
  the   `const get(properties_tag)` member, and `grf_size_automatic` is
  added only when `UseGrf` is true
- drop the `reqd_sub_group_size` attribute from the kernel lambda and the
  functor operator(); no attribute remains on the kernel
- `lauch_kernel`: collapse the `use_large_grf` branch into a single
  `parallel_for(rng, props_kernel<..., use_large_grf>{kernel})`

No behaviour change: on-device (bmg-g31) the property applies the same
sub-group size as the attribute (16), versus 32 by default.

Assisted-by: pi (unsloth/Qwen3.8-27B-GGUF:Q4_K_M via llama.cpp)
Correct incorrect, potentially misleading arguments in cmake
`find_path()`.

This revert my change in 0bb287db0 to the original behaviour.

@arthw arthw 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.

It's good job!

Thank you!

@Titaniumtown Titaniumtown changed the title Sycl build: fix compiler warnings - missing prototypes, dead code, extra semicolons sycl: fix compiler warnings - missing prototypes, dead code, extra semicolons Oct 20, 2026
@Titaniumtown
Titaniumtown merged commit dd002b9 into ggml-org:master Oct 20, 2026
27 of 35 checks passed
@Titaniumtown

Copy link
Copy Markdown
Contributor

Thank you @mctylr-gh!

@mctylr-gh
mctylr-gh deleted the sycl-cmake branch October 20, 2026 05:50
edwardyoon pushed a commit to edwardyoon/focus-llama that referenced this pull request Oct 24, 2026
auRose94 pushed a commit to auRose94/llama.cpp that referenced this pull request Oct 28, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

build Compilation issues ggml changes relating to the ggml tensor library for machine learning SYCL https://en.wikipedia.org/wiki/SYCL - GPU programming language

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants