Skip to content

Commit 016cb2a

Browse files
vmaksimovmaksimo
authored
Merge branch 'sycl' into pulldown_ww07-08_oldPM
2 parents 4addbb8 + 43ca44a commit 016cb2a

18 files changed

+443
-36
lines changed

‎clang/include/clang/Sema/Sema.h

Lines changed: 11 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -10222,6 +10222,16 @@ class Sema final {
1022210222
SYCLIntelNumSimdWorkItemsAttr *
1022310223
MergeSYCLIntelNumSimdWorkItemsAttr(Decl *D,
1022410224
const SYCLIntelNumSimdWorkItemsAttr &A);
10225+
void AddSYCLIntelSchedulerTargetFmaxMhzAttr(Decl *D,
10226+
const AttributeCommonInfo &CI,
10227+
Expr *E);
10228+
SYCLIntelSchedulerTargetFmaxMhzAttr *MergeSYCLIntelSchedulerTargetFmaxMhzAttr(
10229+
Decl *D, const SYCLIntelSchedulerTargetFmaxMhzAttr &A);
10230+
void AddSYCLIntelNoGlobalWorkOffsetAttr(Decl *D,
10231+
const AttributeCommonInfo &CI,
10232+
Expr *E);
10233+
SYCLIntelNoGlobalWorkOffsetAttr *MergeSYCLIntelNoGlobalWorkOffsetAttr(
10234+
Decl *D, const SYCLIntelNoGlobalWorkOffsetAttr &A);
1022510235

1022610236
/// AddAlignedAttr - Adds an aligned attribute to a particular declaration.
1022710237
void AddAlignedAttr(Decl *D, const AttributeCommonInfo &CI, Expr *E,
@@ -13104,8 +13114,7 @@ void Sema::addIntelSingleArgAttr(Decl *D, const AttributeCommonInfo &CI,
1310413114
return;
1310513115
}
1310613116
}
13107-
if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelSchedulerTargetFmaxMhz ||
13108-
CI.getParsedKind() == ParsedAttr::AT_IntelFPGAPrivateCopies) {
13117+
if (CI.getParsedKind() == ParsedAttr::AT_IntelFPGAPrivateCopies) {
1310913118
if (ArgInt < 0) {
1311013119
Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer)
1311113120
<< CI << /*non-negative*/ 1;

‎clang/lib/Driver/Driver.cpp

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3791,6 +3791,20 @@ class OffloadingActionBuilder final {
37913791
getDeviceDependences(OffloadAction::DeviceDependences &DA,
37923792
phases::ID CurPhase, phases::ID FinalPhase,
37933793
PhasesTy &Phases) override {
3794+
if (CurPhase == phases::Preprocess) {
3795+
// Do not perform the host compilation when doing preprocessing only
3796+
// with -fsycl-device-only.
3797+
bool IsPreprocessOnly =
3798+
Args.getLastArg(options::OPT_E) ||
3799+
Args.getLastArg(options::OPT__SLASH_EP, options::OPT__SLASH_P) ||
3800+
Args.getLastArg(options::OPT_M, options::OPT_MM);
3801+
if (Args.hasArg(options::OPT_fsycl_device_only) && IsPreprocessOnly) {
3802+
for (Action *&A : SYCLDeviceActions)
3803+
A = C.getDriver().ConstructPhaseAction(C, Args, CurPhase, A,
3804+
AssociatedOffloadKind);
3805+
return ABRT_Ignore_Host;
3806+
}
3807+
}
37943808

37953809
// FIXME: This adds the integration header generation pass before the
37963810
// Host compilation pass so the Host can use the header generated. This

‎clang/lib/Sema/SemaDecl.cpp

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2622,6 +2622,10 @@ static bool mergeDeclAttribute(Sema &S, NamedDecl *D,
26222622
NewAttr = S.MergeIntelReqdSubGroupSizeAttr(D, *A);
26232623
else if (const auto *A = dyn_cast<SYCLIntelNumSimdWorkItemsAttr>(Attr))
26242624
NewAttr = S.MergeSYCLIntelNumSimdWorkItemsAttr(D, *A);
2625+
else if (const auto *A = dyn_cast<SYCLIntelSchedulerTargetFmaxMhzAttr>(Attr))
2626+
NewAttr = S.MergeSYCLIntelSchedulerTargetFmaxMhzAttr(D, *A);
2627+
else if (const auto *A = dyn_cast<SYCLIntelNoGlobalWorkOffsetAttr>(Attr))
2628+
NewAttr = S.MergeSYCLIntelNoGlobalWorkOffsetAttr(D, *A);
26252629
else if (Attr->shouldInheritEvenIfAlreadyPresent() || !DeclHasAttr(D, Attr))
26262630
NewAttr = cast<InheritableAttr>(Attr->clone(S.Context));
26272631

‎clang/lib/Sema/SemaDeclAttr.cpp

Lines changed: 115 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -3366,19 +3366,71 @@ static void handleUseStallEnableClustersAttr(Sema &S, Decl *D,
33663366
}
33673367

33683368
// Handle scheduler_target_fmax_mhz
3369-
static void handleSchedulerTargetFmaxMhzAttr(Sema &S, Decl *D,
3370-
const ParsedAttr &AL) {
3371-
if (D->isInvalidDecl())
3372-
return;
3369+
void Sema::AddSYCLIntelSchedulerTargetFmaxMhzAttr(Decl *D,
3370+
const AttributeCommonInfo &CI,
3371+
Expr *E) {
3372+
if (!E->isValueDependent()) {
3373+
// Validate that we have an integer constant expression and then store the
3374+
// converted constant expression into the semantic attribute so that we
3375+
// don't have to evaluate it again later.
3376+
llvm::APSInt ArgVal;
3377+
ExprResult Res = VerifyIntegerConstantExpression(E, &ArgVal);
3378+
if (Res.isInvalid())
3379+
return;
3380+
E = Res.get();
33733381

3374-
Expr *E = AL.getArgAsExpr(0);
3382+
// This attribute requires a non-negative value.
3383+
if (ArgVal < 0) {
3384+
Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer)
3385+
<< CI << /*non-negative*/ 1;
3386+
return;
3387+
}
3388+
// Check to see if there's a duplicate attribute with different values
3389+
// already applied to the declaration.
3390+
if (const auto *DeclAttr =
3391+
D->getAttr<SYCLIntelSchedulerTargetFmaxMhzAttr>()) {
3392+
// If the other attribute argument is instantiation dependent, we won't
3393+
// have converted it to a constant expression yet and thus we test
3394+
// whether this is a null pointer.
3395+
const auto *DeclExpr = dyn_cast<ConstantExpr>(DeclAttr->getValue());
3396+
if (DeclExpr && ArgVal != DeclExpr->getResultAsAPSInt()) {
3397+
Diag(CI.getLoc(), diag::warn_duplicate_attribute) << CI;
3398+
Diag(DeclAttr->getLoc(), diag::note_previous_attribute);
3399+
return;
3400+
}
3401+
}
3402+
}
33753403

3376-
if (D->getAttr<SYCLIntelSchedulerTargetFmaxMhzAttr>())
3377-
S.Diag(AL.getLoc(), diag::warn_duplicate_attribute) << AL;
3404+
D->addAttr(::new (Context)
3405+
SYCLIntelSchedulerTargetFmaxMhzAttr(Context, CI, E));
3406+
}
33783407

3408+
SYCLIntelSchedulerTargetFmaxMhzAttr *
3409+
Sema::MergeSYCLIntelSchedulerTargetFmaxMhzAttr(
3410+
Decl *D, const SYCLIntelSchedulerTargetFmaxMhzAttr &A) {
3411+
// Check to see if there's a duplicate attribute with different values
3412+
// already applied to the declaration.
3413+
if (const auto *DeclAttr =
3414+
D->getAttr<SYCLIntelSchedulerTargetFmaxMhzAttr>()) {
3415+
const auto *DeclExpr = dyn_cast<ConstantExpr>(DeclAttr->getValue());
3416+
const auto *MergeExpr = dyn_cast<ConstantExpr>(A.getValue());
3417+
if (DeclExpr && MergeExpr &&
3418+
DeclExpr->getResultAsAPSInt() != MergeExpr->getResultAsAPSInt()) {
3419+
Diag(DeclAttr->getLoc(), diag::warn_duplicate_attribute) << &A;
3420+
Diag(A.getLoc(), diag::note_previous_attribute);
3421+
return nullptr;
3422+
}
3423+
}
3424+
return ::new (Context)
3425+
SYCLIntelSchedulerTargetFmaxMhzAttr(Context, A, A.getValue());
3426+
}
3427+
3428+
static void handleSYCLIntelSchedulerTargetFmaxMhzAttr(Sema &S, Decl *D,
3429+
const ParsedAttr &AL) {
33793430
S.CheckDeprecatedSYCLAttributeSpelling(AL);
33803431

3381-
S.addIntelSingleArgAttr<SYCLIntelSchedulerTargetFmaxMhzAttr>(D, AL, E);
3432+
Expr *E = AL.getArgAsExpr(0);
3433+
S.AddSYCLIntelSchedulerTargetFmaxMhzAttr(D, AL, E);
33823434
}
33833435

33843436
// Handles max_global_work_dim.
@@ -5708,17 +5760,66 @@ static bool checkForDuplicateAttribute(Sema &S, Decl *D,
57085760
return false;
57095761
}
57105762

5711-
static void handleNoGlobalWorkOffsetAttr(Sema &S, Decl *D,
5712-
const ParsedAttr &A) {
5713-
checkForDuplicateAttribute<SYCLIntelNoGlobalWorkOffsetAttr>(S, D, A);
5763+
void Sema::AddSYCLIntelNoGlobalWorkOffsetAttr(Decl *D,
5764+
const AttributeCommonInfo &CI,
5765+
Expr *E) {
5766+
if (!E->isValueDependent()) {
5767+
// Validate that we have an integer constant expression and then store the
5768+
// converted constant expression into the semantic attribute so that we
5769+
// don't have to evaluate it again later.
5770+
llvm::APSInt ArgVal;
5771+
ExprResult Res = VerifyIntegerConstantExpression(E, &ArgVal);
5772+
if (Res.isInvalid())
5773+
return;
5774+
E = Res.get();
5775+
5776+
// Check to see if there's a duplicate attribute with different values
5777+
// already applied to the declaration.
5778+
if (const auto *DeclAttr = D->getAttr<SYCLIntelNoGlobalWorkOffsetAttr>()) {
5779+
// If the other attribute argument is instantiation dependent, we won't
5780+
// have converted it to a constant expression yet and thus we test
5781+
// whether this is a null pointer.
5782+
const auto *DeclExpr = dyn_cast<ConstantExpr>(DeclAttr->getValue());
5783+
if (DeclExpr && ArgVal != DeclExpr->getResultAsAPSInt()) {
5784+
Diag(CI.getLoc(), diag::warn_duplicate_attribute) << CI;
5785+
Diag(DeclAttr->getLoc(), diag::note_previous_attribute);
5786+
return;
5787+
}
5788+
}
5789+
}
5790+
5791+
D->addAttr(::new (Context) SYCLIntelNoGlobalWorkOffsetAttr(Context, CI, E));
5792+
}
5793+
5794+
SYCLIntelNoGlobalWorkOffsetAttr *Sema::MergeSYCLIntelNoGlobalWorkOffsetAttr(
5795+
Decl *D, const SYCLIntelNoGlobalWorkOffsetAttr &A) {
5796+
// Check to see if there's a duplicate attribute with different values
5797+
// already applied to the declaration.
5798+
if (const auto *DeclAttr = D->getAttr<SYCLIntelNoGlobalWorkOffsetAttr>()) {
5799+
const auto *DeclExpr = dyn_cast<ConstantExpr>(DeclAttr->getValue());
5800+
const auto *MergeExpr = dyn_cast<ConstantExpr>(A.getValue());
5801+
if (DeclExpr && MergeExpr &&
5802+
DeclExpr->getResultAsAPSInt() != MergeExpr->getResultAsAPSInt()) {
5803+
Diag(DeclAttr->getLoc(), diag::warn_duplicate_attribute) << &A;
5804+
Diag(A.getLoc(), diag::note_previous_attribute);
5805+
return nullptr;
5806+
}
5807+
}
5808+
return ::new (Context)
5809+
SYCLIntelNoGlobalWorkOffsetAttr(Context, A, A.getValue());
5810+
}
5811+
5812+
static void handleSYCLIntelNoGlobalWorkOffsetAttr(Sema &S, Decl *D,
5813+
const ParsedAttr &A) {
57145814
S.CheckDeprecatedSYCLAttributeSpelling(A);
57155815

57165816
// If no attribute argument is specified, set to default value '1'.
57175817
Expr *E = A.isArgExpr(0)
57185818
? A.getArgAsExpr(0)
57195819
: IntegerLiteral::Create(S.Context, llvm::APInt(32, 1),
57205820
S.Context.IntTy, A.getLoc());
5721-
S.addIntelSingleArgAttr<SYCLIntelNoGlobalWorkOffsetAttr>(D, A, E);
5821+
5822+
S.AddSYCLIntelNoGlobalWorkOffsetAttr(D, A, E);
57225823
}
57235824

57245825
/// Handle the [[intelfpga::doublepump]] and [[intelfpga::singlepump]] attributes.
@@ -9023,13 +9124,13 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
90239124
handleSYCLIntelNumSimdWorkItemsAttr(S, D, AL);
90249125
break;
90259126
case ParsedAttr::AT_SYCLIntelSchedulerTargetFmaxMhz:
9026-
handleSchedulerTargetFmaxMhzAttr(S, D, AL);
9127+
handleSYCLIntelSchedulerTargetFmaxMhzAttr(S, D, AL);
90279128
break;
90289129
case ParsedAttr::AT_SYCLIntelMaxGlobalWorkDim:
90299130
handleMaxGlobalWorkDimAttr(S, D, AL);
90309131
break;
90319132
case ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset:
9032-
handleNoGlobalWorkOffsetAttr(S, D, AL);
9133+
handleSYCLIntelNoGlobalWorkOffsetAttr(S, D, AL);
90339134
break;
90349135
case ParsedAttr::AT_SYCLIntelUseStallEnableClusters:
90359136
handleUseStallEnableClustersAttr(S, D, AL);

‎clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

Lines changed: 22 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -663,6 +663,26 @@ static void instantiateSYCLIntelNumSimdWorkItemsAttr(
663663
S.AddSYCLIntelNumSimdWorkItemsAttr(New, *A, Result.getAs<Expr>());
664664
}
665665

666+
static void instantiateSYCLIntelSchedulerTargetFmaxMhzAttr(
667+
Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,
668+
const SYCLIntelSchedulerTargetFmaxMhzAttr *A, Decl *New) {
669+
EnterExpressionEvaluationContext Unevaluated(
670+
S, Sema::ExpressionEvaluationContext::ConstantEvaluated);
671+
ExprResult Result = S.SubstExpr(A->getValue(), TemplateArgs);
672+
if (!Result.isInvalid())
673+
S.AddSYCLIntelSchedulerTargetFmaxMhzAttr(New, *A, Result.getAs<Expr>());
674+
}
675+
676+
static void instantiateSYCLIntelNoGlobalWorkOffsetAttr(
677+
Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,
678+
const SYCLIntelNoGlobalWorkOffsetAttr *A, Decl *New) {
679+
EnterExpressionEvaluationContext Unevaluated(
680+
S, Sema::ExpressionEvaluationContext::ConstantEvaluated);
681+
ExprResult Result = S.SubstExpr(A->getValue(), TemplateArgs);
682+
if (!Result.isInvalid())
683+
S.AddSYCLIntelNoGlobalWorkOffsetAttr(New, *A, Result.getAs<Expr>());
684+
}
685+
666686
template <typename AttrName>
667687
static void instantiateIntelSYCLFunctionAttr(
668688
Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,
@@ -866,7 +886,7 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs,
866886
}
867887
if (const auto *SYCLIntelSchedulerTargetFmaxMhz =
868888
dyn_cast<SYCLIntelSchedulerTargetFmaxMhzAttr>(TmplAttr)) {
869-
instantiateIntelSYCLFunctionAttr<SYCLIntelSchedulerTargetFmaxMhzAttr>(
889+
instantiateSYCLIntelSchedulerTargetFmaxMhzAttr(
870890
*this, TemplateArgs, SYCLIntelSchedulerTargetFmaxMhz, New);
871891
continue;
872892
}
@@ -884,7 +904,7 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs,
884904
}
885905
if (const auto *SYCLIntelNoGlobalWorkOffset =
886906
dyn_cast<SYCLIntelNoGlobalWorkOffsetAttr>(TmplAttr)) {
887-
instantiateIntelSYCLFunctionAttr<SYCLIntelNoGlobalWorkOffsetAttr>(
907+
instantiateSYCLIntelNoGlobalWorkOffsetAttr(
888908
*this, TemplateArgs, SYCLIntelNoGlobalWorkOffset, New);
889909
continue;
890910
}

‎clang/test/Driver/sycl-device.cpp

Lines changed: 21 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -35,3 +35,24 @@
3535
// RUN: %clang_cl -### -fsycl -fsycl-device-only -o dummy.out %s 2>&1 \
3636
// RUN: | FileCheck -check-prefix=CHECK-OUTPUT-FILE %s
3737
// CHECK-OUTPUT-FILE: clang{{.*}} "-o" "dummy.out"
38+
39+
/// -fsycl-device-only with preprocessing should only do the device compile
40+
// RUN: %clang -ccc-print-phases -E -fsycl -fsycl-device-only %s 2>&1 \
41+
// RUN: | FileCheck -check-prefix=PHASES-PREPROCESS %s
42+
// RUN: %clang_cl -ccc-print-phases -E -fsycl -fsycl-device-only %s 2>&1 \
43+
// RUN: | FileCheck -check-prefix=PHASES-PREPROCESS %s
44+
// RUN: %clang_cl -ccc-print-phases -P -fsycl -fsycl-device-only %s 2>&1 \
45+
// RUN: | FileCheck -check-prefix=PHASES-PREPROCESS %s
46+
// RUN: %clang_cl -ccc-print-phases -EP -fsycl -fsycl-device-only %s 2>&1 \
47+
// RUN: | FileCheck -check-prefix=PHASES-PREPROCESS %s
48+
// PHASES-PREPROCESS: 0: input, {{.*}}, c++, (device-sycl)
49+
// PHASES-PREPROCESS: 1: preprocessor, {0}, c++-cpp-output, (device-sycl)
50+
// PHASES-PREPROCESS: 2: offload, "device-sycl (spir64-unknown-unknown-sycldevice)" {1}, c++-cpp-output
51+
52+
// RUN: %clang -ccc-print-phases -MM -fsycl -fsycl-device-only %s 2>&1 \
53+
// RUN: | FileCheck -check-prefix=PHASES-PREPROC-DEPS %s
54+
// RUN: %clang -ccc-print-phases -M -fsycl -fsycl-device-only %s 2>&1 \
55+
// RUN: | FileCheck -check-prefix=PHASES-PREPROC-DEPS %s
56+
// PHASES-PREPROC-DEPS: 0: input, {{.*}}, c++, (device-sycl)
57+
// PHASES-PROPROC-DEPS: 1: preprocessor, {0}, dependencies, (device-sycl)
58+
// PHASES-PREPROC-DEPS: 2: offload, "device-sycl (spir64-unknown-unknown-sycldevice)" {1}, dependencies

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -50,9 +50,9 @@ int main() {
5050
[[intel::no_global_work_offset(1)]] int a;
5151
});
5252

53-
// expected-warning@+2{{attribute 'no_global_work_offset' is already applied}}
5453
h.single_task<class test_kernel7>(
55-
[]() [[intel::no_global_work_offset(0), intel::no_global_work_offset(1)]]{});
54+
[]() [[intel::no_global_work_offset(0), // expected-note {{previous attribute is here}}
55+
intel::no_global_work_offset(1)]]{}); // expected-warning{{attribute 'no_global_work_offset' is already applied with different parameters}}
5656
});
5757
return 0;
5858
}

‎clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp

Lines changed: 8 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,12 @@
66
[[intelfpga::scheduler_target_fmax_mhz(2)]] void
77
func() {}
88

9+
[[intel::scheduler_target_fmax_mhz(12)]] void bar();
10+
[[intel::scheduler_target_fmax_mhz(12)]] void bar() {} // OK
11+
12+
[[intel::scheduler_target_fmax_mhz(12)]] void baz(); // expected-note {{previous attribute is here}}
13+
[[intel::scheduler_target_fmax_mhz(100)]] void baz(); // expected-warning {{attribute 'scheduler_target_fmax_mhz' is already applied with different parameters}}
14+
915
template <int N>
1016
[[intel::scheduler_target_fmax_mhz(N)]] void zoo() {}
1117

@@ -47,5 +53,6 @@ int main() {
4753
[]() [[intel::scheduler_target_fmax_mhz(-4)]]{}); // expected-error{{'scheduler_target_fmax_mhz' attribute requires a non-negative integral compile time constant expression}}
4854

4955
cl::sycl::kernel_single_task<class test_kernel6>(
50-
[]() [[intel::scheduler_target_fmax_mhz(1), intel::scheduler_target_fmax_mhz(2)]]{}); // expected-warning{{attribute 'scheduler_target_fmax_mhz' is already applied with different arguments}}
56+
[]() [[intel::scheduler_target_fmax_mhz(1), // expected-note {{previous attribute is here}}
57+
intel::scheduler_target_fmax_mhz(2)]]{}); // expected-warning{{attribute 'scheduler_target_fmax_mhz' is already applied with different arguments}}
5158
}

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

Lines changed: 21 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -36,6 +36,18 @@ int main() {
3636
KernelFunctor<1>();
3737
}
3838

39+
[[intel::no_global_work_offset]] void func3 ();
40+
[[intel::no_global_work_offset(1)]] void func3() {} // OK
41+
42+
[[intel::no_global_work_offset(0)]] void func4(); // expected-note {{previous attribute is here}}
43+
[[intel::no_global_work_offset]] void func4(); // expected-warning{{attribute 'no_global_work_offset' is already applied with different parameters}}
44+
45+
[[intel::no_global_work_offset(1)]] void func5();
46+
[[intel::no_global_work_offset(1)]] void func5() {} // OK
47+
48+
[[intel::no_global_work_offset(0)]] void func6(); // expected-note {{previous attribute is here}}
49+
[[intel::no_global_work_offset(1)]] void func6(); // expected-warning{{attribute 'no_global_work_offset' is already applied with different parameters}}
50+
3951
// CHECK: ClassTemplateDecl {{.*}} {{.*}} KernelFunctor
4052
// CHECK: ClassTemplateSpecializationDecl {{.*}} {{.*}} class KernelFunctor definition
4153
// CHECK: CXXRecordDecl {{.*}} {{.*}} implicit class KernelFunctor
@@ -48,14 +60,20 @@ int main() {
4860

4961
// Test that checks template parameter suppport on function.
5062
template <int N>
51-
[[intel::no_global_work_offset(N)]] void func3() {}
63+
[[intel::no_global_work_offset(N)]] void func6() {}
64+
65+
template <int N>
66+
[[intel::no_global_work_offset(0)]] void func7(); // expected-note {{previous attribute is here}}
67+
template <int N>
68+
[[intel::no_global_work_offset(N)]] void func7() {} // expected-warning {{attribute 'no_global_work_offset' is already applied with different parameters}}
5269

5370
int check() {
54-
func3<1>();
71+
func6<1>();
72+
func7<1>(); //expected-note {{in instantiation of function template specialization 'func7<1>' requested here}}
5573
return 0;
5674
}
5775

58-
// CHECK: FunctionDecl {{.*}} {{.*}} func3 'void ()'
76+
// CHECK: FunctionDecl {{.*}} {{.*}} func6 'void ()'
5977
// CHECK: TemplateArgument integral 1
6078
// CHECK: SYCLIntelNoGlobalWorkOffsetAttr {{.*}}
6179
// CHECK-NEXT: ConstantExpr {{.*}} 'int'

0 commit comments

Comments
 (0)