[clang][Sema] Split SYCLKernel back into its own attribute #162868
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Signed-off-by: Sarnie, Nick <[email protected]>
llvmbot
added
clang
|
@llvm/pr-subscribers-clang Author: Nick Sarnie (sarnex) ChangesBased on feedback from #161905, partially revert #137882 so that Full diff: https://github.com/llvm/llvm-project/pull/162868.diff 7 Files Affected:
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 3cde249e286fa..22e60aa9fe312 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1599,8 +1599,15 @@ def CUDAShared : InheritableAttr {
}
def : MutualExclusions<[CUDAConstant, CUDAShared, HIPManaged]>;
+def SYCLKernel : InheritableAttr {
+ let Spellings = [Clang<"sycl_kernel">];
+ let Subjects = SubjectList<[FunctionTmpl]>;
+ let LangOpts = [SYCLDevice];
+ let Documentation = [SYCLKernelDocs];
+}
+
def DeviceKernel : DeclOrTypeAttr {
- let Spellings = [Clang<"device_kernel">, Clang<"sycl_kernel">,
+ let Spellings = [Clang<"device_kernel">,
Clang<"nvptx_kernel">, Clang<"amdgpu_kernel">,
CustomKeyword<"__kernel">, CustomKeyword<"kernel">];
let Documentation = [DeviceKernelDocs];
@@ -1624,15 +1631,6 @@ def DeviceKernel : DeclOrTypeAttr {
if(!A) return false;
return isNVPTXSpelling(*A);
}
- static inline bool isSYCLSpelling(const AttributeCommonInfo& A) {
- return A.getAttributeSpellingListIndex() == GNU_sycl_kernel ||
- A.getAttributeSpellingListIndex() == CXX11_clang_sycl_kernel ||
- A.getAttributeSpellingListIndex() == C23_clang_sycl_kernel;
- }
- static inline bool isSYCLSpelling(const AttributeCommonInfo* A) {
- if(!A) return false;
- return isSYCLSpelling(*A);
- }
static inline bool isOpenCLSpelling(const AttributeCommonInfo& A) {
// Tablegen trips underscores from spellings to build the spelling
// list, but here we have the same spelling with unscores and without,
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index 20a52b49a8f10..e0bbda083b5cf 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -396,13 +396,10 @@ any option of a multiversioned function is undefined.
}];
}
-def DeviceKernelDocs : Documentation {
+def SYCLKernelDocs : Documentation {
let Category = DocCatFunction;
- let Heading = "device_kernel, sycl_kernel, nvptx_kernel, amdgpu_kernel, "
- "kernel, __kernel";
+ let Heading = "sycl_kernel";
let Content = [{
-These attributes specify that the function represents a kernel for device offloading.
-The specific semantics depend on the offloading language, target, and attribute spelling.
The ``sycl_kernel`` attribute specifies that a function template will be used
to outline device code and to generate an OpenCL kernel.
Here is a code example of the SYCL program, which demonstrates the compiler's
@@ -476,6 +473,21 @@ The SYCL kernel in the previous code sample meets these expectations.
}];
}
+def DeviceKernelDocs : Documentation {
+ let Category = DocCatFunction;
+ let Heading = "device_kernel, nvptx_kernel, amdgpu_kernel, "
+ "kernel, __kernel";
+ let Content = [{
+These attributes specify that the function represents a kernel for device offloading.
+The specific semantics depend on the offloading language, target, and attribute spelling.
+Here is a code example using the attribute to mark a function as a kernel:
+
+.. code-block:: c++
+
+ [[clang::device_kernel]] int foo(int x) { return ++x; }
+ }];
+}
+
def SYCLExternalDocs : Documentation {
let Category = DocCatFunction;
let Heading = "sycl_external";
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 0e83c20b27c22..8ac09c4d30f1a 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -20797,7 +20797,7 @@ Sema::FunctionEmissionStatus Sema::getEmissionStatus(const FunctionDecl *FD,
// SYCL functions can be template, so we check if they have appropriate
// attribute prior to checking if it is a template.
- if (LangOpts.SYCLIsDevice && FD->hasAttr<DeviceKernelAttr>())
+ if (LangOpts.SYCLIsDevice && FD->hasAttr<SYCLKernelAttr>())
return FunctionEmissionStatus::Emitted;
// Templates are emitted when they're instantiated.
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 328ccf6694073..3107876565e8e 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -5204,16 +5204,7 @@ static void handleCallConvAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
static void handleDeviceKernelAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
const auto *FD = dyn_cast_or_null<FunctionDecl>(D);
bool IsFunctionTemplate = FD && FD->getDescribedFunctionTemplate();
- if (S.getLangOpts().SYCLIsDevice) {
- if (!IsFunctionTemplate) {
- S.Diag(AL.getLoc(), diag::warn_attribute_wrong_decl_type_str)
- << AL << AL.isRegularKeywordAttribute() << "function templates";
- } else {
- S.SYCL().handleKernelAttr(D, AL);
- }
- } else if (DeviceKernelAttr::isSYCLSpelling(AL)) {
- S.Diag(AL.getLoc(), diag::warn_attribute_ignored) << AL;
- } else if (S.getASTContext().getTargetInfo().getTriple().isNVPTX()) {
+ if (S.getASTContext().getTargetInfo().getTriple().isNVPTX()) {
handleGlobalAttr(S, D, AL);
} else {
// OpenCL C++ will throw a more specific error.
@@ -7100,6 +7091,9 @@ ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, const ParsedAttr &AL,
case ParsedAttr::AT_EnumExtensibility:
handleEnumExtensibilityAttr(S, D, AL);
break;
+ case ParsedAttr::AT_SYCLKernel:
+ S.SYCL().handleKernelAttr(D, AL);
+ break;
case ParsedAttr::AT_SYCLExternal:
handleSimpleAttribute<SYCLExternalAttr>(S, D, AL);
break;
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 2f97f6290f0e8..b981c35c8083f 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -199,7 +199,7 @@ void SemaSYCL::handleKernelAttr(Decl *D, const ParsedAttr &AL) {
return;
}
- handleSimpleAttribute<DeviceKernelAttr>(*this, D, AL);
+ handleSimpleAttribute<SYCLKernelAttr>(*this, D, AL);
}
void SemaSYCL::handleKernelEntryPointAttr(Decl *D, const ParsedAttr &AL) {
diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
index 3819f775811e5..85e3d207b2cf2 100644
--- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
@@ -710,9 +710,9 @@ static void instantiateDependentAMDGPUMaxNumWorkGroupsAttr(
// This doesn't take any template parameters, but we have a custom action that
// needs to happen when the kernel itself is instantiated. We need to run the
// ItaniumMangler to mark the names required to name this kernel.
-static void instantiateDependentDeviceKernelAttr(
+static void instantiateDependentSYCLKernelAttr(
Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,
- const DeviceKernelAttr &Attr, Decl *New) {
+ const SYCLKernelAttr &Attr, Decl *New) {
New->addAttr(Attr.clone(S.getASTContext()));
}
@@ -966,8 +966,8 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs,
continue;
}
- if (auto *A = dyn_cast<DeviceKernelAttr>(TmplAttr)) {
- instantiateDependentDeviceKernelAttr(*this, TemplateArgs, *A, New);
+ if (auto *A = dyn_cast<SYCLKernelAttr>(TmplAttr)) {
+ instantiateDependentSYCLKernelAttr(*this, TemplateArgs, *A, New);
continue;
}
diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp
index bee613aa5f1c5..a9e7c34de94f4 100644
--- a/clang/lib/Sema/SemaType.cpp
+++ b/clang/lib/Sema/SemaType.cpp
@@ -3780,12 +3780,10 @@ static CallingConv getCCForDeclaratorChunk(
}
}
}
- if (!S.getLangOpts().isSYCL()) {
- for (const ParsedAttr &AL : D.getDeclSpec().getAttributes()) {
- if (AL.getKind() == ParsedAttr::AT_DeviceKernel) {
- CC = CC_DeviceKernel;
- break;
- }
+ for (const ParsedAttr &AL : D.getDeclSpec().getAttributes()) {
+ if (AL.getKind() == ParsedAttr::AT_DeviceKernel) {
+ CC = CC_DeviceKernel;
+ break;
}
}
return CC;
|
jhuber6
reviewed
Oct 10, 2025
| def SYCLKernel : InheritableAttr { | ||
| let Spellings = [Clang<"sycl_kernel">]; | ||
| let Subjects = SubjectList<[FunctionTmpl]>; | ||
| let LangOpts = [SYCLDevice]; |
There was a problem hiding this comment.
Since a SYCL compilation does both, would using this attribute always result in an attribute ignored warning for the host?
There was a problem hiding this comment.
Just tried and it seems yes:
"/llvm/build/bin/clang-22" -cc1 -triple x86_64-unknown-linux-gnu -fsycl-is-host -sycl-std=2020 -emit-obj -dumpdir a- -disable-free -clear-ast-before-backend -main-file-name foo.cpp -mrelocation-model pic -pic-level 2 -pic-is-pie -mframe-pointer=all -fmath-errno -ffp-contract=on -fno-rounding-math -mconstructor-aliases -funwind-tables=2 -target-cpu x86-64 -tune-cpu generic -debugger-tuning=gdb -v -resource-dir /llvm/build/lib/clang/22 -internal-isystem /usr/lib/gcc/x86_64-redhat-linux/11/../../../../include/c++/11 -internal-isystem /usr/lib/gcc/x86_64-redhat-linux/11/../../../../include/c++/11/x86_64-redhat-linux -internal-isystem /usr/lib/gcc/x86_64-redhat-linux/11/../../../../include/c++/11/backward -internal-isystem /usr/lib/gcc/x86_64-redhat-linux/11/../../../../include/c++/11 -internal-isystem /usr/lib/gcc/x86_64-redhat-linux/11/../../../../include/c++/11/x86_64-redhat-linux -internal-isystem /usr/lib/gcc/x86_64-redhat-linux/11/../../../../include/c++/11/backward -internal-isystem /llvm/build/lib/clang/22/include -internal-isystem /usr/local/include -internal-isystem /usr/lib/gcc/x86_64-redhat-linux/11/../../../../x86_64-redhat-linux/include -internal-externc-isystem /include -internal-externc-isystem /usr/include -internal-isystem /llvm/build/lib/clang/22/include -internal-isystem /usr/local/include -internal-isystem /usr/lib/gcc/x86_64-redhat-linux/11/../../../../x86_64-redhat-linux/include -internal-externc-isystem /include -internal-externc-isystem /usr/include -fdeprecated-macro -ferror-limit 19 -fmessage-length=262 -fgnuc-version=4.2.1 -fskip-odr-check-in-gmf -fcxx-exceptions -fexceptions -fcolor-diagnostics -fembed-offload-object=/tmp/foo-b73420.out -faddrsig -D__GCC_HAVE_DWARF2_CFI_ASM=1 -o /tmp/foo-053756.o -x c++ foo.cpp
clang -cc1 version 22.0.0git based upon LLVM 22.0.0git default target x86_64-unknown-linux-gnu
ignoring nonexistent directory "/usr/lib/gcc/x86_64-redhat-linux/11/../../../../x86_64-redhat-linux/include"
ignoring nonexistent directory "/include"
ignoring nonexistent directory "/usr/lib/gcc/x86_64-redhat-linux/11/../../../../x86_64-redhat-linux/include"
ignoring nonexistent directory "/include"
ignoring duplicate directory "/usr/lib/gcc/x86_64-redhat-linux/11/../../../../include/c++/11"
ignoring duplicate directory "/usr/lib/gcc/x86_64-redhat-linux/11/../../../../include/c++/11/x86_64-redhat-linux"
ignoring duplicate directory "/usr/lib/gcc/x86_64-redhat-linux/11/../../../../include/c++/11/backward"
ignoring duplicate directory "/llvm/build/lib/clang/22/include"
ignoring duplicate directory "/usr/local/include"
ignoring duplicate directory "/usr/include"
#include "..." search starts here:
#include <...> search starts here:
/usr/lib/gcc/x86_64-redhat-linux/11/../../../../include/c++/11
/usr/lib/gcc/x86_64-redhat-linux/11/../../../../include/c++/11/x86_64-redhat-linux
/usr/lib/gcc/x86_64-redhat-linux/11/../../../../include/c++/11/backward
/llvm/build/lib/clang/22/include
/usr/local/include
/usr/include
End of search list.
foo.cpp:2:3: warning: 'clang::sycl_kernel' attribute ignored [-Wignored-attributes]
2 | [[clang::sycl_kernel]]
No warning for the device compile.
In our downstream fork we handle cases like this by using an ifdef.
There was a problem hiding this comment.
You should just require SYCL and probably ignore it on the CPU, I think that's how CUDA's __global__ works.
There was a problem hiding this comment.
@jhuber6, yes, the attribute intentionally has no effect for host compilation and is expected to be hidden behind a macro for device compilation. We are designing differently for the sycl_kernel_entry_point attribute that is intended to replace sycl_kernel.
There was a problem hiding this comment.
@tahonermann Should we keep throwing the warning on the host? We can easily ignore it and not warn if we want to.
There was a problem hiding this comment.
Please don't assume that what makes sense for __global__, __kernel, etc... makes sense for sycl_kernel. The goal for this PR should be to put sycl_kernel back the way it was before #137882. Don't try to improve it.
There was a problem hiding this comment.
Also about the test, yes we have a few tests, SemaSYCL/kernel-attribute-on-non-sycl.cpp and SemaSYCL/kernel-attribute.cpp that are thorough.
There was a problem hiding this comment.
@tahonermann Okay, then this PR should do that as-is.
erichkeane
approved these changes
Oct 10, 2025
tahonermann
approved these changes
Oct 11, 2025
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
5 participants
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Based on feedback from #161905, partially revert #137882 so that
sycl_kernelis a separate attribute and not just a spelling ofdevice_kernel.