The term "kernel" is used for multiple purposes

[tahonermann]

Consider this text from SYCL 2020 section 4.12.4, "Rules for parameter passing to kernels"

A SYCL application passes parameters to a kernel in different ways depending on whether the kernel is a named function object or a lambda function. If the kernel is a named function object, the operator() member function (or other member functions that it calls) may reference member variables inside the same named function object. Any such member variables become parameters to the kernel. If the kernel is a lambda function, any variables captured by the lambda become parameters to the kernel.

The text is actually nonsensical as written. How can member variables (more precisely, non-static data members) of the kernel become parameters to the kernel?

Some uses of the term "kernel" refer to the named function object or lambda function (more precisely, closure object) that is passed to a kernel invocation function. Others appear to refer to something else; presumably an implementation dependent device entry point function; something the SYCL specification shouldn't have to concern itself with. Note that the term "entry point" appears exactly one time in the SYCL 2020 specification; in the definition of "SYCL kernel function" in the glossary.

SYCL kernel function

A type which is callable with operator() that takes an id, item, nd-item or work-group, and an optional kernel_handler as its last parameter. This type can be passed to kernel enqueue member functions of the command group handler. A SYCL kernel function defines an entry point to a kernel. The function object can be a named device copyable type or lambda function.

This description conflates types and objects.

[gmlueck]

This is also related somewhat to #524.

[TApplencourt]

I can try to have a go at it. I remember getting confused too

[TApplencourt]

The text is actually nonsensical as written. How can member variables (more precisely, non-static data members) of the kernel become parameters to the kernel?

I think the sentence you quote is there to describe functors.

template <typename TAccessorW, typename TAccessorR>
class memcopy_kernel {
public:
  memcopy_kernel(TAccessorW accessorW_, TAccessorR accessorR_)
      : accessorW{accessorW_}, accessorR{accessorR_} {}
  void operator()(sycl::item<1> idx) const {
    accessorW[idx.get_id()] = accessorR[idx.get_id()] + 1;
  }

private:
  TAccessorW accessorW;
  TAccessorR accessorR;
};

// And then
  cgh.parallel_for(sycl::range<1>(global_range),
                   memcopy_kernel(accessorW, accessorR));

But I agree, we should just not care. And always talk about "SYCL Kernek function" IMO. AFAIK lambda are just a nameless functor (with as you point out, so mechanism for capturing argument, hence the confusing between arguments and member variable)

[TApplencourt]

Ok after more digging, it's even more fun. I found 3 Definitions of "SYCL Kernel Function"

• A function object that can execute on a device exposed by a SYCL backend API is called a SYCL kernel function.

• A SYCL kernel function is the scoped block of code that will be compiled using a device compiler. This code may be defined by the body of a lambda function or by the operator() function of a function object. Each instance of the SYCL kernel function will be executed as a single, though not necessarily entirely independent, flow of execution and has to adhere to restrictions on what operations may be allowed to enable device compilers to safely compile it to a range of underlying devices.

• A type which is callable with operator() that takes an id, item, nd-item or work-group, and an optional kernel_handler as its last parameter. This type can be passed to kernel enqueue member functions of the command group handler. A SYCL kernel function defines an entry point to a kernel. The function object can be a named device copyable type or lambda function.

Note that they are conflicting them self as @tahonermann . One is a type, and the other is the block of code. And we never define the instance of the class.

And kernel (that we use everywhere) is defined only one:

• A kernel represents a SYCL kernel function that has been compiled for a device, including all of the device functions it calls. A kernel is implicitly created when a SYCL kernel function is submitted to a device via a kernel invocation command. However, a kernel can also be created manually by pre-compiling a kernel bundle (see Section 4.11).

So lets take a example:

template <typename TAccessorW, typename TAccessorR>
class memcopy_kernel {  // No name for thas? SYCL Kernel Function Class?
public:
  memcopy_kernel(TAccessorW accessorW_, TAccessorR accessorR_)
      : accessorW{accessorW_}, accessorR{accessorR_} {}

  void operator()(sycl::item<1> idx) const {                         // This is the SYCL Kernel Function?
    accessorW[idx.get_id()] = accessorR[idx.get_id()] + 1; //
  }                                                                                          //

private:
  TAccessorW accessorW;
  TAccessorR accessorR;
};

 memcopy_kernel foo(accessorW, accessorR); // Creation of the Instance of class who defines our SYCL Kernel function. We don't have a name for it yet. 
  cgh.parallel_for(sycl::range<1>(global_range),  foo); 
    // I pass `foo`, a instance of ` function object`. It will be implicitly "compiled" into a `kerne`l. And then called. What is called, the `kernel` of the `instance` of the function object?  or both?

I'm not smart enough to understand how to use Kernel bundle. And the difference between kernel, and sycl::kernel. And between.

So I short:

• We should define more clearly SYCL function object and differentiate between the Object Type, the code in the operator(), the instance of it.
• The distinction between kernel and sycl::kernel seems unclear to me too.