[SYCL] Remove need to mark free functions with SYCL_EXTERNAL attribute (#14170)

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.

Author: rdeodhar

clang/lib/Sema/SemaDecl.cpp

@@ -16318,10 +16318,7 @@ Decl *Sema::ActOnFinishFunctionBody(Decl *dcl, Stmt *Body,
     checkTypeSupport(FD->getType(), FD->getLocation(), FD);
 
   // Handle free functions.
-  if (LangOpts.SYCLIsDevice && FD->hasAttr<SYCLDeviceAttr>() && Body &&
-      (FD->getTemplatedKind() == FunctionDecl::TK_NonTemplate ||
-       FD->getTemplatedKind() ==
-           FunctionDecl::TK_FunctionTemplateSpecialization))
+  if (LangOpts.SYCLIsDevice && Body && !FD->isDependentContext())
     SYCL().ProcessFreeFunction(FD);
 
   return dcl;

clang/lib/Sema/SemaSYCL.cpp

@@ -1065,10 +1065,8 @@ static target getAccessTarget(QualType FieldTy,
       AccTy->getTemplateArgs()[3].getAsIntegral().getExtValue());
 }
 
-// FIXME: Free functions must have void return type, be declared at file scope,
-// outside any namespaces, and with the SYCL_DEVICE attribute. If the
-// SYCL_DEVICE attribute is not specified this function is not entered since the
-// possibility of the function being a free function is ruled out already.
+// FIXME: Free functions must have void return type and be declared at file
+// scope, outside any namespaces.
 static bool isFreeFunction(SemaSYCL &SemaSYCLRef, const FunctionDecl *FD) {
   for (auto *IRAttr : FD->specific_attrs<SYCLAddIRAttributesFunctionAttr>()) {
     SmallVector<std::pair<std::string, std::string>, 4> NameValuePairs =

clang/test/CodeGenSYCL/free_function_int_header.cpp

@@ -8,7 +8,6 @@
 #include "sycl.hpp"
 
 // First overload of function ff_2.
-__attribute__((sycl_device))
 [[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel",
                                               2)]] void
 ff_2(int *ptr, int start, int end) {
@@ -17,7 +16,6 @@ ff_2(int *ptr, int start, int end) {
 }
 
 // Second overload of function ff_2.
-__attribute__((sycl_device))
 [[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel",
   2)]] void
   ff_2(int* ptr, int start, int end, int value) {
@@ -27,7 +25,6 @@ __attribute__((sycl_device))
 
 // Templated definition of function ff_3.
 template <typename T>
-__attribute__((sycl_device))
 [[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]] void
 ff_3(T *ptr, T start, T end) {
   for (int i = start; i <= end; i++)

clang/test/Sema/free_function_attribute_check.cpp

@@ -0,0 +1,58 @@
+//==---- free_function_implicit_sycl_extern.cpp ----------------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+// RUN: %clang_cc1 -fsycl-is-device -triple -spir64-unknown-unknown -verify %s
+
+// expected-no-diagnostics
+
+// This test confirms that functions or methods with add_ir_attributes_function
+// attribute values in dependent contexts can be handled without assertions when
+// checking for the presence of free function properties.
+
+template <typename T> constexpr int value() { return 5; }
+
+// In this struct the function the add_ir_attributes_function values for "S()"
+// are as follows. Note that the "value" is represented as a CallExpr.
+// `-SYCLAddIRAttributesFunctionAttr 0x562ec6c13390 < col:5, col : 67 >
+// | -ConstantExpr 0x562ec6c13440 < col:49 > 'const char[5]' lvalue
+// | |-value: LValue <todo>
+// | `-StringLiteral 0x562ec6c13160 < col:49 > 'const char[5]' lvalue "name"
+// `-CallExpr 0x562ec6c13220 < col:57, col : 66 > '<dependent type>'
+//   `-UnresolvedLookupExpr 0x562ec6c131a8 < col:57, col : 64 > '<dependent type>' lvalue(ADL) = 'value' 0x562ec6bea700
+//     `-TemplateArgument type 'T':'type-parameter-0-0'
+//       `-TemplateTypeParmType 0x562ec6bea8b0 'T' dependent depth 0 index 0
+//         `-TemplateTypeParm 0x562ec6bea860 'T'
+
+template <typename T> struct S {
+#if defined(__SYCL_DEVICE_ONLY__)
+  [[__sycl_detail__::add_ir_attributes_function("name", value<T>())]]
+#endif
+  S() {
+  }
+};
+
+// For the free function "f" the add_ir_attributes_function values are:
+// | -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')
+
+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);

sycl/test-e2e/KernelAndProgram/free_function_apis.cpp

@@ -12,7 +12,6 @@
 
 using namespace sycl;
 
-SYCL_EXTERNAL
 SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
     (ext::oneapi::experimental::nd_range_kernel<2>))
 void ff_2(int *ptr, int start) {
@@ -25,8 +24,9 @@ void ff_2(int *ptr, int start) {
 
 // Templated free function definition.
 template <typename T>
-SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((
-    ext::oneapi::experimental::single_task_kernel)) void ff_3(T *ptr, T start) {
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
+    (ext::oneapi::experimental::single_task_kernel))
+void ff_3(T *ptr, T start) {
   int(&ptr2D)[4][4] = *reinterpret_cast<int(*)[4][4]>(ptr);
   nd_item<2> Item = ext::oneapi::this_work_item::get_nd_item<2>();
   id<2> GId = Item.get_global_id();

sycl/test-e2e/KernelAndProgram/free_function_kernels.cpp

@@ -46,7 +46,7 @@ bool checkUSM(int *usmPtr, int size, int *Result) {
   return false;
 }
 
-extern "C" SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
+extern "C" SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
     (ext::oneapi::experimental::single_task_kernel)) void ff_0(int *ptr,
                                                                int start,
                                                                int end) {
@@ -96,7 +96,6 @@ bool test_0(queue Queue) {
 }
 
 // Overloaded free function definition.
-SYCL_EXTERNAL
 SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
     (ext::oneapi::experimental::nd_range_kernel<1>))
 void ff_1(int *ptr, int start, int end) {
@@ -147,7 +146,6 @@ bool test_1(queue Queue) {
 }
 
 // Overloaded free function definition.
-SYCL_EXTERNAL
 SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
     (ext::oneapi::experimental::nd_range_kernel<2>))
 void ff_1(int *ptr, int start) {
@@ -203,8 +201,9 @@ bool test_2(queue Queue) {
 
 // Templated free function definition.
 template <typename T>
-SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((
-    ext::oneapi::experimental::nd_range_kernel<2>)) void ff_3(T *ptr, T start) {
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
+    (ext::oneapi::experimental::nd_range_kernel<2>))
+void ff_3(T *ptr, T start) {
   int(&ptr2D)[4][4] = *reinterpret_cast<int(*)[4][4]>(ptr);
   nd_item<2> Item = ext::oneapi::this_work_item::get_nd_item<2>();
   id<2> GId = Item.get_global_id();

sycl/test/extensions/free_function_errors.cpp

@@ -0,0 +1,51 @@
+//==---- free_function_errors.cpp --------------------------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+// RUN: %clangxx  -fsyntax-only -fsycl-device-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s
+
+#include <array>
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+
+struct S {
+  int i;
+  float f;
+};
+
+union U {
+  int i;
+  float f;
+};
+
+using accType = accessor<int, 1, access::mode::read_write>;
+
+// expected-error@+3 {{'struct S' cannot be used as the type of a kernel parameter}}
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
+    (ext::oneapi::experimental::single_task_kernel))
+void ff(struct S s) {}
+
+// expected-error@+3 {{'union U' cannot be used as the type of a kernel parameter}}
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
+    (ext::oneapi::experimental::single_task_kernel))
+void ff(union U u) {}
+
+// expected-error@+3 {{'accType' (aka 'accessor<int, 1, access::mode::read_write>') cannot be used as the type of a kernel parameter}}
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
+    (ext::oneapi::experimental::single_task_kernel))
+void ff(accType acc) {}
+
+// expected-error@+3 {{'std::array<int, 10>' cannot be used as the type of a kernel parameter}}
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
+    (ext::oneapi::experimental::single_task_kernel))
+void ff(std::array<int, 10> a) {}
+
+// expected-error@+3 {{'int &' cannot be used as the type of a kernel parameter}}
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
+    (ext::oneapi::experimental::single_task_kernel))
+void ff(int &ip) {}