Skip to content
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{{$}}