From 6bbe958c19153b5b2583f3ad7ae396f7184dd1f2 Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Fri, 9 May 2025 13:52:30 +0200 Subject: [PATCH 1/3] [SYCL] the first kernel function declaration should be added with attribute --- .../clang/Basic/DiagnosticSemaKinds.td | 3 +- clang/lib/Sema/SemaSYCL.cpp | 34 ++++++-- .../test/SemaSYCL/free_function_negative.cpp | 16 +++- .../free_functions/redeclaration.cpp | 82 +++++++++++++++++++ 4 files changed, 126 insertions(+), 9 deletions(-) create mode 100644 sycl/test-e2e/Experimental/free_functions/redeclaration.cpp diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 4594c55382b1..dd20f4dcbabc 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12676,7 +12676,8 @@ def err_free_function_variadic_args : Error< "free function kernel cannot be a variadic function">; def err_free_function_return_type : Error< "SYCL free function kernel should have return type 'void'">; - +def err_free_function_first_occurrence_missing_attr: Error< + "the first occurrence of kernel free function should be declared with attribute">; // SYCL kernel entry point diagnostics def err_sycl_entry_point_invalid : Error< diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index ab0864aaf41f..7d0d3a0ef8c2 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1160,15 +1160,35 @@ static target getAccessTarget(QualType FieldTy, } bool SemaSYCL::isFreeFunction(const FunctionDecl *FD) { - for (auto *IRAttr : FD->specific_attrs()) { - SmallVector, 4> NameValuePairs = - IRAttr->getAttributeNameValuePairs(getASTContext()); - for (const auto &NameValuePair : NameValuePairs) { - if (NameValuePair.first == "sycl-nd-range-kernel" || - NameValuePair.first == "sycl-single-task-kernel") { - return true; + llvm::SmallVector Redecls; + for (clang::FunctionDecl *Redecl : FD->redecls()) { + Redecls.push_back(Redecl); + } + bool FirstDecl = true; + clang::SourceLocation Loc = FD->getLocation(); + while (!Redecls.empty()) { + auto *Redecl = Redecls.back(); + Redecls.pop_back(); + if (FirstDecl) + Loc = Redecl->getLocation(); // Save the location of the first decl to use + // in diagnostics. + for (auto *IRAttr : + Redecl->specific_attrs()) { + SmallVector, 4> NameValuePairs = + IRAttr->getAttributeNameValuePairs(getASTContext()); + for (const auto &NameValuePair : NameValuePairs) { + if (NameValuePair.first == "sycl-nd-range-kernel" || + NameValuePair.first == "sycl-single-task-kernel") { + if (FirstDecl) + return true; + else { + Diag(Loc, diag::err_free_function_first_occurrence_missing_attr); + return false; + } + } } } + FirstDecl = false; } return false; } diff --git a/clang/test/SemaSYCL/free_function_negative.cpp b/clang/test/SemaSYCL/free_function_negative.cpp index 3c8e5a07851b..7223befc8af7 100644 --- a/clang/test/SemaSYCL/free_function_negative.cpp +++ b/clang/test/SemaSYCL/free_function_negative.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -verify=expected -fsycl-int-header=%t.h %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -verify=expected -verify-ignore-unexpected=note -fsycl-int-header=%t.h %s #include "sycl.hpp" @@ -10,6 +10,20 @@ foo(int start, ...) { // expected-error {{free function kernel cannot be a varia foo1(int start, ...) { // expected-error {{free function kernel cannot be a variadic function}} } +[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 1)]] void +foo2(int start); + +// expected-error@+1 {{attribute 'add_ir_attributes_function' is already applied with different arguments}} +[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] void +foo2(int start) { +} + +// expected-error@+1 {{the first occurrence of kernel free function should be declared with attribute}} +void foo3(int start, int *ptr); + +[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] void +foo3(int start, int *ptr){} + // expected-error@+2 {{a function with a default argument value cannot be used to define SYCL free function kernel}} [[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] void singleTaskKernelDefaultValues(int Value = 1) { diff --git a/sycl/test-e2e/Experimental/free_functions/redeclaration.cpp b/sycl/test-e2e/Experimental/free_functions/redeclaration.cpp new file mode 100644 index 000000000000..3e4d9c66cd04 --- /dev/null +++ b/sycl/test-e2e/Experimental/free_functions/redeclaration.cpp @@ -0,0 +1,82 @@ +// REQUIRES: aspect-usm_shared_allocations +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +// The name mangling for free function kernels currently does not work with PTX. +// UNSUPPORTED: cuda, hip +// UNSUPPORTED-INTENDED: Not implemented yet for Nvidia/AMD backends. + +#include +#include +#include +#include +#include + +namespace syclext = sycl::ext::oneapi; +namespace syclexp = sycl::ext::oneapi::experimental; + +static constexpr size_t NUM = 1024; +static constexpr size_t WGSIZE = 16; + +int check_result(int *ptr) { + for (size_t i = 0; i < NUM; ++i) { + const float expected = 3 + static_cast(i); + if (ptr[i] != expected) { + std::cout << "Kernel execution did not produce the expected result\n"; + return 1; + } + } + return 0; +} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void free_func(int start, int *ptr); + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void free_func(int start, int *ptr); + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void free_func1(int start, int *ptr); + +void free_func1(int start, int *ptr); + +static int call_kernel_code(sycl::queue &q, sycl::kernel &kernel) { + int *ptr = sycl::malloc_shared(NUM, q); + q.submit([&](sycl::handler &cgh) { + cgh.set_args(3, ptr); + sycl::nd_range ndr{{NUM}, {WGSIZE}}; + cgh.parallel_for(ndr, kernel); + }).wait(); + const int ret = check_result(ptr); + sycl::free(ptr, q); + return ret; +} + +template +int test_declarations(sycl::queue &q, sycl::context &ctxt) { + auto exe_bndl = + syclexp::get_kernel_bundle(ctxt); + sycl::kernel k_func = exe_bndl.template ext_oneapi_get_kernel(); + return call_kernel_code(q, k_func); +} + +#define KERNEL_CODE(start, ptr) \ + size_t id = \ + syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); \ + ptr[id] = start + static_cast(id); + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void free_func(int start, int *ptr) { KERNEL_CODE(start, ptr); } + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void free_func1(int start, int *ptr) { KERNEL_CODE(start, ptr); } + +int main() { + sycl::queue q; + sycl::context ctxt = q.get_context(); + + int result{0}; + result |= test_declarations(q, ctxt); + result |= test_declarations(q, ctxt); + return result; +} From 403ef068fb47d9d8a87a270d0e174dce8f7db05a Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Wed, 14 May 2025 14:06:46 +0200 Subject: [PATCH 2/3] [SYCL] rework method to detect if kernel free function --- .../clang/Basic/DiagnosticSemaKinds.td | 2 +- clang/lib/Sema/SemaSYCL.cpp | 43 +++++++++---------- .../test/SemaSYCL/free_function_negative.cpp | 6 ++- 3 files changed, 26 insertions(+), 25 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index dd20f4dcbabc..2439a1592f05 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12677,7 +12677,7 @@ def err_free_function_variadic_args : Error< def err_free_function_return_type : Error< "SYCL free function kernel should have return type 'void'">; def err_free_function_first_occurrence_missing_attr: Error< - "the first occurrence of kernel free function should be declared with attribute">; + "the first occurrence of kernel free function should be declared with attribute add_ir_attributes_function with 'sycl-nd-range-kernel' or 'sycl-single-task-kernel'">; // SYCL kernel entry point diagnostics def err_sycl_entry_point_invalid : Error< diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 7d0d3a0ef8c2..0298f0089c9e 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1160,35 +1160,34 @@ static target getAccessTarget(QualType FieldTy, } bool SemaSYCL::isFreeFunction(const FunctionDecl *FD) { - llvm::SmallVector Redecls; - for (clang::FunctionDecl *Redecl : FD->redecls()) { - Redecls.push_back(Redecl); - } - bool FirstDecl = true; clang::SourceLocation Loc = FD->getLocation(); - while (!Redecls.empty()) { - auto *Redecl = Redecls.back(); - Redecls.pop_back(); - if (FirstDecl) - Loc = Redecl->getLocation(); // Save the location of the first decl to use - // in diagnostics. + bool NextDeclaredWithAttr = false; + for (clang::FunctionDecl *Redecl : FD->redecls()) { + bool IsFreeFunctionAttr = false; for (auto *IRAttr : Redecl->specific_attrs()) { SmallVector, 4> NameValuePairs = IRAttr->getAttributeNameValuePairs(getASTContext()); - for (const auto &NameValuePair : NameValuePairs) { - if (NameValuePair.first == "sycl-nd-range-kernel" || - NameValuePair.first == "sycl-single-task-kernel") { - if (FirstDecl) - return true; - else { - Diag(Loc, diag::err_free_function_first_occurrence_missing_attr); - return false; - } - } + const auto it = std::find_if( + NameValuePairs.begin(), NameValuePairs.end(), + [](const auto &NameValuePair) { + return NameValuePair.first == "sycl-nd-range-kernel" || + NameValuePair.first == "sycl-single-task-kernel"; + }); + IsFreeFunctionAttr = it != NameValuePairs.end(); + } + if (Redecl->isFirstDecl()) { + if (IsFreeFunctionAttr) + return true; + if (NextDeclaredWithAttr) { + Diag(Loc, diag::err_free_function_first_occurrence_missing_attr); + Diag(Redecl->getLocation(), diag::note_previous_declaration); + return false; } + } else { + Loc = Redecl->getLocation(); + NextDeclaredWithAttr = IsFreeFunctionAttr; } - FirstDecl = false; } return false; } diff --git a/clang/test/SemaSYCL/free_function_negative.cpp b/clang/test/SemaSYCL/free_function_negative.cpp index 7223befc8af7..23d811c6b4ca 100644 --- a/clang/test/SemaSYCL/free_function_negative.cpp +++ b/clang/test/SemaSYCL/free_function_negative.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -verify=expected -verify-ignore-unexpected=note -fsycl-int-header=%t.h %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -verify=expected -fsycl-int-header=%t.h %s #include "sycl.hpp" @@ -10,6 +10,7 @@ foo(int start, ...) { // expected-error {{free function kernel cannot be a varia foo1(int start, ...) { // expected-error {{free function kernel cannot be a variadic function}} } +// expected-note@+1 {{conflicting attribute is here}} [[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 1)]] void foo2(int start); @@ -18,9 +19,10 @@ foo2(int start); foo2(int start) { } -// expected-error@+1 {{the first occurrence of kernel free function should be declared with attribute}} +// expected-note@+1 {{previous declaration is here}} void foo3(int start, int *ptr); +// expected-error@+2 {{the first occurrence of kernel free function should be declared with attribute add_ir_attributes_function with 'sycl-nd-range-kernel' or 'sycl-single-task-kernel'}} [[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] void foo3(int start, int *ptr){} From 9ba920037c265818f87cad80571c96976e6ac905 Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Thu, 15 May 2025 14:06:20 +0200 Subject: [PATCH 3/3] [SYCL] add more tests of redeclarations --- .../clang/Basic/DiagnosticSemaKinds.td | 2 +- clang/lib/Sema/SemaSYCL.cpp | 4 +- .../test/SemaSYCL/free_function_negative.cpp | 16 ++++- .../free_functions/redeclaration.cpp | 68 ++++++++++++------- 4 files changed, 62 insertions(+), 28 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 2439a1592f05..9886b77325d7 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12677,7 +12677,7 @@ def err_free_function_variadic_args : Error< def err_free_function_return_type : Error< "SYCL free function kernel should have return type 'void'">; def err_free_function_first_occurrence_missing_attr: Error< - "the first occurrence of kernel free function should be declared with attribute add_ir_attributes_function with 'sycl-nd-range-kernel' or 'sycl-single-task-kernel'">; + "the first occurrence of SYCL kernel free function should be declared with 'sycl-nd-range-kernel' or 'sycl-single-task-kernel' compile time properties">; // SYCL kernel entry point diagnostics def err_sycl_entry_point_invalid : Error< diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 0298f0089c9e..06de71700b92 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1160,9 +1160,9 @@ static target getAccessTarget(QualType FieldTy, } bool SemaSYCL::isFreeFunction(const FunctionDecl *FD) { - clang::SourceLocation Loc = FD->getLocation(); + SourceLocation Loc = FD->getLocation(); bool NextDeclaredWithAttr = false; - for (clang::FunctionDecl *Redecl : FD->redecls()) { + for (FunctionDecl *Redecl : FD->redecls()) { bool IsFreeFunctionAttr = false; for (auto *IRAttr : Redecl->specific_attrs()) { diff --git a/clang/test/SemaSYCL/free_function_negative.cpp b/clang/test/SemaSYCL/free_function_negative.cpp index 23d811c6b4ca..639259f853dd 100644 --- a/clang/test/SemaSYCL/free_function_negative.cpp +++ b/clang/test/SemaSYCL/free_function_negative.cpp @@ -22,10 +22,24 @@ foo2(int start) { // expected-note@+1 {{previous declaration is here}} void foo3(int start, int *ptr); -// expected-error@+2 {{the first occurrence of kernel free function should be declared with attribute add_ir_attributes_function with 'sycl-nd-range-kernel' or 'sycl-single-task-kernel'}} +// expected-error@+2 {{the first occurrence of SYCL kernel free function should be declared with 'sycl-nd-range-kernel' or 'sycl-single-task-kernel' compile time properties}} [[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] void foo3(int start, int *ptr){} +// expected-note@+1 {{previous declaration is here}} +void foo4(float start, float *ptr); + +// expected-error@+2 {{the first occurrence of SYCL kernel free function should be declared with 'sycl-nd-range-kernel' or 'sycl-single-task-kernel' compile time properties}} +[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] void +foo4(float start, float *ptr); + +[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] void +foo4(float start, float *ptr); + +[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] void +foo4(float start, float *ptr){} + + // expected-error@+2 {{a function with a default argument value cannot be used to define SYCL free function kernel}} [[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] void singleTaskKernelDefaultValues(int Value = 1) { diff --git a/sycl/test-e2e/Experimental/free_functions/redeclaration.cpp b/sycl/test-e2e/Experimental/free_functions/redeclaration.cpp index 3e4d9c66cd04..6f6f37c94339 100644 --- a/sycl/test-e2e/Experimental/free_functions/redeclaration.cpp +++ b/sycl/test-e2e/Experimental/free_functions/redeclaration.cpp @@ -18,9 +18,9 @@ namespace syclexp = sycl::ext::oneapi::experimental; static constexpr size_t NUM = 1024; static constexpr size_t WGSIZE = 16; -int check_result(int *ptr) { +template int check_result(T *ptr, T value) { for (size_t i = 0; i < NUM; ++i) { - const float expected = 3 + static_cast(i); + const T expected = value + static_cast(i); if (ptr[i] != expected) { std::cout << "Kernel execution did not produce the expected result\n"; return 1; @@ -30,53 +30,73 @@ int check_result(int *ptr) { } SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) -void free_func(int start, int *ptr); +void free_func(int *ptr, int start); SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) -void free_func(int start, int *ptr); +void free_func(int *ptr, int start); SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) -void free_func1(int start, int *ptr); +void free_func1(int *ptr, int start); -void free_func1(int start, int *ptr); +void free_func1(int *ptr, int start); -static int call_kernel_code(sycl::queue &q, sycl::kernel &kernel) { - int *ptr = sycl::malloc_shared(NUM, q); +template +static int call_kernel_code(sycl::queue &q, sycl::kernel &kernel, T value) { + T *ptr = sycl::malloc_shared(NUM, q); q.submit([&](sycl::handler &cgh) { - cgh.set_args(3, ptr); + if (value == 0) + cgh.set_args(ptr); + else + cgh.set_args(ptr, value); sycl::nd_range ndr{{NUM}, {WGSIZE}}; cgh.parallel_for(ndr, kernel); }).wait(); - const int ret = check_result(ptr); + const int ret = check_result(ptr, value); sycl::free(ptr, q); return ret; } -template -int test_declarations(sycl::queue &q, sycl::context &ctxt) { - auto exe_bndl = - syclexp::get_kernel_bundle(ctxt); - sycl::kernel k_func = exe_bndl.template ext_oneapi_get_kernel(); - return call_kernel_code(q, k_func); -} - -#define KERNEL_CODE(start, ptr) \ +#define KERNEL_CODE(start, ptr, type) \ size_t id = \ syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); \ - ptr[id] = start + static_cast(id); + ptr[id] = static_cast(start) + static_cast(id); + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void free_func(int *ptr, int start) { KERNEL_CODE(start, ptr, int); } + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void free_func1(int *ptr, int start) { KERNEL_CODE(start, ptr, int); } + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void free_func2(int *ptr, int start) { KERNEL_CODE(start, ptr, int); } SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) -void free_func(int start, int *ptr) { KERNEL_CODE(start, ptr); } +void free_func2(float *ptr, float start) { KERNEL_CODE(start, ptr, float); } SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) -void free_func1(int start, int *ptr) { KERNEL_CODE(start, ptr); } +void free_func2(int *ptr) { KERNEL_CODE(0, ptr, int); } + +template +int test_declarations(sycl::queue &q, sycl::context &ctxt, T value) { + auto exe_bndl = + syclexp::get_kernel_bundle(ctxt); + sycl::kernel k_func = exe_bndl.template ext_oneapi_get_kernel(); + return call_kernel_code(q, k_func, value); +} int main() { sycl::queue q; sycl::context ctxt = q.get_context(); int result{0}; - result |= test_declarations(q, ctxt); - result |= test_declarations(q, ctxt); + result |= test_declarations(q, ctxt, 3); + result |= test_declarations(q, ctxt, 3); + result |= + test_declarations(free_func2), int>( + q, ctxt, 3); + result |= test_declarations(free_func2), + float>(q, ctxt, 3.14f); + result |= test_declarations(free_func2), int>( + q, ctxt, 0); return result; }