Skip to content

Commit a5fde5a

Browse files
smanna12smanna12
authored
[SYCL] Add template parameter support for no_global_work_offset attribute (#2839)
This patch adds support for template parameter on [[intel:: no_global_work_offset())]] attribute where valid values are 0 and 1 and attribute parameter is optional, so [[intelfpga::no_global_work_offset]] means the same as [[intelfpga::no_global_work_offset(1)]]. updates sema/codegen tests with mock headers on device. uses existing function "sema::addIntelSYCLSingleArgFunctionAttr" from other single argument function attributes such as num_simd_work_items, max_global_work_dim, and intel_reqd_sub_group_size to avoid source codes duplication and reuse for the template parameter support. Signed-off-by: Soumi Manna <[email protected]>
1 parent a5065ab commit a5fde5a

11 files changed

+181
-95
lines changed

‎clang/include/clang/Basic/Attr.td

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1306,7 +1306,7 @@ def SYCLIntelMaxGlobalWorkDim : InheritableAttr {
13061306
def SYCLIntelNoGlobalWorkOffset : InheritableAttr {
13071307
let Spellings = [CXX11<"intelfpga","no_global_work_offset">,
13081308
CXX11<"intel","no_global_work_offset">];
1309-
let Args = [BoolArgument<"Enabled", 1>];
1309+
let Args = [ExprArgument<"Value", /*optional*/1>];
13101310
let LangOpts = [SYCLIsDevice, SYCLIsHost];
13111311
let Subjects = SubjectList<[Function], ErrorDiag>;
13121312
let Documentation = [SYCLIntelNoGlobalWorkOffsetAttrDocs];

‎clang/include/clang/Basic/DiagnosticGroups.td

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -667,10 +667,8 @@ def NSReturnsMismatch : DiagGroup<"nsreturns-mismatch">;
667667
def IndependentClassAttribute : DiagGroup<"IndependentClass-attribute">;
668668
def UnknownAttributes : DiagGroup<"unknown-attributes">;
669669
def IgnoredAttributes : DiagGroup<"ignored-attributes">;
670-
def AdjustedAttributes : DiagGroup<"adjusted-attributes">;
671670
def Attributes : DiagGroup<"attributes", [UnknownAttributes,
672-
IgnoredAttributes,
673-
AdjustedAttributes]>;
671+
IgnoredAttributes]>;
674672
def UnknownSanitizers : DiagGroup<"unknown-sanitizers">;
675673
def UnnamedTypeTemplateArgs : DiagGroup<"unnamed-type-template-args",
676674
[CXX98CompatUnnamedTypeTemplateArgs]>;

‎clang/include/clang/Basic/DiagnosticSemaKinds.td

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -11143,9 +11143,6 @@ def err_sycl_function_attribute_mismatch : Error<
1114311143
"SYCL kernel without %0 attribute can't call a function with this attribute">;
1114411144
def err_sycl_x_y_z_arguments_must_be_one : Error<
1114511145
"%0 X-, Y- and Z- sizes must be 1 when %1 attribute is used with value 0">;
11146-
def warn_boolean_attribute_argument_is_not_valid: Warning<
11147-
"The value of %0 attribute should be 0 or 1. Adjusted to 1">,
11148-
InGroup<AdjustedAttributes>;
1114911146
def err_sycl_attibute_cannot_be_applied_here
1115011147
: Error<"%0 attribute cannot be applied to a "
1115111148
"static function or function in an anonymous namespace">;

‎clang/lib/CodeGen/CodeGenFunction.cpp

Lines changed: 6 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -694,7 +694,12 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD,
694694

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

‎clang/lib/Sema/SemaDeclAttr.cpp

Lines changed: 7 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -5276,24 +5276,18 @@ static void handleNoGlobalWorkOffsetAttr(Sema &S, Decl *D,
52765276

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

5279-
uint32_t Enabled = 1;
5280-
if (Attr.getNumArgs()) {
5281-
const Expr *E = Attr.getArgAsExpr(0);
5282-
if (!checkUInt32Argument(S, Attr, E, Enabled, 0,
5283-
/*StrictlyUnsigned=*/true))
5284-
return;
5285-
}
5286-
if (Enabled > 1)
5287-
S.Diag(Attr.getLoc(), diag::warn_boolean_attribute_argument_is_not_valid)
5288-
<< Attr;
5289-
52905279
if (Attr.getKind() == ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset &&
52915280
checkDeprecatedSYCLAttributeSpelling(S, Attr))
52925281
S.Diag(Attr.getLoc(), diag::note_spelling_suggestion)
52935282
<< "'intel::no_global_work_offset'";
52945283

5295-
D->addAttr(::new (S.Context)
5296-
SYCLIntelNoGlobalWorkOffsetAttr(S.Context, Attr, Enabled));
5284+
// If no attribute argument is specified, set to default value '1'.
5285+
Expr *E = Attr.isArgExpr(0)
5286+
? Attr.getArgAsExpr(0)
5287+
: IntegerLiteral::Create(S.Context, llvm::APInt(32, 1),
5288+
S.Context.IntTy, Attr.getLoc());
5289+
S.addIntelSYCLSingleArgFunctionAttr<SYCLIntelNoGlobalWorkOffsetAttr>(D, Attr,
5290+
E);
52975291
}
52985292

52995293
/// Handle the [[intelfpga::doublepump]] and [[intelfpga::singlepump]] attributes.

‎clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -775,6 +775,12 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs,
775775
*this, TemplateArgs, SYCLIntelMaxGlobalWorkDim, New);
776776
continue;
777777
}
778+
if (const auto *SYCLIntelNoGlobalWorkOffset =
779+
dyn_cast<SYCLIntelNoGlobalWorkOffsetAttr>(TmplAttr)) {
780+
instantiateIntelSYCLFunctionAttr<SYCLIntelNoGlobalWorkOffsetAttr>(
781+
*this, TemplateArgs, SYCLIntelNoGlobalWorkOffset, New);
782+
continue;
783+
}
778784
// Existing DLL attribute on the instantiation takes precedence.
779785
if (TmplAttr->getKind() == attr::DLLExport ||
780786
TmplAttr->getKind() == attr::DLLImport) {

‎clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp

Lines changed: 36 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -1,28 +1,49 @@
1-
// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s
1+
// 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
2+
3+
#include "sycl.hpp"
4+
5+
using namespace cl::sycl;
6+
queue q;
27

38
class Foo {
49
public:
510
[[intel::no_global_work_offset(1)]] void operator()() const {}
611
};
712

8-
template <typename name, typename Func>
9-
__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
10-
kernelFunc();
11-
}
13+
template <int SIZE>
14+
class Functor {
15+
public:
16+
[[intel::no_global_work_offset(SIZE)]] void operator()() const {}
17+
};
18+
19+
template <int N>
20+
[[intel::no_global_work_offset(N)]] void func() {}
21+
22+
int main() {
23+
q.submit([&](handler &h) {
24+
Foo boo;
25+
h.single_task<class kernel_name1>(boo);
26+
27+
h.single_task<class kernel_name2>(
28+
[]() [[intel::no_global_work_offset]]{});
1229

13-
void bar() {
14-
Foo boo;
15-
kernel<class kernel_name1>(boo);
30+
h.single_task<class kernel_name3>(
31+
[]() [[intel::no_global_work_offset(0)]]{});
1632

17-
kernel<class kernel_name2>(
18-
[]() [[intel::no_global_work_offset]]{});
33+
Functor<1> f;
34+
h.single_task<class kernel_name4>(f);
1935

20-
kernel<class kernel_name3>(
21-
[]() [[intel::no_global_work_offset(0)]]{});
36+
h.single_task<class kernel_name5>([]() {
37+
func<1>();
38+
});
39+
});
40+
return 0;
2241
}
2342

24-
// CHECK: define spir_kernel void @{{.*}}kernel_name1() {{.*}} !no_global_work_offset ![[NUM5:[0-9]+]]
25-
// CHECK: define spir_kernel void @{{.*}}kernel_name2() {{.*}} !no_global_work_offset ![[NUM5]]
26-
// CHECK: define spir_kernel void @{{.*}}kernel_name3() {{.*}} ![[NUM4:[0-9]+]]
43+
// CHECK: define spir_kernel void @{{.*}}kernel_name1"() #0 {{.*}} !no_global_work_offset ![[NUM5:[0-9]+]]
44+
// CHECK: define spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !no_global_work_offset ![[NUM5]]
45+
// CHECK: define spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} ![[NUM4:[0-9]+]]
46+
// CHECK: define spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} !no_global_work_offset ![[NUM5]]
47+
// CHECK: define spir_kernel void @{{.*}}kernel_name5"() #0 {{.*}} !no_global_work_offset ![[NUM5]]
2748
// CHECK-NOT: ![[NUM4]] = !{i32 0}
2849
// CHECK: ![[NUM5]] = !{}

‎clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -46,7 +46,7 @@ void invoke_foo2() {
4646
// CHECK-LABEL: FunctionDecl {{.*}} invoke_foo2 'void ()'
4747
// CHECK: `-FunctionDecl {{.*}}KernelName 'void ()'
4848
// CHECK: -IntelReqdSubGroupSizeAttr {{.*}}
49-
// CHECK: `-SYCLIntelNoGlobalWorkOffsetAttr {{.*}} Enabled
49+
// CHECK: `-SYCLIntelNoGlobalWorkOffsetAttr {{.*}}
5050
parallel_for<class KernelName>([]() {});
5151
#else
5252
parallel_for<class KernelName>([]() {}); // expected-error 2 {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}}

‎clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp

Lines changed: 40 additions & 41 deletions
Original file line numberDiff line numberDiff line change
@@ -1,51 +1,50 @@
1-
// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-return-type -fcxx-exceptions -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s
1+
// 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
2+
3+
#include "sycl.hpp"
4+
5+
using namespace cl::sycl;
6+
queue q;
27

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

9-
template <typename name, typename Func>
10-
void kernel(Func kernelFunc) {
11-
kernelFunc();
12-
}
13-
1414
int main() {
15-
// CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}}Enabled
16-
kernel<class test_kernel1>([]() {
17-
FuncObj();
18-
});
19-
20-
// CHECK: SYCLIntelNoGlobalWorkOffsetAttr
21-
// CHECK-NOT: Enabled
22-
kernel<class test_kernel2>(
23-
[]() [[intel::no_global_work_offset(0)]]{});
24-
25-
// CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}}Enabled
26-
// expected-warning@+2{{'no_global_work_offset' attribute should be 0 or 1. Adjusted to 1}}
27-
kernel<class test_kernel3>(
28-
[]() [[intel::no_global_work_offset(42)]]{});
29-
30-
// expected-error@+2{{'no_global_work_offset' attribute requires a non-negative integral compile time constant expression}}
31-
kernel<class test_kernel4>(
32-
[]() [[intel::no_global_work_offset(-1)]]{});
33-
34-
// expected-error@+2{{'no_global_work_offset' attribute requires parameter 0 to be an integer constant}}
35-
kernel<class test_kernel5>(
36-
[]() [[intel::no_global_work_offset("foo")]]{});
37-
38-
kernel<class test_kernel6>([]() {
39-
// expected-error@+1{{'no_global_work_offset' attribute only applies to functions}}
40-
[[intel::no_global_work_offset(1)]] int a;
15+
q.submit([&](handler &h) {
16+
// CHECK: SYCLIntelNoGlobalWorkOffsetAttr {{.*}}
17+
// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}}
18+
h.single_task<class test_kernel1>(FuncObj());
19+
20+
// CHECK: SYCLIntelNoGlobalWorkOffsetAttr {{.*}}
21+
// CHECK-NEXT: IntegerLiteral{{.*}}0{{$}}
22+
h.single_task<class test_kernel2>(
23+
[]() [[intel::no_global_work_offset(0)]]{});
24+
25+
// CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}}
26+
// CHECK-NEXT: IntegerLiteral{{.*}}42{{$}}
27+
h.single_task<class test_kernel3>(
28+
[]() [[intel::no_global_work_offset(42)]]{});
29+
30+
// CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}}
31+
// CHECK-NEXT: UnaryOperator{{.*}} 'int' prefix '-'
32+
// CHECK-NEXT-NEXT: IntegerLiteral{{.*}}1{{$}}
33+
h.single_task<class test_kernel4>(
34+
[]() [[intel::no_global_work_offset(-1)]]{});
35+
36+
// expected-error@+2{{'no_global_work_offset' attribute requires an integer constant}}
37+
h.single_task<class test_kernel5>(
38+
[]() [[intel::no_global_work_offset("foo")]]{});
39+
40+
h.single_task<class test_kernel6>([]() {
41+
// expected-error@+1{{'no_global_work_offset' attribute only applies to functions}}
42+
[[intel::no_global_work_offset(1)]] int a;
43+
});
44+
45+
// expected-warning@+2{{attribute 'no_global_work_offset' is already applied}}
46+
h.single_task<class test_kernel7>(
47+
[]() [[intel::no_global_work_offset(0), intel::no_global_work_offset(1)]]{});
4148
});
42-
43-
// CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}}
44-
// CHECK-NOT: Enabled
45-
// CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}}Enabled
46-
// expected-warning@+2{{attribute 'no_global_work_offset' is already applied}}
47-
kernel<class test_kernel7>(
48-
[]() [[intel::no_global_work_offset(0), intel::no_global_work_offset(1)]]{});
49-
5049
return 0;
5150
}

‎clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp

Lines changed: 23 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -1,8 +1,11 @@
1-
// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -triple spir64 -Wno-sycl-2017-compat -verify
2-
// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -triple spir64 -DTRIGGER_ERROR -Wno-sycl-2017-compat -verify
3-
// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -triple spir64 -Wno-sycl-2017-compat | FileCheck %s
1+
// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -Wno-sycl-2017-compat -verify
2+
// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -DTRIGGER_ERROR -Wno-sycl-2017-compat -verify
3+
// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -Wno-sycl-2017-compat | FileCheck %s
44

5-
#include "Inputs/sycl.hpp"
5+
#include "sycl.hpp"
6+
7+
using namespace cl::sycl;
8+
queue q;
69

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

4851
int main() {
52+
q.submit([&](handler &h) {
4953
#ifndef TRIGGER_ERROR
50-
// CHECK-LABEL: FunctionDecl {{.*}} main 'int ()'
51-
// CHECK: `-FunctionDecl {{.*}}test_kernel1 'void ()'
52-
// CHECK: -SYCLIntelMaxWorkGroupSizeAttr {{.*}} Inherited 4 4 4
53-
// CHECK: -SYCLIntelNoGlobalWorkOffsetAttr {{.*}} Inherited Enabled
54-
// CHECK: `-ReqdWorkGroupSizeAttr {{.*}} 2 2 2
55-
cl::sycl::kernel_single_task<class test_kernel1>(
56-
[]() { func1(); });
54+
// CHECK-LABEL: FunctionDecl {{.*}} main 'int ()'
55+
// CHECK: `-FunctionDecl {{.*}}test_kernel1 'void ()'
56+
// CHECK: -SYCLIntelMaxWorkGroupSizeAttr {{.*}} Inherited 4 4 4
57+
// CHECK: -SYCLIntelNoGlobalWorkOffsetAttr {{.*}}
58+
// CHECK: `-ReqdWorkGroupSizeAttr {{.*}} 2 2 2
59+
h.single_task<class test_kernel1>(
60+
[]() { func1(); });
5761

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

62-
cl::sycl::kernel_single_task<class test_kernel3>(
63-
[]() { func3(); });
66+
h.single_task<class test_kernel3>(
67+
[]() { func3(); });
6468

65-
cl::sycl::kernel_single_task<class test_kernel4>(
66-
[]() { func4(); });
69+
h.single_task<class test_kernel4>(
70+
[]() { func4(); });
6771
#endif
72+
});
73+
return 0;
6874
}

‎clang/test/SemaSYCL/sycl-device-intel-fpga-no-global-work-offset.cpp

Lines changed: 60 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,60 @@
1+
// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s
2+
3+
// Test that checks template parameter support for 'no_global_work_offset' attribute on sycl device.
4+
5+
// Test that checks wrong function template instantiation and ensures that the type
6+
// is checked properly when instantiating from the template definition.
7+
template <typename Ty>
8+
// expected-error@+1{{'no_global_work_offset' attribute requires an integer constant}}
9+
[[intel::no_global_work_offset(Ty{})]] void func() {}
10+
11+
struct S {};
12+
void var() {
13+
//expected-note@+1{{in instantiation of function template specialization 'func<S>' requested here}}
14+
func<S>();
15+
}
16+
17+
// Test that checks expression is not a constant expression.
18+
int foo();
19+
// expected-error@+1{{'no_global_work_offset' attribute requires an integer constant}}
20+
[[intel::no_global_work_offset(foo() + 12)]] void func1();
21+
22+
// Test that checks expression is a constant expression.
23+
constexpr int bar() { return 0; }
24+
[[intel::no_global_work_offset(bar() + 12)]] void func2(); // OK
25+
26+
// Test that checks template parameter suppport on member function of class template.
27+
template <int SIZE>
28+
class KernelFunctor {
29+
public:
30+
[[intel::no_global_work_offset(SIZE)]] void operator()() {}
31+
};
32+
33+
int main() {
34+
KernelFunctor<1>();
35+
}
36+
37+
// CHECK: ClassTemplateDecl {{.*}} {{.*}} KernelFunctor
38+
// CHECK: ClassTemplateSpecializationDecl {{.*}} {{.*}} class KernelFunctor definition
39+
// CHECK: CXXRecordDecl {{.*}} {{.*}} implicit class KernelFunctor
40+
// CHECK: SYCLIntelNoGlobalWorkOffsetAttr {{.*}}
41+
// CHECK: SubstNonTypeTemplateParmExpr {{.*}}
42+
// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}}
43+
// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}}
44+
45+
// Test that checks template parameter suppport on function.
46+
template <int N>
47+
[[intel::no_global_work_offset(N)]] void func3() {}
48+
49+
int check() {
50+
func3<1>();
51+
return 0;
52+
}
53+
54+
// CHECK: FunctionTemplateDecl {{.*}} {{.*}} func3
55+
// CHECK: NonTypeTemplateParmDecl {{.*}} {{.*}} referenced 'int' depth 0 index 0 N
56+
// CHECK: FunctionDecl {{.*}} {{.*}} func3 'void ()'
57+
// CHECK: SYCLIntelNoGlobalWorkOffsetAttr {{.*}}
58+
// CHECK: SubstNonTypeTemplateParmExpr {{.*}}
59+
// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}}
60+
// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}}

0 commit comments

Comments
 (0)