[SYCL] Remove need to mark free functions with SYCL_EXTERNAL attribute

[rdeodhar]

This change removes the need to mark free functions with the SYCL_EXTERNAL attribute.

Within clang some instances of add_ir_attribute_function contain non-constexpr expressions (for example, introduced by sycl/ext/oneapi/matrix/matrix-unified.hpp). This occurs when there are non-instantiated functions/methods used in defining attribute values. These expressions are converted into constexprs later in the compilation pipeline.

Calling getAttributeNameValuePairs unconditionally leads to compile-time asserts. To account for the non-constexpr attribute values, a check is first made to rule out instances of the problematic values.

[elizabethandrews]

FE tests should not invoke the driver

[elizabethandrews]

Diagnostic tests should be in test/SemaSYCL

[elizabethandrews]

This test still does not show the FE bug you have fixed. Like I mentioned in the previous comments, please provide a test where SYCLAddIRAttributesFunction attribute has the values which caused a crash when used with free functions earlier. I assume you can just write a new test in test/SemaSYCL with the problematic case from add_ir_attributes_function (which you have now fixed) to illustrate this.

[rdeodhar]

Any test that includes sycl.hpp fails without the hasDependent expr check I added. This is because sycl.hpp includes sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp, which has in it:

#if defined(__SYCL_DEVICE_ONLY__)
  [[__sycl_detail__::add_ir_attributes_function(
      "sycl-joint-matrix-type", "sycl-joint-matrix-use",
      "sycl-joint-matrix-rows", "sycl-joint-matrix-cols",
      sycl::detail::convertTypeToMatrixTypeString<T>(),
      sycl::detail::convertMatrixUseEnumToString(Use), Rows, Cols)]]
#endif // defined(__SYCL_DEVICE_ONLY__)

That amounts to >334 tests.

A test-case with simply #include <sycl/sycl.hpp> fails without the check I added. So what is the point of adding yet another test that duplicates the conditions in hundreds of existing tests?

Or perhaps if I included "sycl.hpp" in the test "free_function_implicit_sycl_extern.cpp" instead of a subset of the SYCL headers as it now does, then "sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp" would get included. Then, the failing conditions would be introduced and the test would demonstrate that the problem has been fixed?

[AaronBallman]

So what is the point of adding yet another test that duplicates the conditions in hundreds of existing tests?

I'm confused.

The title of the PR is about removing a need to mark free functions with SYCL_EXTERNAL. The code changes in the PR are talking about constexpr but then implement something about dependency. So we asked you for test coverage to help us understand what it was you were fixing, and now you're pointing us to a constructor in a header file and saying that provides the test coverage. How? It's not a free function.

Also, "this gigantic pile of header files exercises the code path and so that should be sufficient" is really unhelpful to reviewers. We want to look at a test case and go "oh, that's why those changes happened." We should not have to dig through tens of thousands of lines of entirely unrelated header file content for something that can be tested in a self-contained manner. Tests in Clang should have as few includes in them as possible (for a whole host of reasons, including readability).

[rdeodhar]

The check for a free function now occurs on every function in a compilation unit. The joint matrix header introduces add_ir_attribute_function values that cannot he handled by simply calling getAttributeNameValuePairs. Calling hasDependentExpr before, filters out the problematic cases.

Do you want to see some atributes added to a free function similar to what is done by the joint matrix header?

Would that modification to the test free_function_implicit_sycl_extern.cpp satisfy the need?

[AaronBallman]

Oh wait, are there two things happening in this review? What I think I'm understanding is:

• We needed to check for dependent arguments to the SYCLAddIRAttributesFunctionAttr attribute when processing a declaration; if we don't check for dependent arguments, we'd hit assertions.
• Once we fix that bug, we can treat free functions the same as member functions so that they don't require a special attribute to say it's available on the device.

Is that accurate?

[elizabethandrews]

Is that accurate?

This is my understanding. The code added in isFreeFunction was problematic before because it wasn't doing the dependency check.

Would that modification to the test free_function_implicit_sycl_extern.cpp satisfy the need?

I would suggest removing all headers from that test and writing a simple test which applies the attribute with the previously problematic arguments, on any function which will hit isFreeFunction. FE lit testing is not the same as E2E tests. We need to see reduced case with the dependency bug which you have now fixed.

[AaronBallman]

This is my understanding.

Thank you!

I would suggest removing all headers from that test and writing a simple test which applies the attribute with the previously problematic arguments, on any function which will hit isFreeFunction.

+1, based on my understanding, you should be able to write a free function template with a dependent attribute argument and is marked SYCL_EXTERNAL; that should reproduce the assertion before your changes, and helps us see that your changes are what removed the assertion.

[rdeodhar]

In trying to create a test-case that shows the problem with non-constexpr values in add_ir_attribute_function I discovered that the "bad" values only occur when the attribute is applied to a member function. See code below.

template <typename T> constexpr int value() { return 5; }
template <typename T> struct S {
#if defined(__SYCL_DEVICE_ONLY__)
  [[__sycl_detail__::add_ir_attributes_function("name", value<T>())]]
#endif
  S() {
  }
};
template <typename T>
__attribute__((sycl_device)) [[__sycl_detail__::add_ir_attributes_function(
    "name",
    value<T>())]] [[__sycl_detail__::
                        add_ir_attributes_function("sycl-single-task-kernel",
                                                   0)]] void
f(T i) {}
template void f(int i);

For the member function S() the attribute looks like this:

`-SYCLAddIRAttributesFunctionAttr 0x55989d69e7c0 <col:5, col:67>
  |-ConstantExpr 0x55989d69e870 <col:49> 'const char[5]' lvalue
  | |-value: LValue <todo>
  | `-StringLiteral 0x55989d69e590 <col:49> 'const char[5]' lvalue "name"
  `-CallExpr 0x55989d69e650 <col:57, col:66> '<dependent type>'
    `-UnresolvedLookupExpr 0x55989d69e5d8 <col:57, col:64> '<dependent type>' lvalue (ADL) = 'value' 0x55989d675b30
      `-TemplateArgument type 'T':'type-parameter-0-0'
        `-TemplateTypeParmType 0x55989d675ce0 'T' dependent depth 0 index 0
          `-TemplateTypeParm 0x55989d675c90 'T'

while for the function 'f' the attribute looks like this:

|-SYCLAddIRAttributesFunctionAttr 0x56361c3c3ea8 <line:37:32, line:39:15>
| |-ConstantExpr 0x56361c3c3f00 <line:38:5> 'const char[5]' lvalue
| | |-value: LValue <todo>
| | `-StringLiteral 0x56361c398cf0 <col:5> 'const char[5]' lvalue "name"
| `-ConstantExpr 0x56361c3c3f60 <line:39:5, col:14> 'int'
|   |-value: Int 5
|   `-CallExpr 0x56361c3c3e88 <col:5, col:14> 'int'
|     `-ImplicitCastExpr 0x56361c3c3e70 <col:5, col:12> 'int (*)()' <FunctionToPointerDecay>
|       `-DeclRefExpr 0x56361c3c3dc0 <col:5, col:12> 'int ()' lvalue Function 0x56361c3c3cc8 'value' 'int ()' (FunctionTemplate 0x56361c398a90 'value')

This leads me to believe that the bad values in the attribute will only occur on member functions. Since free functions cannot be member functions filtering out member functions before checking for a function being a free function avoids the problem and is a good idea anyway. With just that change in SemaDecl.cpp all tests pass.
On the other hand, perhaps the bad attribute values could be introduced in some other way. In which case the hasDependentExpr check that I've added in SemaSYCL.cpp would still be necessary.

I've updated the PR to keep both the filtering out of member functions and the hasDependentExpr check. Its a "belt and suspenders" solution. I'd like your opinion whether to keep it that way or remove the hasDependentExpr check.

[elizabethandrews]

Based on your tests below and explanation in PR comments, it sounds like this issue does not have anything to do with constexpr values. The issue here is we have non-instantiated functions/methods being checked when they shouldn't be. Can you try replacing this check with LangOpts.SYCLIsDevice && Body && FD->isDependentContext() without the TemplatedKind checks. You also do not need the if (FD) since L16305 above checks this already.

I think you should also replace the hasDependentExpr checks you added in SemaSYCL to asserts instead because we should not be having dependent arguments at that point.

[AaronBallman]

Slight correction, that should be !FD->isDependentContext() -- basically, if the function isn't dependent, it's safe to check the arguments; if the function IS dependent somehow, we need to wait to do those checks until the function has been instantiated (regardless of whether it's a member function or free function).

[rdeodhar]

The hasDependentExpr check I added was a check before calling getAttributeNameValuePairs because that would assert if all expressions in the attribute were not constexprs. So there is already an assert there. Perhaps an assert before calling getAttributeNameValuePairs is not needed?

Making the other change in SemaDecl.cpp was fine. The tests all pass.

[elizabethandrews]

Ok. You can remove all that the code you added then if one already exists in getAttributeNameValuePairs

[rdeodhar]

Done.

[elizabethandrews]

Can you please update this comment?

You can say something along the lines of "This test ensures that the compiler does not crash when functions or methods with add_ir_attributes_function attribute in dependent contexts are checked for the presence of the free function property"

[rdeodhar]

OK, I've updated the comment.

[dm-vodopyanov]

@rdeodhar tag in PR title should be in [ and ]

https://github.com/intel/llvm/blob/sycl/sycl/doc/developer/ContributeToDPCPP.md#commit-message

[rdeodhar]

@rdeodhar tag in PR title should be in [ and ]

https://github.com/intel/llvm/blob/sycl/sycl/doc/developer/ContributeToDPCPP.md#commit-message

OK, done.