Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

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

Merged
merged 57 commits into from
Aug 14, 2024
Merged
Show file tree
Hide file tree
Changes from 55 commits
Commits
Show all changes
57 commits
Select commit Hold shift + click to select a range
efb93a5
Support for free function traits.
rdeodhar May 23, 2024
2cd0d1a
Updated tests.
rdeodhar Jun 11, 2024
9003bf0
Merge branch 'sycl' of https://github.com/intel/llvm into freefunc_tr…
rdeodhar Jun 11, 2024
c5c2540
Updated a test.
rdeodhar Jun 11, 2024
318c528
Merge branch 'sycl' of https://github.com/intel/llvm into freefunc_tr…
rdeodhar Jun 12, 2024
c091fca
Added header file for clang testing.
rdeodhar Jun 12, 2024
d78a903
Deferred HIP support to a future PR.
rdeodhar Jun 12, 2024
a0ed967
Correction to test.
rdeodhar Jun 13, 2024
26b7b09
Experimental test change.
rdeodhar Jun 13, 2024
fae8dd3
Enable HIP.
rdeodhar Jun 13, 2024
44ef070
Add debug print.
rdeodhar Jun 13, 2024
51f7072
Adjusted a test to add USM requirement.
rdeodhar Jun 14, 2024
89ef3dc
Removed need for SYCL_EXTERNAL on free functions.
rdeodhar Jun 17, 2024
9c0b8f0
Simplifications mostly, and support for extern "C".
rdeodhar Jun 20, 2024
42478c8
Remove debug code.
rdeodhar Jun 20, 2024
a926383
Changes to namespaces.
rdeodhar Jun 21, 2024
be9c177
Merge branch 'sycl' of https://github.com/intel/llvm into freefunc_tr…
rdeodhar Jun 24, 2024
e9f50cc
Removed some unneeded code and unneeded namespace qualifiers.
rdeodhar Jun 24, 2024
2dc28f7
Merge branch 'sycl' of https://github.com/intel/llvm into freefunc_tr…
rdeodhar Jun 24, 2024
1e1aa2c
Fixes for templated free functions.
rdeodhar Jun 27, 2024
0a5569b
Removed debug code.
rdeodhar Jun 27, 2024
07ebf0d
Merge branch 'sycl' of https://github.com/intel/llvm into freefunc_tr…
rdeodhar Jun 28, 2024
be6c168
Removed unneeded code.
rdeodhar Jun 28, 2024
ef8b031
Merge branch 'freefunc_traits' of https://github.com/rdeodhar/llvm in…
rdeodhar Jul 1, 2024
40fb924
Removed need to specify SYCL_EXTERNAL.
rdeodhar Jul 8, 2024
bbdf983
Merge branch 'sycl' of https://github.com/intel/llvm into freefunc3
rdeodhar Jul 8, 2024
38ad179
Merge branch 'sycl' of https://github.com/intel/llvm into freefunc3
rdeodhar Jul 9, 2024
85fe22a
Merge branch 'sycl' of https://github.com/intel/llvm into freefunc_tr…
rdeodhar Jul 9, 2024
46d3e5d
Merge branch 'sycl' of https://github.com/intel/llvm into freefunc_tr…
rdeodhar Jul 10, 2024
66665af
Merge branch 'sycl' of https://github.com/intel/llvm into freefunc3
rdeodhar Jul 12, 2024
9f4eded
Changed templated forward declaration for shim.
rdeodhar Jul 12, 2024
b8e1665
Correction to forward declaration of templated free functions.
rdeodhar Jul 12, 2024
670cc61
Merge branch 'freefunc_traits' of https://github.com/rdeodhar/llvm in…
rdeodhar Jul 12, 2024
f753c37
Changed the way non-templated functions are forward declared.
rdeodhar Jul 15, 2024
36ba16a
Merge branch 'sycl' of https://github.com/intel/llvm into freefunc_tr…
rdeodhar Jul 17, 2024
7e6aacf
Merge branch 'sycl' of https://github.com/intel/llvm into freefunc_tr…
rdeodhar Jul 17, 2024
c57f9cd
Corrected a test.
rdeodhar Jul 17, 2024
7b0dcf9
Minor change to API test.
rdeodhar Jul 22, 2024
a0f4ddc
Merge branch 'sycl' of https://github.com/intel/llvm into freefunc_tr…
rdeodhar Jul 22, 2024
06f91cf
Merge branch 'sycl' of https://github.com/intel/llvm into freefunc3
rdeodhar Jul 22, 2024
b83a080
Merge branch 'freefunc_traits' of https://github.com/rdeodhar/llvm in…
rdeodhar Jul 23, 2024
cbb3f23
Correction to where free functions are placed.
rdeodhar Jul 29, 2024
db857f2
Removed "device" marking from test.
rdeodhar Jul 29, 2024
0475c35
Merge branch 'sycl' of https://github.com/intel/llvm into freefunc3
rdeodhar Jul 29, 2024
e39f44a
Merge branch 'sycl' of https://github.com/intel/llvm into freefunc3
rdeodhar Jul 30, 2024
2a90137
Simplified checking of free function attributes.
rdeodhar Jul 31, 2024
74b5d82
Merge branch 'sycl' of https://github.com/intel/llvm into freefunc3
rdeodhar Jul 31, 2024
6d7a31a
Added test for free function parameter errors.
rdeodhar Aug 1, 2024
e5edf4c
Merge branch 'sycl' of https://github.com/intel/llvm into freefunc3
rdeodhar Aug 2, 2024
e0f5b82
Added test to ensure missing SYCL_EXTERN on free function does not ge…
rdeodhar Aug 6, 2024
5603aef
Fixed comment.
rdeodhar Aug 6, 2024
d803209
Exclude member functions as free function candidates and add a test.
rdeodhar Aug 9, 2024
3b71b2c
Added add_ir_atribute_check test.
rdeodhar Aug 10, 2024
af52693
Comment correction.
rdeodhar Aug 10, 2024
9d1a87e
Merge branch 'sycl' of https://github.com/intel/llvm into freefunc3
rdeodhar Aug 10, 2024
2d2ef56
Simplification.
rdeodhar Aug 13, 2024
7189420
Updated a comment.
rdeodhar Aug 14, 2024
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
14 changes: 9 additions & 5 deletions clang/lib/Sema/SemaDecl.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -16318,11 +16318,15 @@ 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))
SYCL().ProcessFreeFunction(FD);
if (FD) {
// Free functions cannot be member functions so skip those.
const auto *MD = dyn_cast<CXXMethodDecl>(FD);
if (LangOpts.SYCLIsDevice && !MD && Body &&
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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).

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done.

(FD->getTemplatedKind() == FunctionDecl::TK_NonTemplate ||
FD->getTemplatedKind() ==
FunctionDecl::TK_FunctionTemplateSpecialization))
SYCL().ProcessFreeFunction(FD);
}

return dcl;
}
Expand Down
20 changes: 16 additions & 4 deletions clang/lib/Sema/SemaSYCL.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -1073,12 +1073,20 @@ static target getAccessTarget(QualType FieldTy,
AccTy->getTemplateArgs()[3].getAsIntegral().getExtValue());
}

// FIXME: Free functions must have void return type, be declared at file scope,
AaronBallman marked this conversation as resolved.
Show resolved Hide resolved
// 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.
auto hasDependentExpr = [](auto Args) {
return llvm::any_of(Args, [](const Expr *E) {
return E->isValueDependent() || E->isTypeDependent();
});
};

// 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>()) {
// Free function properties are all compiletime constants, so skip checking
// any attribute values that use dependent expressions.
if (hasDependentExpr(IRAttr->args()))
continue;
SmallVector<std::pair<std::string, std::string>, 4> NameValuePairs =
IRAttr->getAttributeNameValuePairs(SemaSYCLRef.getASTContext());
for (const auto &NameValuePair : NameValuePairs) {
Expand All @@ -1100,6 +1108,10 @@ static bool isFreeFunction(SemaSYCL &SemaSYCLRef, const FunctionDecl *FD) {
static int getFreeFunctionRangeDim(SemaSYCL &SemaSYCLRef,
const FunctionDecl *FD) {
for (auto *IRAttr : FD->specific_attrs<SYCLAddIRAttributesFunctionAttr>()) {
// Free function properties are all compiletime constants, so skip checking
// any attribute values that use dependent expressions.
if (hasDependentExpr(IRAttr->args()))
continue;
SmallVector<std::pair<std::string, std::string>, 4> NameValuePairs =
IRAttr->getAttributeNameValuePairs(SemaSYCLRef.getASTContext());
for (const auto &NameValuePair : NameValuePairs) {
Expand Down
3 changes: 0 additions & 3 deletions clang/test/CodeGenSYCL/free_function_int_header.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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) {
Expand All @@ -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) {
Expand All @@ -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++)
Expand Down
59 changes: 59 additions & 0 deletions clang/test/Sema/free_function_attribute_check.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,59 @@
//==---- free_function_attribute_check.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 checks that non-constexpr values appearing in
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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"

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

OK, I've updated the comment.

// add_ir_attributes_function can be handled when checking whether a function is
// a free function. When all values in the attribute can be handled, then we can
// safely test any candidate function for being a free function.

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);
6 changes: 3 additions & 3 deletions sycl/test-e2e/KernelAndProgram/free_function_apis.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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) {
Expand All @@ -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();
Expand Down
9 changes: 4 additions & 5 deletions sycl/test-e2e/KernelAndProgram/free_function_kernels.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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) {
Expand Down Expand Up @@ -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) {
Expand Down Expand Up @@ -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) {
Expand Down Expand Up @@ -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();
Expand Down
51 changes: 51 additions & 0 deletions sycl/test/extensions/free_function_errors.cpp
Original file line number Diff line number Diff line change
@@ -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) {}
Loading