how to pass parameters by value to kernel?

[KestutisMa]

How to to pass struct, int or other type parameters by value to kernel?
Is it possible only by creating buffers?

CUDA call of "kernel_1" equivalent would be kernel_1<<<grid, block>>>(dev_n, dev_m, dev_h, dev_I, dev_V_YX, i, i_reg, dt_sim, cellCountX, cellCountY); where "dev_n, dev_m, dev_h, dev_I, dev_V_YX" are device pointers, and "i, i_reg, dt_sim, cellCountX, cellCountY)" are variables passed by values.

[mmha]

Hello @KestutisMa,

With SYCL your kernels are not free functions like in CUDA, but function objects. The function object's member variables are your kernel arguments.

Your declaration would roughly translate to this code below.

struct some_other_kernel_arg {
  float a;
  int b;
};

struct kernel_1 {
    cl::sycl::accessor<float, 2,
        cl::sycl::access::mode::read_write,
        cl::sycl::access::target::global_buffer> dev_n, dev_m, dev_h, dev_I, dev_V_YX;
    int i, i_reg, dt_sim, cellCountX, cellCountY;
    some_other_kernel_arg arg; // Structs are no problem

    void operator()(cl::sycl::nd_item<2> idx) const {

    }
};

extern cl::sycl::range<2> grid, block;

int main() {
    cl::sycl::queue q;
    q.submit([](cl::sycl::handler &h) {
        cl::sycl::nd_range<2> range{grid, block};
        kernel_1 k{ /* set kernel args in constructor here */};
        h.parallel_for(range, k);
    });
}

To reduce verbosity you can also use lambdas and their capture list:

struct kernel_1;

int main() {
    cl::sycl::queue q;
    q.submit([](cl::sycl::handler &h) {
        cl::sycl::nd_range<2> range{grid, block};
        cl::sycl::accessor<float, 2,
            cl::sycl::access::mode::read_write,
            cl::sycl::access::target::global_buffer>
        dev_n{ ... }, dev_m{ ... }, dev_h{ ... }, dev_I{ ... }, dev_V_YX{ ... };

        int i, i_reg, dt_sim, cellCountX, cellCountY;
        
        h.parallel_for<kernel_1>(range, [=](cl::sycl::nd_item<2> idx) {
//                                       ^-- Note the copy
            auto gid = idx.get_global_id();
            float x = dev_n[gid];
        });
    });
}

In C++14 and later you can use generalized lambda captures to further decrease verbosity:

h.parallel_for<kernel_1>(range, [
    dev_n = nbuf.get_access(...),
    i = 42,
    // ...
  ](cl::sycl::nd_item<2> idx) {

});

In the spec you can find the exact restrictions on what may be passed or used otherwise:

A kernel can be defined as a named function object type. These function objects provide the same functionality as any C++ function object, with the restriction that they need to follow C++11 standard layout rules.

Structures containing pointers may be shared but the value of any pointer passed between SYCL devices or between the host and a SYCL device is undefined

Variables with thread storage duration are not allowed to be odr-used in kernel code

Variables with static storage duration that are odr-used inside a kernel must be const or constexpr and zero-initialized or constant-initialized.

Let me know if that helps :)

[KestutisMa]

Thank you, Morris, for detailed answer. I will try this.