Skip to content

fix merge after readding syclkernel attr #20363

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.

Sign up for GitHub

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

Merged
merged 2 commits into from
Oct 15, 2025
Merged

fix merge after readding syclkernel attr #20363

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
9940ddb
fix merge after readding syclkernel attr
sarnex Oct 14, 2025
9a28bce
remove extra DeviceKernelAttr check
sarnex Oct 15, 2025
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
3 changes: 2 additions & 1 deletion clang/include/clang/AST/GlobalDecl.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -164,7 +164,8 @@ class GlobalDecl {
}

static KernelReferenceKind getDefaultKernelReference(const FunctionDecl *D) {
return (D->hasAttr<DeviceKernelAttr>() || D->getLangOpts().CUDAIsDevice)
return (D->hasAttr<DeviceKernelAttr>() || D->getLangOpts().CUDAIsDevice ||
D->hasAttr<SYCLKernelAttr>())
Copy link
Contributor

Choose a reason for hiding this comment

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

I'm not sure that checking for SYCLKernelAttr is correct here. We didn't check for it prior to llvm/llvm-project#137882 landing.

Copy link
Contributor Author

@sarnex sarnex Oct 15, 2025 •
edited
Loading

Choose a reason for hiding this comment

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

Right, upstream didnt, but we did in syclos before we pulled down that commit, see here https://github.com/intel/llvm/blob/nightly-2025-07-18/clang/include/clang/AST/GlobalDecl.h#L168

Copy link
Contributor

Choose a reason for hiding this comment

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

Thanks, yes, added by commit 3b170e3 in May, the pulldown happened in commit fcd95a9, and that change was still present in the parent commit 6b2006a (from the previous pulldown). Isn't git archeology a joy.

? KernelReferenceKind::Kernel
: KernelReferenceKind::Stub;
}
Expand Down
7 changes: 3 additions & 4 deletions clang/include/clang/Basic/Attr.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1640,7 +1640,6 @@ def CUDAShared : InheritableAttr {
}
def : MutualExclusions<[CUDAConstant, CUDAShared, HIPManaged]>;

<<<<<<< HEAD
def GlobalStorageNonLocalVar : SubsetSubject<Var,
[{S->hasGlobalStorage() &&
!S->isLocalVarDeclOrParm()}],
Expand All @@ -1667,13 +1666,13 @@ def SYCLGlobalVar : InheritableAttr {
let LangOpts = [SYCLIsDevice];
// Only used internally by the SYCL implementation
let Documentation = [Undocumented];
=======
}

def SYCLKernel : InheritableAttr {
let Spellings = [Clang<"sycl_kernel">];
let Subjects = SubjectList<[FunctionTmpl]>;
let LangOpts = [SYCLDevice];
let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
let Documentation = [SYCLKernelDocs];
>>>>>>> 1db148cc946eb95fefd5399766e379fc030eef78
}

def DeviceKernel : DeclOrTypeAttr {
Expand Down
3 changes: 2 additions & 1 deletion clang/lib/AST/Decl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5569,7 +5569,8 @@ FunctionDecl *FunctionDecl::CreateDeserialized(ASTContext &C, GlobalDeclID ID) {

bool FunctionDecl::isReferenceableKernel() const {
return hasAttr<CUDAGlobalAttr>() ||
DeviceKernelAttr::isOpenCLSpelling(getAttr<DeviceKernelAttr>()) || hasAttr<DeviceKernelAttr>();
DeviceKernelAttr::isOpenCLSpelling(getAttr<DeviceKernelAttr>()) ||
hasAttr<DeviceKernelAttr>() || hasAttr<SYCLKernelAttr>();
}

BlockDecl *BlockDecl::Create(ASTContext &C, DeclContext *DC, SourceLocation L) {
Expand Down
2 changes: 1 addition & 1 deletion clang/lib/CodeGen/CodeGenFunction.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -772,7 +772,7 @@ void CodeGenFunction::EmitKernelMetadata(const FunctionDecl *FD,
}

bool IsKernelOrDevice =
FD->hasAttr<DeviceKernelAttr>() || FD->hasAttr<SYCLDeviceAttr>();
FD->hasAttr<SYCLKernelAttr>() || FD->hasAttr<SYCLDeviceAttr>();
const IntelReqdSubGroupSizeAttr *ReqSubGroup =
FD->getAttr<IntelReqdSubGroupSizeAttr>();

Expand Down
2 changes: 1 addition & 1 deletion clang/lib/CodeGen/CodeGenModule.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6953,7 +6953,7 @@ CodeGenModule::getLLVMLinkageForDeclarator(const DeclaratorDecl *D,
// with the SYCL_EXTERNAL macro. For any function or variable that does not
// have this, linkonce_odr suffices. If -fno-sycl-rdc is passed, we know there
// is only one translation unit and can so mark them internal.
if (getLangOpts().SYCLIsDevice && !D->hasAttr<DeviceKernelAttr>() &&
if (getLangOpts().SYCLIsDevice && !D->hasAttr<SYCLKernelAttr>() &&
!D->hasAttr<SYCLDeviceAttr>() && !D->hasAttr<SYCLExternalAttr>() &&
!SemaSYCL::isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
D->getType()))
Expand Down
2 changes: 1 addition & 1 deletion clang/lib/Sema/SemaChecking.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4408,7 +4408,7 @@ void Sema::checkCall(NamedDecl *FDecl, const FunctionProtoType *Proto,
if (FD)
diagnoseArgDependentDiagnoseIfAttrs(FD, ThisArg, Args, Loc);

if (FD && FD->hasAttr<DeviceKernelAttr>() && getLangOpts().isSYCL())
if (FD && FD->hasAttr<SYCLKernelAttr>())
SYCL().CheckSYCLKernelCall(FD, Args);

// Diagnose variadic calls in SYCL.
Expand Down
14 changes: 3 additions & 11 deletions clang/lib/Sema/SemaDecl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -20957,7 +20957,7 @@ Sema::DeviceDiagnosticReason Sema::getEmissionReason(const FunctionDecl *FD) {
// FIXME: This should really be a bitwise-or of the language modes.
if (FD->hasAttr<SYCLSimdAttr>())
return Sema::DeviceDiagnosticReason::Esimd;
if (FD->hasAttr<SYCLDeviceAttr>() || FD->hasAttr<DeviceKernelAttr>())
if (FD->hasAttr<SYCLDeviceAttr>() || FD->hasAttr<SYCLKernelAttr>())
return getLangOpts().SYCLCUDACompat
? Sema::DeviceDiagnosticReason::SyclCudaCompat
: Sema::DeviceDiagnosticReason::Sycl;
Expand All @@ -20978,20 +20978,12 @@ Sema::FunctionEmissionStatus Sema::getEmissionStatus(const FunctionDecl *FD,
bool Final) {
assert(FD && "Expected non-null FunctionDecl");

<<<<<<< HEAD
=======
// 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<SYCLKernelAttr>())
return FunctionEmissionStatus::Emitted;

>>>>>>> 1db148cc946eb95fefd5399766e379fc030eef78
// Templates are emitted when they're instantiated.
if (FD->isDependentContext())
return FunctionEmissionStatus::TemplateDiscarded;

if (LangOpts.SYCLIsDevice &&
(FD->hasAttr<SYCLDeviceAttr>() || FD->hasAttr<DeviceKernelAttr>()) &&
(FD->hasAttr<SYCLDeviceAttr>() || FD->hasAttr<SYCLKernelAttr>()) &&
!FD->hasAttr<ArtificialAttr>())
return FunctionEmissionStatus::Emitted;

Expand Down Expand Up @@ -21077,7 +21069,7 @@ Sema::FunctionEmissionStatus Sema::getEmissionStatus(const FunctionDecl *FD,
}

if (getLangOpts().SYCLIsDevice) {
if (!FD->hasAttr<SYCLDeviceAttr>() && !FD->hasAttr<DeviceKernelAttr>())
if (!FD->hasAttr<SYCLDeviceAttr>() && !FD->hasAttr<SYCLKernelAttr>())
return FunctionEmissionStatus::Unknown;

// Check whether this function is externally visible -- if so, it's
Expand Down
17 changes: 1 addition & 16 deletions clang/lib/Sema/SemaDeclAttr.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5266,20 +5266,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();
<<<<<<< HEAD
if (S.getLangOpts().isSYCL()) {
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()) {
>>>>>>> 1db148cc946eb95fefd5399766e379fc030eef78
handleGlobalAttr(S, D, AL);
Comment on lines 5269 to 5270
Copy link
Contributor

Choose a reason for hiding this comment

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

I'm not sure this is right. Prior to llvm/llvm-project#137882, processing of NVPTXKernelAttr and CUDAGlobalAttr were not dependent on an NVPTX target. I'm not sure if this has ramifications for the SYCL-CUDA compat mode.

Copy link
Contributor Author

@sarnex sarnex Oct 15, 2025 •
edited
Loading

Choose a reason for hiding this comment

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

For NVPTXKernelAttr, it looks like it is target dependent in the attr definition:

https://github.com/llvm/llvm-project/blob/87f0227cb60147a26a1eeb4fb06e3b505e9c7261/clang/include/clang/Basic/Attr.td#L1470

def NVPTXKernel : InheritableAttr, **TargetSpecificAttr<TargetNVPTX>** {

so if the target doesn't match the the attr is never added, the check for the target agains the attr is in ParsedAttr, before we create that NVPTXKernel object

and this function can't be hit for CUDAGlobalAttr because it's guarded by a check for DeviceKernelAttr.

so I think the cases handleGlobalAttr sees NVPTXKernelAttr or CUDAGlobalAttr are the exact same before and after

Copy link
Contributor

Choose a reason for hiding this comment

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

Ok, I think I agree with that, but shouldn't there also be a check for isNVPTXSpelling()? If the spelling is "__kernel" or "kernel", I would expect the else branch to be the desired control flow.

Copy link
Contributor Author

@sarnex sarnex Oct 15, 2025 •
edited
Loading

Choose a reason for hiding this comment

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

upstream Joseph was really against target-specific behavior, and even if seems confusing he would prefer for all the remaining spellings of DeviceKernelAttr to be aliases, expect for OpenCL because it's very particular, and if we do make a change like that i would expect it to be upstream

Copy link
Contributor

Choose a reason for hiding this comment

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

I understand that he would like to have one attribute for kernels that are strongly tied to a particular target (though I don't necessarily agree with him). With this code as is though, wouldn't use of the OpenCL spellings end up going down the wrong branch when targeting NVPTX?

Copy link
Contributor Author

@sarnex sarnex Oct 15, 2025 •
edited
Loading

Choose a reason for hiding this comment

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

Yeah I think you're right, but it doesn't seem to be causing any problems either upstream or here.
Is it okay if we move forward with this PR in syclos and I make a PR upstream to add a check so we don't go down this branch for OpenCL spellings?

Copy link
Contributor

Choose a reason for hiding this comment

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

Sure, I think that is reasonable.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

llvm/llvm-project#163859

} else {
// OpenCL C++ will throw a more specific error.
Expand Down Expand Up @@ -7184,13 +7171,11 @@ ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, const ParsedAttr &AL,
case ParsedAttr::AT_EnumExtensibility:
handleEnumExtensibilityAttr(S, D, AL);
break;
<<<<<<< HEAD
case ParsedAttr::AT_SYCLSimd:
handleSimpleAttribute<SYCLSimdAttr>(S, D, AL);
=======
break;
case ParsedAttr::AT_SYCLKernel:
S.SYCL().handleKernelAttr(D, AL);
>>>>>>> 1db148cc946eb95fefd5399766e379fc030eef78
break;
case ParsedAttr::AT_SYCLExternal:
handleSimpleAttribute<SYCLExternalAttr>(S, D, AL);
Expand Down
26 changes: 4 additions & 22 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -911,7 +911,7 @@ class SingleDeviceFunctionTracker {
// a SYCLKernel or SYCLDevice attribute on it, add it to the set of
// routines potentially reachable on device. This is to diagnose such
// cases later in finalizeSYCLDelayedAnalysis().
if (!CurrentDecl->isDefined() && !CurrentDecl->hasAttr<DeviceKernelAttr>() &&
if (!CurrentDecl->isDefined() && !CurrentDecl->hasAttr<SYCLKernelAttr>() &&
!CurrentDecl->hasAttr<SYCLDeviceAttr>())
Parent.SemaSYCLRef.addFDToReachableFromSyclDevice(CurrentDecl,
CallStack.back());
Expand Down Expand Up @@ -971,7 +971,7 @@ class SingleDeviceFunctionTracker {
if (isSYCLKernelBodyFunction(CurrentDecl)) {
// This is a direct callee of the kernel.
if (CallStack.size() == 1 &&
CallStack.back()->hasAttr<DeviceKernelAttr>()) {
CallStack.back()->hasAttr<SYCLKernelAttr>()) {
assert(!KernelBody && "inconsistent call graph - only one kernel body "
"function can be called");
KernelBody = CurrentDecl;
Expand Down Expand Up @@ -3009,7 +3009,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
// to TransformStmt in replaceWithLocalClone can diagnose something that got
// diagnosed on the actual kernel.
KernelDecl->addAttr(
DeviceKernelAttr::CreateImplicit(SemaSYCLRef.getASTContext()));
SYCLKernelAttr::CreateImplicit(SemaSYCLRef.getASTContext()));

SemaSYCLRef.addSyclDeviceDecl(KernelDecl);
}
Expand Down Expand Up @@ -6093,7 +6093,7 @@ void SemaSYCL::finalizeSYCLDelayedAnalysis(const FunctionDecl *Caller,
return;

// If Callee has a SYCL attribute, no diagnostic needed.
if (Callee->hasAttr<SYCLDeviceAttr>() || Callee->hasAttr<DeviceKernelAttr>())
if (Callee->hasAttr<SYCLDeviceAttr>() || Callee->hasAttr<SYCLKernelAttr>())
return;

// If Callee has a CUDA device attribute, no diagnostic needed.
Expand Down Expand Up @@ -7716,24 +7716,6 @@ void SemaSYCL::performSYCLDelayedAttributesAnalaysis(const FunctionDecl *FD) {
diag::warn_sycl_incorrect_use_attribute_non_kernel_function)
<< KernelAttr;
}
<<<<<<< HEAD
=======

// Function must have at least one argument.
if (getFunctionOrMethodNumParams(D) != 1) {
Diag(FT->getLocation(), diag::warn_sycl_kernel_num_of_function_params);
return;
}

// Function must return void.
QualType RetTy = getFunctionOrMethodResultType(D);
if (!RetTy->isVoidType()) {
Diag(FT->getLocation(), diag::warn_sycl_kernel_return_type);
return;
}

handleSimpleAttribute<SYCLKernelAttr>(*this, D, AL);
>>>>>>> 1db148cc946eb95fefd5399766e379fc030eef78
}

void SemaSYCL::handleKernelEntryPointAttr(Decl *D, const ParsedAttr &AL) {
Expand Down
2 changes: 1 addition & 1 deletion clang/lib/Sema/SemaSYCLDeclAttr.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -55,7 +55,7 @@ void SemaSYCL::handleKernelAttr(Decl *D, const ParsedAttr &AL) {
return;
}

handleSimpleAttribute<DeviceKernelAttr>(*this, D, AL);
handleSimpleAttribute<SYCLKernelAttr>(*this, D, AL);
}

// Returns a DupArgResult value; Same means the args have the same value,
Expand Down
2 changes: 1 addition & 1 deletion clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7659,7 +7659,7 @@ static void processFunctionInstantiation(Sema &S,
DefinitionRequired, AtEndOfTU);
if (!FD->isDefined())
return;
if (S.LangOpts.SYCLIsDevice && FD->hasAttr<DeviceKernelAttr>())
if (S.LangOpts.SYCLIsDevice && FD->hasAttr<SYCLKernelAttr>())
S.SYCL().ConstructOpenCLKernel(FD, MC);
FD->setInstantiationIsPending(false);
}
Expand Down
5 changes: 0 additions & 5 deletions clang/lib/Sema/SemaType.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3797,12 +3797,7 @@ static CallingConv getCCForDeclaratorChunk(
}
}
for (const ParsedAttr &AL : D.getDeclSpec().getAttributes()) {
<<<<<<< HEAD
if (AL.getKind() == ParsedAttr::AT_DeviceKernel &&
!DeviceKernelAttr::isSYCLSpelling(AL)) {
=======
if (AL.getKind() == ParsedAttr::AT_DeviceKernel) {
>>>>>>> 1db148cc946eb95fefd5399766e379fc030eef78
CC = CC_DeviceKernel;
break;
}
Expand Down
4 changes: 2 additions & 2 deletions clang/test/SemaSYCL/sycl-esimd-ast.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -14,15 +14,15 @@ int main() {
deviceQueue.submit([&](sycl::handler &h) {
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel1
// CHECK: SYCLSimdAttr {{.*}} Implicit
// CHECK-NEXT: DeviceKernelAttr {{.*}} Implicit
// CHECK-NEXT: SYCLKernelAttr {{.*}} Implicit
// CHECK-NEXT: AsmLabelAttr {{.*}} Implicit
// CHECK-NEXT: SYCLSimdAttr {{.*}}
h.single_task<class test_kernel1>(
FuncObj());

// CHECK-LABEL: FunctionDecl {{.*}}test_kernel2
// CHECK: SYCLSimdAttr {{.*}} Implicit
// CHECK-NEXT: DeviceKernelAttr {{.*}} Implicit
// CHECK-NEXT: SYCLKernelAttr {{.*}} Implicit
// CHECK-NEXT: AsmLabelAttr {{.*}} Implicit
// CHECK-NEXT: SYCLSimdAttr {{.*}}
h.single_task<class test_kernel2>(
Expand Down