[SYCL] Diagnostics should be generated at point of kernel invocation

[elizabethandrews]

Kernel location is set differently when kernel is defined using a lambda vs a functor. Diagnostics are emitted at point of invocation when lambdas are used but emitted at functor declaration when functors are used. It should be generated at point of invocation for functor as well.

[srividya-sundaram]

Diagnostics are emitted at point of invocation when lambdas are used

I am not sure this is entirely true.
In unnamed-kernel.cpp , if I have the lambda defined outside of the kernel-invocation-function like below:

class InvalidKernelName1 {};
q.submit([&](cl::sycl::handler &h) {
  auto x = [] { int a = 1; };
  h.single_task<InvalidKernelName1>(x);
});

The diagnostic emitted is :

unnamed-kernel.cpp:30:16: error: kernel needs to have a globally-visible name
      auto x = [] { int a = 1; };
               ^

In the case where the lambda function is defined in-line :

class InvalidKernelName1 {};
    q.submit([&](cl::sycl::handler &h) {
      h.single_task<InvalidKernelName1>([] {});
    });

The diagnostic emitted is :

unnamed-kernel.cpp:31:41: error: kernel needs to have a globally-visible name
      h.single_task<InvalidKernelName1>([]{});
                                           ^

The diagnostics appears to be emitted at point of kernel-invocation-function, but it really is pointing to the line where the lambda function is defined and the kernel-invocation-function happens to be in the same line.

Also, in ConstructOpenCLKernel , KernelCallerFunc->getLocation() points to kernel_single_task

And KernelObj is the lambda object, and by the definition of getLocation(), it seems to return the location where the object was first declared(and defined?)
In SemaSYCL.cpp, KernelCallerFunc is used to generate OpenCL kernel. But this does not correspond to the actual kernel invocation function in user’s code.

I am not sure how to get the actual kernel invocation function mapping in ConstructOpenCLKernel

@Fznamznon @erichkeane

[erichkeane]

Without the CallExpr, you can't find the call location. Unfortuantely, that is long gone by the time we get to ConstructOpenCLKernel, since we're only doing the instantiation at that point. AND, we don't have a good way of finding that call-expr, since there isn't a good place to record it.

[erichkeane]

Actually, browsing the FunctionDecl interfae, getPointOfInstantiation (so KernelCallerFunc->getPointOfInstantiation) seems to say that it returns the location where it was first forced to be instantiated. So you COULD use that, and I think it would be at least at the call location there. (so it might be the 'h' 'h.single_task(lambda)'. IF that isn't empty(and it is a default argument to the setFunctionTemplateSpecialization function) it is probably about as good as we can do.

[srividya-sundaram]

KernelCallerFunc->getPointOfInstantiation does not point to kernel-invocation-function in user code.

The diagnostic looks like this:

/forked-repo/llvm/clang/test/SemaSYCL/unnamed-kernel.cpp:3:
/iusers/sundaram/forked-repo/llvm/clang/test/SemaSYCL/Inputs/sycl.hpp:204:5: error: kernel needs to have a globally-visible name
    kernel_single_task<NameT>(kernelFunc);
    ^

[erichkeane]

Right, as I said, that gets you the 'call location'. I think that is about as good as you're going to be able to get here. It isn't perfect, but at least it is consistent. Like I said, any attempt to find the actual callexpr is going to be a pretty massive, difficult, and error-prone task. Call expressions are created everywhere, so it would be tough to track them all down. Point of instantiation is at least the correct line all the time.

The other thought I have, is perhaps use the existing Function-object location, then use a note like, 'used in kernel here' which uses the point-of-instantiation. It seems like it would at least LOOK not terrible. We wouldn't get the template-instantiation stack out of it, (so it might simply point into library code sometimes), but maybe it would be more useful?

[erichkeane]

To clarify my first point:

The only solution that gets us the actual call-expr is to diagnose it when we are diagnosing the normal call-expr (instead of in constructOpenCLKernel). See how 'deprecated' and the 'availability' attributes are implemented (though neither of those depend on the argument-types passed, so I'm not sure that info is available there).

They are actually pretty error-prone, particularly when it comes to the template instantiations. AND since SYCL kernel instantiations are special, you'd need to make sure you don't cause the template to be instantiated EARLY. We count on the template instantiation being a delayed instantiation in order for it it work.

Quick edit: I looked through the clang attributes list and found 2 more that might be better htan availability/deprecated to investigate:
https://clang.llvm.org/docs/AttributeReference.html#format
https://clang.llvm.org/docs/AttributeReference.html#diagnose-if

Both interrogate the arguments and diagnose based on them. I'm not sure how they work with templates (again, the 1st is supposed to be for printf-C functions, and the 2nd is not particularly well maintained), but they might have easier enough implementations to improve your investigation.

Either way, you'd need to move these types of diagnostics out of SemaSYCL and into the rest of the compiler.

[elizabethandrews]

The other thought I have, is perhaps use the existing Function-object location, then use a note like, 'used in kernel here' which uses the point-of-instantiation.
@erichkeane When would this point to the library code and not user code where single task is instantiated?

[erichkeane]

This when the call to the sycl_kernel is located in the library :)

In the case above, single_task is NOT the kernel, it is a wrapper for it. The ACTUAL one is kernel_single_task, which is called by single_task.

PERHAPS we could do some sort of instantiation-tree walking until we hit non-library code (as hideous as that would be), but the other suggestion I had (diagnosing it at the CheckCall) seems superior to me on diagnostics where that is possible (that is, it doesn't require seeing the definition of any functions CALLED by the kernel).