[sycl free func] The example in sycl_ext_oneapi_free_function_kernels.asciidoc can't be compiled successfully

[NeoZhangJianyu]

Describe the bug

I try to build the example code in https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_free_function_kernels.asciidoc. The code can't be compiled successfully. User can't learn how to use the feature. The previous release of example in this doc works well.

oneAPI:

icpx --version
Intel(R) oneAPI DPC++/C++ Compiler 2025.3.0 (2025.3.0.20251010)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /opt/intel/oneapi/compiler/2025.3/bin/compiler
Configuration file: /opt/intel/oneapi/compiler/2025.3/bin/compiler/../icpx.cfg

Build cmd and error log:

cpx -fsycl sycl_free_base.cpp 
sycl_free_base.cpp:23:39: error: no member named 'kernel_function' in
      namespace 'sycl::ext::oneapi::experimental'
   23 |   syclexp::nd_launch(q, ndr, syclexp::kernel_function<iota>, 3.14f, ptr);
      |                                       ^~~~~~~~~~~~~~~
sycl_free_base.cpp:23:60: error: expected expression
   23 |   syclexp::nd_launch(q, ndr, syclexp::kernel_function<iota>, 3.14f, ptr);
      |                                                            ^

Example code:

#include <sycl/sycl.hpp>
namespace syclext = sycl::ext::oneapi;
namespace syclexp = sycl::ext::oneapi::experimental;

static constexpr size_t NUM = 1024;
static constexpr size_t WGSIZE = 16;

SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
void iota(float start, float *ptr) {
  // Get the ID of this kernel iteration.
  size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();

  ptr[id] = start + static_cast<float>(id);
}

int main() {
  sycl::queue q;
  sycl::context ctxt = q.get_context();

  float *ptr = sycl::malloc_shared<float>(NUM, q);

  sycl::nd_range ndr{{NUM}, {WGSIZE}};
  syclexp::nd_launch(q, ndr, syclexp::kernel_function<iota>, 3.14f, ptr);

  q.wait();
}

To reproduce

1. Include a code snippet that is as short as possible
2. Specify the command which should be used to compile the program
3. Specify the command which should be used to launch the program
4. Indicate what is wrong and what was expected

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

[steffenlarsen]

Hi @NeoZhangJianyu, thank you for using the extension and letting us know! The free function kernel extension is still in the "proposed" state, which generally means that it is at best partially implemented. We are currently working on bringing it up-to-speed with the most recent changes.

[NeoZhangJianyu]

Yes, I see.

I suggest adding the correct release (internal or external) of compiler to support this example successfully. Or provide a right example for latest release, like 2025.3.1.

[gmlueck]

Note that the syclexp::kernel_function feature that you are trying to use is being added in #20698.