[SYCL][SemaSYCL] Functor Inheritance Incorrect Outline/CodeGen

[agozillon]

The example is a little contrived in this case, but from what I can tell it's all legal SYCL (please correct me if I am wrong though as my Spec-foo still needs a lot of work). I've been looking into it as I'm trying to find a solution for it myself ASAP, although I imagine if I find one it might not be ideal.

#include <CL/sycl.hpp>

struct empty_base_class {};

struct functor : empty_base_class {
  uint16_t foo;
  uint32_t bar;

  void operator()() {
    auto tmp_foo = foo;
    auto tmp_bar = bar;
  }
};

int main() {
  cl::sycl::queue q;

  q.submit([&](cl::sycl::handler &cgh) {
     functor f{};
     cgh.single_task(f);
  });

  return 0;
}

So the above test will result in:

error: Explicit load/store type does not match pointee type of pointer operand (Producer: 'LLVM9.0.0svn' Reader: 'LLVM 9.0.0svn')

From what I can tell, this seems to be a side affect of moving from the illegal assignment style of AST generation to the more legal Initialization method. The InitExprList isn't setup to handle the Base classes initialization, so when it gets to the function VisitInitListExpr: https://github.com/intel/llvm/blob/sycl/clang/lib/CodeGen/CGExprAgg.cpp#L1502

It results in a sort of off by 1 (or off by the number of base classes I suppose) error in the above case where an invalid store is generated because you're matching the wrong Initialization Expr against the wrong destination type. In the above case it'll be something like the store of an i16 vs i32* because of the two uint scalar types being initialized.

You can change the uint types to be the same size/type and you won't get the error, but I believe the IR will not generate all of the required stores (not looked too deeply into that) and I imagine if any type is different it'll result in a similar load/store error.

Worth noting that the error is from a Release mode compilation, in Debug it will probably crash/assert here:
https://github.com/intel/llvm/blob/sycl/clang/lib/CodeGen/CGExprAgg.cpp#L1846

[mikerice1969]

You are correct that we are not handling functors with base classes correctly. In the prior assignment-based code the base class portion of the functor would have been left uninitialized and code that used anything from that base would have potential problems. Of course there is nothing to initialize in an empty base so this specific example didn't have a problem.

My first thought is add (in CreateOpenCLKernelBody) an initializer to the list for each base before the fields are initialized. We likely need to use the same base initializers that the functor copy constructor uses.

[agozillon]

I imagine that would work; at the moment I'm doing a workaround that I think is similar, I add an InitExpr that initializes the base classes using their default constructor prior to the creation of the field initialization, it seems to work at least for the use case I currently have (which doesn't really make much use of data in the base class at the moment) but I imagine it's quite brittle.

I'm wondering if just adding a base class initializer is enough for more complex cases though, mainly because the OpenCL kernel initializes the functor based on parameters passed into the kernel, so I'm wondering if there is certain cases where the base class will also add some parameter dependencies to the kernel and also need to be considered. Perhaps I'm misunderstanding part of the process and over thinking it though!

[mikerice1969]

Right. It is not clear how to correctly initialize the bases here. Consider this:

struct Base {
  Base(int i) : i(i) {}
  int i;
};

struct Functor : Base {
  Functor(int i) : Base(i) {}
  Functor(const Functor& F) : Base(F.i) {}
  void operator()() {}
};

There is no default constructor for Base. In the Functor copy constructor the Base(int) constructor is used. In this case F.i is used but it could use a constant, global variable, function call, etc. To find the correct constructor we’d need to analyze the Functor copy constructor and we could only do that if the body were actually available in the translation unit. And as you mention the parameters needed to initialize the bases may not be passed into the kernel.

So I don’t really have a good idea how to fix this in general. We probably need some help from the SYCL language spec to deal with this. Maybe it should be disallowing functors with base classes.

How important is it to you to have base classes for your functor? Would adding the base class as a member work instead?

[agozillon]

Yes a catch all solution for this does look like it may be difficult, I'm unsure what the specifications standpoint is on it but it would be nice to have some clarity.

Your questions might be better answered by @keryell in this case! However, at least for the immediate future the workaround I have in place seems to work for the use case we currently have, long-term that probably won't be the case though.

[kbobrovs]

I'm wondering if there is certain cases where the base class will also add some parameter dependencies to the kernel and also need to be considered.

Yes, for example consider this case. Most functors used in practice have accessor fields. The Base class can have it too, and the accessor can be used in Functor's () operator. For that to work the accessor in the base should be passed as a separate parameter (3 parameters actually).