Skip to content

[SYCL] Add template parameter support for no_global_work_offset attribute #2839

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.

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 13 commits into from
Dec 17, 2020
Merged
2 changes: 1 addition & 1 deletion clang/include/clang/Basic/Attr.td
Original file line number Diff line number Diff line change
Expand Up @@ -1303,7 +1303,7 @@ def SYCLIntelMaxGlobalWorkDim : InheritableAttr {
def SYCLIntelNoGlobalWorkOffset : InheritableAttr {
let Spellings = [CXX11<"intelfpga","no_global_work_offset">,
CXX11<"intel","no_global_work_offset">];
let Args = [BoolArgument<"Enabled", 1>];
let Args = [ExprArgument<"Value", /*optional*/1>];
let LangOpts = [SYCLIsDevice, SYCLIsHost];
let Subjects = SubjectList<[Function], ErrorDiag>;
let Documentation = [SYCLIntelNoGlobalWorkOffsetAttrDocs];
Expand Down
4 changes: 1 addition & 3 deletions clang/include/clang/Basic/DiagnosticGroups.td
Original file line number Diff line number Diff line change
Expand Up @@ -667,10 +667,8 @@ def NSReturnsMismatch : DiagGroup<"nsreturns-mismatch">;
def IndependentClassAttribute : DiagGroup<"IndependentClass-attribute">;
def UnknownAttributes : DiagGroup<"unknown-attributes">;
def IgnoredAttributes : DiagGroup<"ignored-attributes">;
def AdjustedAttributes : DiagGroup<"adjusted-attributes">;
def Attributes : DiagGroup<"attributes", [UnknownAttributes,
IgnoredAttributes,
AdjustedAttributes]>;
IgnoredAttributes]>;
def UnknownSanitizers : DiagGroup<"unknown-sanitizers">;
def UnnamedTypeTemplateArgs : DiagGroup<"unnamed-type-template-args",
[CXX98CompatUnnamedTypeTemplateArgs]>;
Expand Down
3 changes: 0 additions & 3 deletions clang/include/clang/Basic/DiagnosticSemaKinds.td
Original file line number Diff line number Diff line change
Expand Up @@ -11134,9 +11134,6 @@ def err_sycl_function_attribute_mismatch : Error<
"SYCL kernel without %0 attribute can't call a function with this attribute">;
def err_sycl_x_y_z_arguments_must_be_one : Error<
"%0 X-, Y- and Z- sizes must be 1 when %1 attribute is used with value 0">;
def warn_boolean_attribute_argument_is_not_valid: Warning<
"The value of %0 attribute should be 0 or 1. Adjusted to 1">,
InGroup<AdjustedAttributes>;
def err_sycl_attibute_cannot_be_applied_here
: Error<"%0 attribute cannot be applied to a "
"static function or function in an anonymous namespace">;
Expand Down
7 changes: 6 additions & 1 deletion clang/lib/CodeGen/CodeGenFunction.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -694,7 +694,12 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD,

if (const SYCLIntelNoGlobalWorkOffsetAttr *A =
FD->getAttr<SYCLIntelNoGlobalWorkOffsetAttr>()) {
if (A->getEnabled())
const Expr *Arg = A->getValue();
assert(Arg && "Got an unexpected null argument");
Optional<llvm::APSInt> ArgVal =
Arg->getIntegerConstantExpr(FD->getASTContext());
assert(ArgVal.hasValue() && "Not an integer constant expression");
if (ArgVal->getBoolValue())
Fn->setMetadata("no_global_work_offset", llvm::MDNode::get(Context, {}));
}

Expand Down
20 changes: 7 additions & 13 deletions clang/lib/Sema/SemaDeclAttr.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -5294,24 +5294,18 @@ static void handleNoGlobalWorkOffsetAttr(Sema &S, Decl *D,

checkForDuplicateAttribute<SYCLIntelNoGlobalWorkOffsetAttr>(S, D, Attr);

uint32_t Enabled = 1;
if (Attr.getNumArgs()) {
const Expr *E = Attr.getArgAsExpr(0);
if (!checkUInt32Argument(S, Attr, E, Enabled, 0,
/*StrictlyUnsigned=*/true))
return;
}
if (Enabled > 1)
S.Diag(Attr.getLoc(), diag::warn_boolean_attribute_argument_is_not_valid)
<< Attr;

if (Attr.getKind() == ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset &&
checkDeprecatedSYCLAttributeSpelling(S, Attr))
S.Diag(Attr.getLoc(), diag::note_spelling_suggestion)
<< "'intel::no_global_work_offset'";

D->addAttr(::new (S.Context)
SYCLIntelNoGlobalWorkOffsetAttr(S.Context, Attr, Enabled));
// If no attribute argument is specified, set to default value '1'.
Expr *E = Attr.isArgExpr(0)
? Attr.getArgAsExpr(0)
: IntegerLiteral::Create(S.Context, llvm::APInt(32, 1),
S.Context.IntTy, Attr.getLoc());
S.addIntelSYCLSingleArgFunctionAttr<SYCLIntelNoGlobalWorkOffsetAttr>(D, Attr,
E);
}

/// Handle the [[intelfpga::doublepump]] and [[intelfpga::singlepump]] attributes.
Expand Down
6 changes: 6 additions & 0 deletions clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -775,6 +775,12 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs,
*this, TemplateArgs, SYCLIntelMaxGlobalWorkDim, New);
continue;
}
if (const auto *SYCLIntelNoGlobalWorkOffset =
dyn_cast<SYCLIntelNoGlobalWorkOffsetAttr>(TmplAttr)) {
instantiateIntelSYCLFunctionAttr<SYCLIntelNoGlobalWorkOffsetAttr>(
*this, TemplateArgs, SYCLIntelNoGlobalWorkOffset, New);
continue;
}
// Existing DLL attribute on the instantiation takes precedence.
if (TmplAttr->getKind() == attr::DLLExport ||
TmplAttr->getKind() == attr::DLLImport) {
Expand Down
51 changes: 36 additions & 15 deletions clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp
Original file line number Diff line number Diff line change
@@ -1,28 +1,49 @@
// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s
// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s

#include "sycl.hpp"

using namespace cl::sycl;
queue q;

class Foo {
public:
[[intel::no_global_work_offset(1)]] void operator()() const {}
};

template <typename name, typename Func>
__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
kernelFunc();
}
template <int SIZE>
class Functor {
public:
[[intel::no_global_work_offset(SIZE)]] void operator()() const {}
};

template <int N>
[[intel::no_global_work_offset(N)]] void func() {}

int main() {
q.submit([&](handler &h) {
Foo boo;
h.single_task<class kernel_name1>(boo);

h.single_task<class kernel_name2>(
[]() [[intel::no_global_work_offset]]{});

void bar() {
Foo boo;
kernel<class kernel_name1>(boo);
h.single_task<class kernel_name3>(
[]() [[intel::no_global_work_offset(0)]]{});

kernel<class kernel_name2>(
[]() [[intel::no_global_work_offset]]{});
Functor<1> f;
h.single_task<class kernel_name4>(f);

kernel<class kernel_name3>(
[]() [[intel::no_global_work_offset(0)]]{});
h.single_task<class kernel_name5>([]() {
func<1>();
});
});
return 0;
}

// CHECK: define spir_kernel void @{{.*}}kernel_name1() {{.*}} !no_global_work_offset ![[NUM5:[0-9]+]]
// CHECK: define spir_kernel void @{{.*}}kernel_name2() {{.*}} !no_global_work_offset ![[NUM5]]
// CHECK: define spir_kernel void @{{.*}}kernel_name3() {{.*}} ![[NUM4:[0-9]+]]
// CHECK: define spir_kernel void @{{.*}}kernel_name1"() #0 {{.*}} !no_global_work_offset ![[NUM5:[0-9]+]]
// CHECK: define spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !no_global_work_offset ![[NUM5]]
// CHECK: define spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} ![[NUM4:[0-9]+]]
// CHECK: define spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} !no_global_work_offset ![[NUM5]]
// CHECK: define spir_kernel void @{{.*}}kernel_name5"() #0 {{.*}} !no_global_work_offset ![[NUM5]]
// CHECK-NOT: ![[NUM4]] = !{i32 0}
// CHECK: ![[NUM5]] = !{}
Original file line number Diff line number Diff line change
Expand Up @@ -46,7 +46,7 @@ void invoke_foo2() {
// CHECK-LABEL: FunctionDecl {{.*}} invoke_foo2 'void ()'
// CHECK: `-FunctionDecl {{.*}}KernelName 'void ()'
// CHECK: -IntelReqdSubGroupSizeAttr {{.*}}
// CHECK: `-SYCLIntelNoGlobalWorkOffsetAttr {{.*}} Enabled
// CHECK: `-SYCLIntelNoGlobalWorkOffsetAttr {{.*}}
parallel_for<class KernelName>([]() {});
#else
parallel_for<class KernelName>([]() {}); // expected-error 2 {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}}
Expand Down
81 changes: 40 additions & 41 deletions clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp
Original file line number Diff line number Diff line change
@@ -1,51 +1,50 @@
// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-return-type -fcxx-exceptions -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s
// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -Wno-return-type -Wno-sycl-2017-compat -fcxx-exceptions -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s

#include "sycl.hpp"

using namespace cl::sycl;
queue q;

struct FuncObj {
//expected-warning@+2 {{attribute 'intelfpga::no_global_work_offset' is deprecated}}
//expected-note@+1 {{did you mean to use 'intel::no_global_work_offset' instead?}}
[[intelfpga::no_global_work_offset]] void operator()() {}
[[intelfpga::no_global_work_offset]] void operator()() const {}
};

template <typename name, typename Func>
void kernel(Func kernelFunc) {
kernelFunc();
}

int main() {
// CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}}Enabled
kernel<class test_kernel1>([]() {
FuncObj();
});

// CHECK: SYCLIntelNoGlobalWorkOffsetAttr
// CHECK-NOT: Enabled
kernel<class test_kernel2>(
[]() [[intel::no_global_work_offset(0)]]{});

// CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}}Enabled
// expected-warning@+2{{'no_global_work_offset' attribute should be 0 or 1. Adjusted to 1}}
kernel<class test_kernel3>(
[]() [[intel::no_global_work_offset(42)]]{});

// expected-error@+2{{'no_global_work_offset' attribute requires a non-negative integral compile time constant expression}}
kernel<class test_kernel4>(
[]() [[intel::no_global_work_offset(-1)]]{});

// expected-error@+2{{'no_global_work_offset' attribute requires parameter 0 to be an integer constant}}
kernel<class test_kernel5>(
[]() [[intel::no_global_work_offset("foo")]]{});

kernel<class test_kernel6>([]() {
// expected-error@+1{{'no_global_work_offset' attribute only applies to functions}}
[[intel::no_global_work_offset(1)]] int a;
q.submit([&](handler &h) {
// CHECK: SYCLIntelNoGlobalWorkOffsetAttr {{.*}}
// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}}
h.single_task<class test_kernel1>(FuncObj());

// CHECK: SYCLIntelNoGlobalWorkOffsetAttr {{.*}}
// CHECK-NEXT: IntegerLiteral{{.*}}0{{$}}
h.single_task<class test_kernel2>(
[]() [[intel::no_global_work_offset(0)]]{});

// CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}}
// CHECK-NEXT: IntegerLiteral{{.*}}42{{$}}
h.single_task<class test_kernel3>(
[]() [[intel::no_global_work_offset(42)]]{});

// CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}}
// CHECK-NEXT: UnaryOperator{{.*}} 'int' prefix '-'
// CHECK-NEXT-NEXT: IntegerLiteral{{.*}}1{{$}}
h.single_task<class test_kernel4>(
[]() [[intel::no_global_work_offset(-1)]]{});

// expected-error@+2{{'no_global_work_offset' attribute requires an integer constant}}
h.single_task<class test_kernel5>(
[]() [[intel::no_global_work_offset("foo")]]{});

h.single_task<class test_kernel6>([]() {
// expected-error@+1{{'no_global_work_offset' attribute only applies to functions}}
[[intel::no_global_work_offset(1)]] int a;
});

// expected-warning@+2{{attribute 'no_global_work_offset' is already applied}}
h.single_task<class test_kernel7>(
[]() [[intel::no_global_work_offset(0), intel::no_global_work_offset(1)]]{});
});

// CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}}
// CHECK-NOT: Enabled
// CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}}Enabled
// expected-warning@+2{{attribute 'no_global_work_offset' is already applied}}
kernel<class test_kernel7>(
[]() [[intel::no_global_work_offset(0), intel::no_global_work_offset(1)]]{});

return 0;
}
40 changes: 23 additions & 17 deletions clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp
Original file line number Diff line number Diff line change
@@ -1,8 +1,11 @@
// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -triple spir64 -Wno-sycl-2017-compat -verify
// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -triple spir64 -DTRIGGER_ERROR -Wno-sycl-2017-compat -verify
// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -triple spir64 -Wno-sycl-2017-compat | FileCheck %s
// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -Wno-sycl-2017-compat -verify
// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -DTRIGGER_ERROR -Wno-sycl-2017-compat -verify
// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -Wno-sycl-2017-compat | FileCheck %s

#include "Inputs/sycl.hpp"
#include "sycl.hpp"

using namespace cl::sycl;
queue q;

#ifndef TRIGGER_ERROR
//first case - good case
Expand Down Expand Up @@ -46,23 +49,26 @@ func4() {} // expected-error {{'max_work_group_size' attribute conflicts with ''
#endif

int main() {
q.submit([&](handler &h) {
#ifndef TRIGGER_ERROR
// CHECK-LABEL: FunctionDecl {{.*}} main 'int ()'
// CHECK: `-FunctionDecl {{.*}}test_kernel1 'void ()'
// CHECK: -SYCLIntelMaxWorkGroupSizeAttr {{.*}} Inherited 4 4 4
// CHECK: -SYCLIntelNoGlobalWorkOffsetAttr {{.*}} Inherited Enabled
// CHECK: `-ReqdWorkGroupSizeAttr {{.*}} 2 2 2
cl::sycl::kernel_single_task<class test_kernel1>(
[]() { func1(); });
// CHECK-LABEL: FunctionDecl {{.*}} main 'int ()'
// CHECK: `-FunctionDecl {{.*}}test_kernel1 'void ()'
// CHECK: -SYCLIntelMaxWorkGroupSizeAttr {{.*}} Inherited 4 4 4
// CHECK: -SYCLIntelNoGlobalWorkOffsetAttr {{.*}}
// CHECK: `-ReqdWorkGroupSizeAttr {{.*}} 2 2 2
h.single_task<class test_kernel1>(
[]() { func1(); });

#else
cl::sycl::kernel_single_task<class test_kernel2>(
[]() { func2(); }); // expected-error {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}}
h.single_task<class test_kernel2>(
[]() { func2(); }); // expected-error {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}}

cl::sycl::kernel_single_task<class test_kernel3>(
[]() { func3(); });
h.single_task<class test_kernel3>(
[]() { func3(); });

cl::sycl::kernel_single_task<class test_kernel4>(
[]() { func4(); });
h.single_task<class test_kernel4>(
[]() { func4(); });
#endif
});
return 0;
}
Original file line number Diff line number Diff line change
@@ -0,0 +1,60 @@
// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s

// Test that checks template parameter support for 'no_global_work_offset' attribute on sycl device.

// Test that checks wrong function template instantiation and ensures that the type
// is checked properly when instantiating from the template definition.
template <typename Ty>
// expected-error@+1{{'no_global_work_offset' attribute requires an integer constant}}
[[intel::no_global_work_offset(Ty{})]] void func() {}

struct S {};
void var() {
//expected-note@+1{{in instantiation of function template specialization 'func<S>' requested here}}
func<S>();
}

// Test that checks expression is not a constant expression.
int foo();
// expected-error@+1{{'no_global_work_offset' attribute requires an integer constant}}
[[intel::no_global_work_offset(foo() + 12)]] void func1();

// Test that checks expression is a constant expression.
constexpr int bar() { return 0; }
[[intel::no_global_work_offset(bar() + 12)]] void func2(); // OK

// Test that checks template parameter suppport on member function of class template.
template <int SIZE>
class KernelFunctor {
public:
[[intel::no_global_work_offset(SIZE)]] void operator()() {}
};

int main() {
KernelFunctor<1>();
}

// CHECK: ClassTemplateDecl {{.*}} {{.*}} KernelFunctor
// CHECK: ClassTemplateSpecializationDecl {{.*}} {{.*}} class KernelFunctor definition
// CHECK: CXXRecordDecl {{.*}} {{.*}} implicit class KernelFunctor
// CHECK: SYCLIntelNoGlobalWorkOffsetAttr {{.*}}
// CHECK: SubstNonTypeTemplateParmExpr {{.*}}
// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}}
// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}}

// Test that checks template parameter suppport on function.
template <int N>
[[intel::no_global_work_offset(N)]] void func3() {}

int check() {
func3<1>();
return 0;
}

// CHECK: FunctionTemplateDecl {{.*}} {{.*}} func3
// CHECK: NonTypeTemplateParmDecl {{.*}} {{.*}} referenced 'int' depth 0 index 0 N
// CHECK: FunctionDecl {{.*}} {{.*}} func3 'void ()'
// CHECK: SYCLIntelNoGlobalWorkOffsetAttr {{.*}}
// CHECK: SubstNonTypeTemplateParmExpr {{.*}}
// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}}
// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}}