Cannot capture 'this' in a SYCL kernel

SYCL development
1

What is the best way of fixing this compilation error?
E.g. in

class Foo {
  int Bar = 123;

public:
  int doSomething() {
    [...]
      cgh.parallel_for<Foo>(nd_range<1>(range<1>(5),range<1>(5)), [=](nd_item<1>) {
        printf("%d", Bar);
      });
    [...]
  }
1 Like
2

This is a deliberate restriction. You don’t want your “this” object to be sent to the accelerator every time a kernel is invoked that uses one of the fields.

The easiest way in this case is to make a local variable that copies the field of ‘this’ and use that in the lambda instead. I.e

int Bar = this->Bar;
cgh.parallel_for<Foo>(nd_range<1>(range<1>(5),range<1>(5)), [=](nd_item<1>) {
  printf("%d", Bar);
});

If Bar was an array it would need to be wrapped inside a struct (that behaves like an array) and then passed to the kernel.
In cases of more complex data structures, ensure that calling the copy constructor does not have any side effects.

In cases where you want to pass in a lot of these fields you could copy the whole of this into a local:

auto self = *this;
cgh.parallel_for<Foo>(nd_range<1>(range<1>(5),range<1>(5)), [=](nd_item<1>) {
  printf("%d", self.Bar);
});

However, if the copy constructor of your object is non-trivial this might have a performance impact.
This approach will also not work if the type of this is abstract.

An alternative is to create a class for your lambda, which contains your field(s) and pass an instantiation to your parallel_for:

class my_function_object {
  int Bar = 123;

 public:
  void operator()(nd_item<1> range) { printf("%d", Bar); }
};

[...]
   
      my_function_object obj;
      cgh.parallel_for(nd_range<1>(range<1>(5),range<1>(5)), obj);

This class should also contain any accessors you want to capture as members.

2 Likes