Skip to content
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

Jump to bottom

[SYCL] Support or diagnose use of namespace std types as kernel type … #1579

Closed
wants to merge 1 commit into from
Closed

[SYCL] Support or diagnose use of namespace std types as kernel type … #1579

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
45 changes: 22 additions & 23 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1649,30 +1649,27 @@ void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D,
auto *NS = dyn_cast_or_null<NamespaceDecl>(DC);

if (!NS) {
if (!DC->isTranslationUnit()) {
const TagDecl *TD = isa<ClassTemplateDecl>(D)
? cast<ClassTemplateDecl>(D)->getTemplatedDecl()
: dyn_cast<TagDecl>(D);

if (TD && !UnnamedLambdaSupport) {
// defined class constituting the kernel name is not globally
// accessible - contradicts the spec
const bool KernelNameIsMissing = TD->getName().empty();
if (KernelNameIsMissing) {
const TagDecl *TD = isa<ClassTemplateDecl>(D)
Copy link
Contributor

Choose a reason for hiding this comment

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

I wonder if this needs to be lower-level than this, since would this work with 'int'? Or the other builtin types? I don't think they are tagdecls, but they are probably named decls.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

For the fundamental types and other builtin types, we don't even get here. We enter here only for those types for which a forward declaration is necessary.

Copy link
Contributor

Choose a reason for hiding this comment

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

I see. @kbobrovs is most familiar with the Int header stuff, so I'll leave the review to him.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Thanks @erichkeane

? cast<ClassTemplateDecl>(D)->getTemplatedDecl()
: dyn_cast<TagDecl>(D);
if (!TD)
break;

const bool KernelNameIsMissing = TD->getName().empty();
if (KernelNameIsMissing)
Diag.Report(KernelLocation, diag::err_sycl_kernel_incorrectly_named)
<< /* kernel name is missing */ 0;
else if (!DC->isTranslationUnit()) {
// defined class constituting the kernel name is not globally
// accessible - contradicts the spec
if (!UnnamedLambdaSupport) {
if (TD->isCompleteDefinition())
Diag.Report(KernelLocation, diag::err_sycl_kernel_incorrectly_named)
<< /* kernel name is missing */ 0;
// Don't emit note if kernel name was completely omitted
} else {
if (TD->isCompleteDefinition())
Diag.Report(KernelLocation,
diag::err_sycl_kernel_incorrectly_named)
<< /* kernel name is not globally-visible */ 1;
else
Diag.Report(KernelLocation, diag::warn_sycl_implicit_decl);
Diag.Report(D->getSourceRange().getBegin(),
diag::note_previous_decl)
<< TD->getName();
}
<< /* kernel name is not globally-visible */ 1;
else
Diag.Report(KernelLocation, diag::warn_sycl_implicit_decl);
Diag.Report(D->getSourceRange().getBegin(), diag::note_previous_decl)
<< TD->getName();
}
}
break;
Expand Down Expand Up @@ -1826,8 +1823,10 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
O << "// This is auto-generated SYCL integration header.\n";
O << "\n";

O << "#include <cstddef>\n";
Copy link
Contributor

Choose a reason for hiding this comment

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

I don't think we need these changes to the generated header (this and the next change). Because there is no promise that std types work as lambda name.

Copy link
Contributor

Choose a reason for hiding this comment

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

@premanandrao, could you address this comment, please?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

@bader , @jtmott-intel is working on it.

O << "#include <CL/sycl/detail/defines.hpp>\n";
O << "#include <CL/sycl/detail/kernel_desc.hpp>\n";
O << "using nullptr_t = std::nullptr_t;\n";

O << "\n";

Expand Down
63 changes: 63 additions & 0 deletions clang/test/CodeGenSYCL/stdtypes_kernel_type.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,63 @@
// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-int-header=%t.h %s
// RUN: FileCheck -input-file=%t.h %s
//
// CHECK: #include <cstddef>
// CHECK-NEXT: #include <CL/sycl/detail/defines.hpp>
// CHECK-NEXT: #include <CL/sycl/detail/kernel_desc.hpp>
// CHECK-NEXT: using nullptr_t = std::nullptr_t;
//
// CHECK: static constexpr
// CHECK-NEXT: const char* const kernel_names[] = {
// CHECK-NEXT: "_ZTSDn"
// CHECK-NEXT: "_ZTSSt4byte"
// CHECK-NEXT: "_ZTSm",
// CHECK-NEXT: "_ZTSl"
// CHECK-NEXT: };
//
// CHECK: static constexpr
// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
// CHECK-NEXT: //--- _ZTSDn
// CHECK-EMPTY:
// CHECK-NEXT: //--- _ZTSSt4byte
// CHECK-EMPTY:
// CHECK-NEXT: //--- _ZTSm
// CHECK-EMPTY:
// CHECK-NEXT: //--- _ZTSl
// CHECK-EMPTY:
// CHECK-NEXT: };
//
// CHECK: static constexpr
// CHECK-NEXT: const unsigned kernel_signature_start[] = {
// CHECK-NEXT: 0, // _ZTSDn
// CHECK-NEXT: 1, // _ZTSSt4byte
// CHECK-NEXT: 2, // _ZTSm
// CHECK-NEXT: 3 // _ZTSl
// CHECK-NEXT: };

// CHECK: template <> struct KernelInfo<nullptr_t> {
// CHECK: template <> struct KernelInfo<::std::byte> {
// CHECK: template <> struct KernelInfo<unsigned long> {
// CHECK: template <> struct KernelInfo<long> {

void usage() {
}

namespace std {
typedef long unsigned int size_t;
typedef long int ptrdiff_t;
typedef decltype(nullptr) nullptr_t;
enum class byte : unsigned char {};
} // namespace std

template <typename name, typename Func>
__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
kernelFunc();
}

int main() {
kernel_single_task<std::nullptr_t>([]() { usage(); });
kernel_single_task<std::byte>([=]() {});
kernel_single_task<std::size_t>([=]() {});
kernel_single_task<std::ptrdiff_t>([=]() {});
return 0;
}
23 changes: 18 additions & 5 deletions clang/test/SemaSYCL/unnamed-kernel.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,16 +1,22 @@
// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -verify %s
// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsycl-unnamed-lambda -fsyntax-only -verify %s
#include <sycl.hpp>

#ifdef __SYCL_UNNAMED_LAMBDA__
// expected-no-diagnostics
#endif

#include <sycl.hpp>

namespace namespace1 {
template <typename T>
class KernelName;
}

namespace std {
typedef struct {
} max_align_t;
} // namespace std

struct MyWrapper {
private:
class InvalidKernelName0 {};
Expand Down Expand Up @@ -41,15 +47,15 @@ struct MyWrapper {

#ifndef __SYCL_UNNAMED_LAMBDA__
// expected-error@+4 {{kernel needs to have a globally-visible name}}
// expected-note@16 {{InvalidKernelName0 declared here}}
// expected-note@21 {{InvalidKernelName0 declared here}}
#endif
q.submit([&](cl::sycl::handler &h) {
h.single_task<InvalidKernelName0>([] {});
});

#ifndef __SYCL_UNNAMED_LAMBDA__
// expected-error@+4 {{kernel needs to have a globally-visible name}}
// expected-note@17 {{InvalidKernelName3 declared here}}
// expected-note@22 {{InvalidKernelName3 declared here}}
#endif
q.submit([&](cl::sycl::handler &h) {
h.single_task<namespace1::KernelName<InvalidKernelName3>>([] {});
Expand All @@ -60,10 +66,17 @@ struct MyWrapper {
h.single_task<ValidAlias>([] {});
});

#ifndef __SYCL_UNNAMED_LAMBDA__
// expected-error@+3 {{kernel name is missing}}
#endif
q.submit([&](cl::sycl::handler &h) {
h.single_task<std::max_align_t>([] {});
});

using InvalidAlias = InvalidKernelName4;
#ifndef __SYCL_UNNAMED_LAMBDA__
// expected-error@+4 {{kernel needs to have a globally-visible name}}
// expected-note@18 {{InvalidKernelName4 declared here}}
// expected-note@23 {{InvalidKernelName4 declared here}}
#endif
q.submit([&](cl::sycl::handler &h) {
h.single_task<InvalidAlias>([] {});
Expand All @@ -72,7 +85,7 @@ struct MyWrapper {
using InvalidAlias1 = InvalidKernelName5;
#ifndef __SYCL_UNNAMED_LAMBDA__
// expected-error@+4 {{kernel needs to have a globally-visible name}}
// expected-note@19 {{InvalidKernelName5 declared here}}
// expected-note@24 {{InvalidKernelName5 declared here}}
#endif
q.submit([&](cl::sycl::handler &h) {
h.single_task<namespace1::KernelName<InvalidAlias1>>([] {});
Expand Down