OIDN (Blender) build broken in nightly

[sarnex]

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

[sarnex]

Likely guility PR #21875

Since the API seems deprecated, maybe OIDN needs to be updated instead of a compiler change @atafra?

FYI @againull

[atafra]

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]

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]

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]

@atafra

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.

@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]

@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]

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!

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 prefix ext_<vendorstring>. DPC++ extensions typically use oneapi for 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]

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]

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]

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]

@atafra Is there an OIDN patch I can apply that moves to the ESIMD kernel so can we fix the CI here? Thanks