Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
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
18 changes: 8 additions & 10 deletions clang/include/clang/Basic/Attr.td
Original file line number Diff line number Diff line change
Expand Up @@ -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];
Copy link
Contributor

Choose a reason for hiding this comment

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

Since a SYCL compilation does both, would using this attribute always result in an attribute ignored warning for the host?

Copy link
Member Author

@sarnex sarnex Oct 10, 2025

Choose a reason for hiding this comment

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

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.

Copy link
Contributor

Choose a reason for hiding this comment

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

You should just require SYCL and probably ignore it on the CPU, I think that's how CUDA's __global__ works.

Copy link
Member Author

Choose a reason for hiding this comment

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

Let me try that

Copy link
Contributor

Choose a reason for hiding this comment

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

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

Copy link
Member Author

Choose a reason for hiding this comment

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

@tahonermann Should we keep throwing the warning on the host? We can easily ignore it and not warn if we want to.

Copy link
Contributor

Choose a reason for hiding this comment

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

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.

Copy link
Member Author

@sarnex sarnex Oct 10, 2025

Choose a reason for hiding this comment

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

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.

Copy link
Member Author

Choose a reason for hiding this comment

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

@tahonermann Okay, then this PR should do that as-is.

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];
Expand All @@ -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,
Expand Down
22 changes: 17 additions & 5 deletions clang/include/clang/Basic/AttrDocs.td
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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";
Expand Down
2 changes: 1 addition & 1 deletion clang/lib/Sema/SemaDecl.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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.
Expand Down
14 changes: 4 additions & 10 deletions clang/lib/Sema/SemaDeclAttr.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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.
Expand Down Expand Up @@ -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;
Expand Down
2 changes: 1 addition & 1 deletion clang/lib/Sema/SemaSYCL.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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) {
Expand Down
8 changes: 4 additions & 4 deletions clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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()));
}

Expand Down Expand Up @@ -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;
}

Expand Down
10 changes: 4 additions & 6 deletions clang/lib/Sema/SemaType.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down