Repository navigation
OIDN (Blender) build broken in nightly #21893
Description
Describe the bug
We are seeing this build failure in OIDN which is a dependency for Blender
https://github.com/intel/llvm/actions/runs/25033592382/job/73322785903
D:/github/_work/llvm/llvm/oidn-src\core\kernel.h:193:28: error: no member named 'store' in 'sycl::sub_group'
193 | item.get_sub_group().store(dst, x);
| ~~~~~~~~~~~~~~~~~~~~ ^
D:/github/_work/llvm/llvm/oidn-src/devices/sycl\../gpu/gpu_input_process.h:144:14: note: in instantiation of function template specialization 'oidn::WorkGroupItem<2>::subgroupStore<sycl::detail::half_impl::half>' requested here
144 | it.subgroupStore(dstPtr + i * dstPaddedC, DstT(dstBlock));
| ^
D:/github/_work/llvm/llvm/oidn-src/devices/sycl\sycl_engine.h:94:80: note: in instantiation of member function 'oidn::GPUInputProcessKernel<sycl::detail::half_impl::half, oidn::TensorLayout::Chw16c, 16>::operator()' requested here
94 | [=](sycl::nd_item<N> it) [[sycl::reqd_sub_group_size(subgroupSize)]] { kernel(it); });
| ^
D:/github/_work/llvm/llvm/oidn-src/devices/sycl\../gpu/gpu_input_process.h:257:24: note: in instantiation of function template specialization 'oidn::SYCLEngine::submitKernel<16, 2, oidn::GPUInputProcessKernel<sycl::detail::half_impl::half, oidn::TensorLayout::Chw16c, 16>>' requested here
257 | engine->template submitKernel<subgroupSize>(numGroups, groupSize, kernel);
| ^
D:/github/_work/llvm/llvm/oidn-src/devices/sycl\../gpu/gpu_input_process.h:211:16: note: in instantiation of function template specialization 'oidn::GPUInputProcess<oidn::SYCLEngine, sycl::detail::half_impl::half, oidn::TensorLayout::Chw16c, 16>::submitImpl<3>' requested here
211 | case 3: submitImpl<3>(); break;
| ^
D:/github/_work/llvm/llvm/oidn-src/devices/sycl\../gpu/gpu_input_process.h:185:5: note: in instantiation of member function 'oidn::GPUInputProcess<oidn::SYCLEngine, sycl::detail::half_impl::half, oidn::TensorLayout::Chw16c, 16>::submitKernels' requested here
185 | GPUInputProcess(EngineT* engine, const InputProcessDesc& desc)
| ^
D:/github/_work/llvm/llvm/oidn-src\core\ref.h:146:23: note: in instantiation of member function 'oidn::GPUInputProcess<oidn::SYCLEngine, sycl::detail::half_impl::half, oidn::TensorLayout::Chw16c, 16>::GPUInputProcess' requested here
146 | return Ref<T>(new T(std::forward<Args>(args)...));
| ^
D:/github/_work/llvm/llvm/oidn-src/devices/sycl/sycl_engine.cpp:101:12: note: in instantiation of function template specialization 'oidn::makeRef<oidn::GPUInputProcess<oidn::SYCLEngine, sycl::detail::half_impl::half, oidn::TensorLayout::Chw16c, 16>, oidn::SYCLEngine *, const oidn::InputProcessDesc &>' requested here
101 | return makeRef<GPUInputProcess<SYCLEngine, half, TensorLayout::Chw16c, 16>>(this, desc);
| ^
In file included from D:/github/_work/llvm/llvm/oidn-src/devices/sycl/sycl_engine.cpp:8:
D:/github/_work/llvm/llvm/oidn-src/devices/sycl\../gpu/gpu_input_process.h:144:14: error: SYCL kernel cannot call an undefined function without SYCL_EXTERNAL attribute
144 | it.subgroupStore(dstPtr + i * dstPaddedC, DstT(dstBlock));
| ^
D:/github/_work/llvm/llvm/oidn-src\core\kernel.h:191:29: note: 'subgroupStore<sycl::detail::half_impl::half>' declared here
191 | oidn_device_inline void subgroupStore(GlobalPtr<T> dst, const T& x) const
| ^
D:/github/_work/llvm/llvm/oidn-src/devices/sycl\../gpu/gpu_input_process.h:82:29: note: called by 'operator()'
82 | oidn_device_inline void operator ()(const oidn_private WorkGroupItem<2>& it) const
To reproduce
No response
Environment
- OS: [e.g Windows/Linux]
- Target device and vendor: [e.g. Intel GPU]
- DPC++ version: [e.g. commit hash or output of
clang++ --version] - Dependencies version: [e.g. the output of
sycl-ls --verbose]
Additional context
No response
Activity
sarnex commentedon Apr 28, 2026
Likely guility PR #21875
Since the API seems deprecated, maybe OIDN needs to be updated instead of a compiler change @atafra?
FYI @againull
atafra commentedon Apr 28, 2026
Do I understand correctly that a stable (non-experimental) API was removed in favor of an experimental one? That seems very strange. I assume the experimental API will also disappear again, requiring another change. If there is a new stable API the replaces this deprecated one, I'm happy to make the change, but until then, I would rather not spend time on this, since it would be temporary anyway.
sarnex commentedon Apr 28, 2026
I will let @againull comment on the substance of the change.
However we recently added Blender (and OIDN) testing in our CI which is why this was caught, so it's unclear to me how we can do that testing if we decide to not implement fixes somewhere (either in the compiler or the downstream code)
atafra commentedon Apr 28, 2026
My preference would be that no stable API should be removed from SYCL while there is no stable replacement for it. Otherwise, the developers are forced to constantly update their code, wasting a lot of time. IMO this is a very developer unfriendly practice. Either this API shouldn't have been removed yet, or there should be a non-experimental version of the new API. SYCL seems to keep around "experimental" APIs for years, which doesn't make much sense to me: if it's being used for such a long time, what's experimental about it? This just suggests a lack of commitment to support these APIs in the long term, without any stable version in sight.
againull commentedon Apr 28, 2026
@atafra
@atafra Sorry for the disruption. The old subgroup load/store API was deprecated about a year and a half ago in favor of the newer experimental extension sycl_ext_oneapi_group_load_store. We removed the deprecated API believing that users had migrated to the new extension. The newer extension cannot be promoted to “supported” today because it depends on properties, which are still experimental, but some work has begun to promote those dependencies (KhronosGroup/SYCL-Docs#980).
The switch to experimental extension sycl_ext_oneapi_group_load_store will allow you to get rid of deprecation warnings, but I understand your concern that it will require additional changes in the future.
I’ve prepared a revert PR - let me know whether you’d prefer that I proceed with the revert, or help you migrate to the experimental sycl_ext_oneapi_group_load_store.
atafra commentedon Apr 28, 2026
@againull I just realized that currently the only SYCL kernel which uses these deprecated functions will be removed in the upcoming OIDN version, because it will be replaced with an ESIMD kernel. So the problem will go away by itself.
But I'm still curious: what's different about this new extension apart from the syntax? What was the reason that a new API was needed? Thanks!
gmlueck commentedon Apr 29, 2026
The original API was added very early in our project before we had a defined process for adding extensions. Since then (and quite a while ago), we got feedback from users that all extensions should be named in some way to make it obvious that they are not part of the core SYCL specification. Therefore, we settled on a naming convention that extensions should either be defined in the
sycl::ext::<vendorstring>namespace (where<vendorstring>is some unique name for each vendor), or they should be named with the prefixext_<vendorstring>. DPC++ extensions typically useoneapifor the<vendorstring>.From our perspective, we'd like to remove the old extension for two reasons. First, it doesn't follow the naming convention, so application programmers may not realize they are using a DPC++ extension. Second, we want to clean up our code, so that we don't have two implementations for the same functionality.
If you are planning to remove the kernel using this extension anyway, are you no longer opposed to us removing it?
atafra commentedon Apr 29, 2026
Thanks, I understand the context now.
Since this won't cause an immediate issue for OIDN, I do not oppose removing the old API anymore. However, I'm still somewhat concerned about having an experimental extension as a replacement. What is the reason for keeping this new API as experimental? Do you expect further breaking changes to this API? Do you have a roadmap for moving this out of
experimental?gmlueck commentedon Apr 29, 2026
We do have a roadmap for formalizing experimental extensions. Our preference for most extensions is to make them into "KHR" extensions. These are extensions that the Khronos SYCL workgroup ratifies and publishes as an appendix in the core SYCL specification.
When we do this, we take this as an opportunity to learn from our experience using the experimental extension. Sometimes that means we change things in the API. For example, we are currently turning sycl_ext_oneapi_properties into a KHR extension. We decided that some of the features in the experimental extension were not being used by anyone, and these features were causing compilation to be slower for everyone. (Those particular features had a lot of C++ metaprogramming in the headers.) So, the KHR version of the extension will not contain those features. Bringing a KHR extension to the WG also exposes it to review by new people, and it is not uncommon for this to result in additional changes. In general, the process results in a better extension.
The sub-group load/store extension uses properties, so we would definitely want the non-experimental version to use the upcoming KHR properties, rather than the experimental properties. We'd probably also want to reconsider some of the features in that extension. The current (experimental) sycl_ext_oneapi_group_load_store had contributors from several people, and a lot of features got added. Maybe these are all important features, but we should spend some time to make sure they are really important and not just bloat. There's also a weird dependency on the group-sort extension, which we would want to investigate. So, you can see that moving the extension out of "experimental" status is more than just simply blessing it and moving the extension spec to a new directory.
Once an extension becomes a KHR, we deprecate the experimental one and ask people to transition to the KHR. After a deprecation period, we would remove it.
atafra commentedon Apr 29, 2026
Thanks. I understand the reasoning behind all this but unfortunately this mechanism makes shipping products using SYCL extremely challenging. The extension says that "Shipping software products should not rely on APIs defined in this specification.". If developers would follow this advice, it would be basically impossible to ship any performant or practical product based on SYCL for years. The reality is that SYCL code which aims to achieve good/usable performance is forced to use a multitude of such experimental extensions, while each of these get independently changed, deprecated and replaced at various different times, all which developers need to keep track of. This not only increases development time considerably but inevitably makes the SYCL code highly non-portable, with just only a couple of compiler versions being able to build a particular version of the codebase.
sarnex commentedon May 1, 2026
@atafra Is there an OIDN patch I can apply that moves to the ESIMD kernel so can we fix the CI here? Thanks
[CI] Use devel branch commit for OIDN when building Blender (#21926)
atafra commentedon May 8, 2026
@sarnex Please apply the following patch to OIDN: subgroup_store.patch
Or you can also switch to the latest OIDN https://github.com/RenderKit/oidn/tree/devel branch until OIDN 2.5.0 gets released.
sarnex commentedon May 8, 2026
Thanks, I've switched to the devel branch for now, it's working fine
sarnex commentedon May 14, 2026
@atafra Do you need anything else from the SYCL team as part of this tracker? I made the tracker since CI was failing and that's fixed now, so I'll close it if there's nothing needed. Thanks!