[SYCL] Fix for detection of free function calls.

[rdeodhar]

This change simplifies the code to detect calls to ""free functions", i.e., functions such as this_item().
The existing code does not work in all cases to detect the Kernel function, while the corrected code uses an already known Kernel function.

Signed-off-by: rdeodhar [email protected]

[premanandrao]

Could you please add a test that shows the cases that didn't work before?

[rdeodhar]

There is an existing test llvm_test_suite_sycl/basic_free_function_queries_free_function_queries.

[elizabethandrews]

The existing code does not work in all cases to detect the Kernel function,

Can you please add a test to illustrate this. Every change is supposed to be accompanied by a test unless it is NFC

[cperkinsintel]

LGTM.

[rdeodhar]

A test has been added for the case:
parallel_for<class X>(range<1>{n}, [=](auto I) {}

[elizabethandrews]

Oh I see. So getOperatorParens(KernelObj) doesn't return operator method when auto is used? Do you understand why? It is not obvious to me why. @premanandrao could you also please take a look?

[premanandrao]

@rdeodhar, is this the case that was failing before?

[rdeodhar]

Yes, this particular construct in an existing runtime test was failing. I have replicated the situation in the existing clang lit test. Owing to some templating of members, the operator() of the KernelObj was not being found. However, the code looking for the operator() is unnecessary because the kernel function is already available.

[elizabethandrews]

Hmm... That means KernelObj is defined differently when parameter type is auto? Do you know why this happens? I'm approving the patch since the answer to this question is probably outside the scope of this patch. However, if you know the answer I am curious :)

[erichkeane]

IMO, there is likely value to instead fix the poorly named getOperatorParens instead.

[elizabethandrews]

@erichkeane would you know why its getting wrapped?

[AaronBallman]

@erichkeane well that's a fun case -- I think you can use decls_begin() and decls_end() to do this:

static const CXXMethodDecl *getOperatorParens(const CXXRecordDecl *Rec) {
  auto MethIter = llvm::find_if(Rec->methods(), [](const CXXMethodDecl *MD) {
    return MD->getOverloadedOperator() == OO_Call;
  });
  if (MethIter != Rec->methods_end())
   return *MethIter;

  using function_template_iterator = specific_decl_iterator<FunctionTemplateDecl>;
  using function_template_range = llvm::iterator_range<specific_decl_iterator<FunctionTemplateDecl>>;
  auto FTIter = llvm::find_if(function_template_range(function_template_iterator(Rec->decls_begin()),
                                                      function_template_iterator(Rec->decls_end())),
                                           [](const FunctionTemplateDecl *FTD) {
                                             if (const auto *MD = dyn_cast<CXXMethodDecl>(FTD->getTemplatedDecl()))
                                               return MD->getOverloadedOperator() == OO_Call;
                                             return false;
                                           });
  if (FTIter != function_template_iterator(Rec->decls_end()))
   return *FTIter;
  return nullptr;
}

(This is totally untested and was written in a web browser, so YMMV.)

[erichkeane]

@erichkeane would you know why its getting wrapped?

Because we ALWAYS wrap templates. We need something in the AST to contain BOTH the uninstantiated and instantiated template.

We should probably update the getOperatorParens, but, IMO, it should probably return an array of valid (non-dependent) operators. I'm guessing this whole idea of using getOperatorParens is flawed though. Trying to figure out the call operator of the type rather than checking the callgraph is likely the wrong way about it. @kbobrovs : Please take a look at how you're using getOperatorParens and make sure it would work in a case where there is multiple operator()s, including instantiated template versions.

IN the case of what @AaronBallman just posted (thanks btw!), I don't know if that works for 2 reasons: First, I think the return *FTIter; line is wrong, since that returns a FunctionTemplateDecl*, instead of a CXXMethodDecl*. Second, it has the same problem of only returning the FIRST operator () :(

[AaronBallman]

Good call, that should be return cast<CXXMethodDecl>(*FTIter) instead, and it definitely only finds the first. As for how to handle multiple operator() definitions -- I agree that the function should either return a container of all the operators or take the container as a parameter.

[kbobrovs]

@kbobrovs : Please take a look at how you're using getOperatorParens and make sure it would work in a case where there is multiple operator()s, including instantiated template versions.

AFAIR, this simplistic approach was based on assumption that the API code is under our control, and we know that there are no other operator '()' calls from the parallel_for_kernel.

[elizabethandrews]

Hmm... That means KernelObj is defined differently when parameter type is auto? Do you know why this happens? I'm approving the patch since the answer to this question is probably outside the scope of this patch. However, if you know the answer I am curious :)

[rdeodhar]

The SYCL handler.hpp makes a special-case of an integral kernel parameter type, like so:

    using TransformedArgType =
        typename std::conditional<std::is_integral<LambdaArgType>::value &&
                                      Dims == 1,
                                  item<Dims>, LambdaArgType>::type;

Now I don't know why that leads to the behavior we see, but it was not worthwhile to pursue it for this bug fix, because the code that doesn't work as expected is not needed at all.

[elizabethandrews]

The SYCL handler.hpp makes a special-case of an integral kernel parameter type, like so:

    using TransformedArgType =
        typename std::conditional<std::is_integral<LambdaArgType>::value &&
                                      Dims == 1,
                                  item<Dims>, LambdaArgType>::type;

Now I don't know why that leads to the behavior we see, but it was not worthwhile to pursue it for this bug fix, because the code that doesn't work as expected is not needed at all.

Ok. Thanks for explaining