Skip to content

[SYCL] Generate warning about kernel argument size on all devices #2509

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
Sep 22, 2020
Merged

[SYCL] Generate warning about kernel argument size on all devices #2509

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
b2e5f3f
[SYCL] Generate warning about kernel argument size on all devices
elizabethandrews Sep 22, 2020
65b716c
Implemented review comments:
elizabethandrews Sep 22, 2020
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
4 changes: 2 additions & 2 deletions clang/include/clang/Basic/DiagnosticSemaKinds.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -11004,8 +11004,8 @@ def err_sycl_restrict : Error<
"nor constant-initialized"
"}0">;
def warn_sycl_kernel_too_big_args : Warning<
"size of kernel arguments (%0 bytes) exceeds supported maximum of %1 bytes "
"on GPU">, InGroup<SyclStrict>;
"size of kernel arguments (%0 bytes) may exceed the supported maximum "
"of %1 bytes on some devices">, InGroup<SyclStrict>;
def err_sycl_virtual_types : Error<
"No class with a vtable can be used in a SYCL kernel or any code included in the kernel">;
def note_sycl_recursive_function_declared_here: Note<"function implemented using recursion declared here">;
Expand Down
10 changes: 4 additions & 6 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -58,7 +58,7 @@ enum KernelInvocationKind {

const static std::string InitMethodName = "__init";
const static std::string FinalizeMethodName = "__finalize";
constexpr unsigned GPUMaxKernelArgsSize = 2048;
constexpr unsigned MaxKernelArgsSize = 2048;

namespace {

Expand Down Expand Up @@ -1697,11 +1697,9 @@ class SyclKernelArgsSizeChecker : public SyclKernelFieldHandler {
: SyclKernelFieldHandler(S), KernelLoc(Loc) {}

~SyclKernelArgsSizeChecker() {
if (SemaRef.Context.getTargetInfo().getTriple().getSubArch() ==
llvm::Triple::SPIRSubArch_gen)
if (SizeOfParams > GPUMaxKernelArgsSize)
SemaRef.Diag(KernelLoc, diag::warn_sycl_kernel_too_big_args)
<< SizeOfParams << GPUMaxKernelArgsSize;
if (SizeOfParams > MaxKernelArgsSize)
SemaRef.Diag(KernelLoc, diag::warn_sycl_kernel_too_big_args)
<< SizeOfParams << MaxKernelArgsSize;
}

bool handleSyclAccessorType(FieldDecl *FD, QualType FieldTy) final {
Expand Down
16 changes: 5 additions & 11 deletions clang/test/SemaSYCL/args-size-overflow.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,7 +1,5 @@
// RUN: %clang_cc1 -fsycl -triple spir64_gen -DGPU -fsycl-is-device -fsyntax-only -verify %s
// RUN: %clang_cc1 -fsycl -triple spir64 -fsycl-is-device -fsyntax-only -verify %s
// RUN: %clang_cc1 -fsycl -triple spir64_gen -Wno-sycl-strict -fsycl-is-device -fsyntax-only -verify %s
// RUN: %clang_cc1 -fsycl -triple spir64_gen -Werror=sycl-strict -DERROR -fsycl-is-device -fsyntax-only -verify %s
// RUN: %clang_cc1 -fsycl -triple spir64 -Werror=sycl-strict -DERROR -fsycl-is-device -fsyntax-only -verify %s

#include "Inputs/sycl.hpp"
class Foo;
Expand All @@ -13,12 +11,10 @@ __attribute__((sycl_kernel)) void kernel(F KernelFunc) {

template <typename Name, typename F>
void parallel_for(F KernelFunc) {
#ifdef GPU
// expected-warning@+6 {{size of kernel arguments (7994 bytes) exceeds supported maximum of 2048 bytes on GPU}}
#elif ERROR
// expected-error@+4 {{size of kernel arguments (7994 bytes) exceeds supported maximum of 2048 bytes on GPU}}
#ifdef ERROR
// expected-error@+4 {{size of kernel arguments (7994 bytes) may exceed the supported maximum of 2048 bytes on some devices}}
#else
// expected-no-diagnostics
// expected-warning@+2 {{size of kernel arguments (7994 bytes) may exceed the supported maximum of 2048 bytes on some devices}}
#endif
kernel<Name>(KernelFunc);
}
Expand All @@ -35,8 +31,6 @@ void use() {
int Array[1991];
} Args;
auto L = [=]() { (void)Args; };
#if defined(GPU) || defined(ERROR)
// expected-note@+2 {{in instantiation of function template specialization 'parallel_for<Foo}}
#endif
// expected-note@+1 {{in instantiation of function template specialization 'parallel_for<Foo}}
parallel_for<Foo>(L);
}