Skip to content

Commit

Permalink
Create isForwardDeclarable
Browse files Browse the repository at this point in the history
  • Loading branch information
@Fznamznon
Fznamznon committed Sep 23, 2024
1 parent 0b88367 commit ce98062
Show file tree
Hide file tree
Showing 2 changed files with 109 additions and 99 deletions.
206 changes: 108 additions & 98 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -489,6 +489,84 @@ void SemaSYCL::checkSYCLDeviceVarDecl(VarDecl *Var) {
checkSYCLType(*this, Ty, Loc, Visited);
}

enum NotForwardDeclarableReason {
UnscopedEnum,
StdNamespace,
UnnamedTag,
NotAtNamespaceScope,
None
};

// This is a helper function which is used to check if a class declaration is:
// * declared within namespace 'std' (at any level)
// e.g., namespace std { namespace literals { class Whatever; } }
// h.single_task<std::literals::Whatever>([]() {});
// * declared within a function
// e.g., void foo() { struct S { int i; };
// h.single_task<S>([]() {}); }
// * declared within another tag
// e.g., struct S { struct T { int i } t; };
// h.single_task<S::T>([]() {});
// User for kernel name types and class/struct types used in free function
// kernel arguments.
static NotForwardDeclarableReason
isForwardDeclarable(const NamedDecl *DeclToCheck, SemaSYCL &S,
bool DiagForFreeFunction = false) {
if (const auto *ED = dyn_cast<EnumDecl>(DeclToCheck);
ED && !ED->isScoped() && !ED->isFixed())
return NotForwardDeclarableReason::UnscopedEnum;

const DeclContext *DeclCtx = DeclToCheck->getDeclContext();
if (DeclCtx) {
while (!DeclCtx->isTranslationUnit() &&
(isa<NamespaceDecl>(DeclCtx) || isa<LinkageSpecDecl>(DeclCtx))) {
const auto *NSDecl = dyn_cast<NamespaceDecl>(DeclCtx);
// We don't report free function kernel parameter case because the
// restriction for the type used there to be forward declarable comes from
// the need to forward declare it in the integration header. We're safe
// to do so because the integration header is an implemention detail and
// is generated by the compiler.
// We do diagnose case with kernel name type since the spec requires us to
// do so.
if (!DiagForFreeFunction && NSDecl && NSDecl->isStdNamespace())
return NotForwardDeclarableReason::StdNamespace;
DeclCtx = DeclCtx->getParent();
}
}

// Check if the we've met a Tag declaration local to a non-namespace scope
// (i.e. Inside a function or within another Tag etc).
if (const auto *Tag = dyn_cast<TagDecl>(DeclToCheck)) {
if (Tag->getIdentifier() == nullptr)
return NotForwardDeclarableReason::UnnamedTag;
if (!DeclCtx->isTranslationUnit()) {
// Diagnose used types without complete definition i.e.
// int main() {
// class KernelName1;
// parallel_for<class KernelName1>(..);
// }
// For kernel name type This case can only be diagnosed during host
// compilation because the integration header is required to distinguish
// between the invalid code (above) and the following valid code:
// int main() {
// parallel_for<class KernelName2>(..);
// }
// The device compiler forward declares both KernelName1 and
// KernelName2 in the integration header as ::KernelName1 and
// ::KernelName2. The problem with the former case is the additional
// declaration 'class KernelName1' in non-global scope. Lookup in this
// case will resolve to ::main::KernelName1 (instead of
// ::KernelName1). Since this is not visible to runtime code that
// submits kernels, this is invalid.
if (Tag->isCompleteDefinition() ||
S.getLangOpts().SYCLEnableIntHeaderDiags || DiagForFreeFunction)
return NotForwardDeclarableReason::NotAtNamespaceScope;
}
}

return NotForwardDeclarableReason::None;
}

// Tests whether given function is a lambda function or '()' operator used as
// SYCL kernel body function (e.g. in parallel_for).
// NOTE: This is incomplete implemenation. See TODO in the FE TODO list for the
Expand Down Expand Up @@ -1964,25 +2042,9 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
// For free functions all struct/class kernel arguments are forward declared
// in integration header, that adds additional restrictions for kernel
// arguments.
// Lambdas are not forward declarable. So, diagnose them properly.
if (RD->isLambda()) {
Diag.Report(PD->getLocation(),
diag::err_bad_free_function_kernel_param_type)
<< ParamTy;
Diag.Report(PD->getLocation(),
diag::note_free_function_kernel_param_type_not_fwd_declarable)
<< ParamTy;
IsInvalid = true;
return isValid();
}

// Check that the type is defined at namespace scope.
const DeclContext *DeclCtx = RD->getDeclContext();
while (!DeclCtx->isTranslationUnit() &&
(isa<NamespaceDecl>(DeclCtx) || isa<LinkageSpecDecl>(DeclCtx)))
DeclCtx = DeclCtx->getParent();

if (!DeclCtx->isTranslationUnit()) {
NotForwardDeclarableReason NFDR =
isForwardDeclarable(RD, SemaSYCLRef, /*DiagForFreeFunction=*/true);
if (NFDR != NotForwardDeclarableReason::None) {
Diag.Report(PD->getLocation(),
diag::err_bad_free_function_kernel_param_type)
<< ParamTy;
Expand Down Expand Up @@ -4875,89 +4937,37 @@ class SYCLKernelNameTypeVisitor
}

void DiagnoseKernelNameType(const NamedDecl *DeclNamed) {
/*
This is a helper function which throws an error if the kernel name
declaration is:
* declared within namespace 'std' (at any level)
e.g., namespace std { namespace literals { class Whatever; } }
h.single_task<std::literals::Whatever>([]() {});
* declared within a function
e.g., void foo() { struct S { int i; };
h.single_task<S>([]() {}); }
* declared within another tag
e.g., struct S { struct T { int i } t; };
h.single_task<S::T>([]() {});
*/

if (const auto *ED = dyn_cast<EnumDecl>(DeclNamed)) {
if (!ED->isScoped() && !ED->isFixed()) {
if (!IsUnnamedKernel) {
NotForwardDeclarableReason NFDR = isForwardDeclarable(DeclNamed, S);
switch (NFDR) {
case NotForwardDeclarableReason::UnscopedEnum:
S.Diag(KernelInvocationFuncLoc, diag::err_sycl_kernel_incorrectly_named)
<< /* unscoped enum requires fixed underlying type */ 1
<< DeclNamed;
IsInvalid = true;
}
}

const DeclContext *DeclCtx = DeclNamed->getDeclContext();
if (DeclCtx && !IsUnnamedKernel) {

// Check if the kernel name declaration is declared within namespace
// "std" (at any level).
while (!DeclCtx->isTranslationUnit() && isa<NamespaceDecl>(DeclCtx)) {
const auto *NSDecl = cast<NamespaceDecl>(DeclCtx);
if (NSDecl->isStdNamespace()) {
S.Diag(KernelInvocationFuncLoc,
diag::err_invalid_std_type_in_sycl_kernel)
<< KernelNameType << DeclNamed;
IsInvalid = true;
return;
}
DeclCtx = DeclCtx->getParent();
}

// Check if the kernel name is a Tag declaration
// local to a non-namespace scope (i.e. Inside a function or within
// another Tag etc).
if (!DeclCtx->isTranslationUnit() && !isa<NamespaceDecl>(DeclCtx)) {
if (const auto *Tag = dyn_cast<TagDecl>(DeclNamed)) {
bool UnnamedLambdaUsed = Tag->getIdentifier() == nullptr;

if (UnnamedLambdaUsed) {
S.Diag(KernelInvocationFuncLoc,
diag::err_sycl_kernel_incorrectly_named)
<< /* unnamed type is invalid */ 2 << KernelNameType;
IsInvalid = true;
return;
}

// Diagnose used types without complete definition i.e.
// int main() {
// class KernelName1;
// parallel_for<class KernelName1>(..);
// }
// This case can only be diagnosed during host compilation because the
// integration header is required to distinguish between the invalid
// code (above) and the following valid code:
// int main() {
// parallel_for<class KernelName2>(..);
// }
// The device compiler forward declares both KernelName1 and
// KernelName2 in the integration header as ::KernelName1 and
// ::KernelName2. The problem with the former case is the additional
// declaration 'class KernelName1' in non-global scope. Lookup in this
// case will resolve to ::main::KernelName1 (instead of
// ::KernelName1). Since this is not visible to runtime code that
// submits kernels, this is invalid.
if (Tag->isCompleteDefinition() ||
S.getLangOpts().SYCLEnableIntHeaderDiags) {
S.Diag(KernelInvocationFuncLoc,
diag::err_sycl_kernel_incorrectly_named)
<< /* kernel name should be forward declarable at namespace
scope */
0 << KernelNameType;
IsInvalid = true;
}
}
return;
case NotForwardDeclarableReason::StdNamespace:
S.Diag(KernelInvocationFuncLoc,
diag::err_invalid_std_type_in_sycl_kernel)
<< KernelNameType << DeclNamed;
IsInvalid = true;
return;
case NotForwardDeclarableReason::UnnamedTag:
S.Diag(KernelInvocationFuncLoc, diag::err_sycl_kernel_incorrectly_named)
<< /* unnamed type is invalid */ 2 << KernelNameType;
IsInvalid = true;
return;
case NotForwardDeclarableReason::NotAtNamespaceScope:
S.Diag(KernelInvocationFuncLoc, diag::err_sycl_kernel_incorrectly_named)
<< /* kernel name should be forward declarable at namespace
scope */
0 << KernelNameType;
IsInvalid = true;
return;
case NotForwardDeclarableReason::None:
default:
// Do nothing, we're fine.
break;
}
}
}
Expand Down
2 changes: 1 addition & 1 deletion clang/test/SemaSYCL/kernelname-enum.cpp
View file
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -sycl-std=2020 -verify %s
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -sycl-std=2020 -fno-sycl-unnamed-lambda -verify %s

// This test verifies that kernel names containing unscoped enums are diagnosed correctly.

Expand Down

0 comments on commit ce98062

Please sign in to comment.