Skip to content

Commit 299796a

Browse files
smanna12smanna12
authored
[SYCL][NFC] Update tests for FPGA attributes (#3632)
This patch 1. updates CodeGen tests for disable_loop_pipelining and initiation_interval attributes. 2. rewrites Sema and CodeGen tests for scheduler-target-fmax-mhz attribute according to FE test guidelines and also separates AST and diagnostic parts for scheduler-target-fmax-mhz attribute to avoid unrelated changes on #3601. Signed-off-by: Soumi Manna <[email protected]>
1 parent d53742a commit 299796a

7 files changed

+240
-170
lines changed

‎clang/test/CodeGenSYCL/disable_loop_pipelining.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -29,8 +29,8 @@ int main() {
2929
return 0;
3030
}
3131

32-
// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel1"() #0 !kernel_arg_buffer_location ![[NUM4:[0-9]+]] !disable_loop_pipelining ![[NUM5:[0-9]+]]
33-
// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel2"() #0 !kernel_arg_buffer_location ![[NUM4]]
34-
// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel3"() #0 !kernel_arg_buffer_location ![[NUM4]] !disable_loop_pipelining ![[NUM5]]
32+
// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel1"() #0 {{.*}} !disable_loop_pipelining ![[NUM5:[0-9]+]]
33+
// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel2"() #0 {{.*}} ![[NUM4:[0-9]+]]
34+
// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel3"() #0 {{.*}} !disable_loop_pipelining ![[NUM5]]
3535
// CHECK: ![[NUM4]] = !{}
3636
// CHECK: ![[NUM5]] = !{i32 1}

‎clang/test/CodeGenSYCL/initiation_interval.cpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -39,10 +39,10 @@ int main() {
3939
return 0;
4040
}
4141

42-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1"() #0 !kernel_arg_buffer_location ![[NUM0:[0-9]+]] !initiation_interval ![[NUM1:[0-9]+]]
43-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 !kernel_arg_buffer_location ![[NUM0]] !initiation_interval ![[NUM42:[0-9]+]]
44-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 !kernel_arg_buffer_location ![[NUM0]] !initiation_interval ![[NUM2:[0-9]+]]
45-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 !kernel_arg_buffer_location ![[NUM0]]
42+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1"() #0 {{.*}} !initiation_interval ![[NUM1:[0-9]+]]
43+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !initiation_interval ![[NUM42:[0-9]+]]
44+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} !initiation_interval ![[NUM2:[0-9]+]]
45+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} ![[NUM0:[0-9]+]]
4646
// CHECK: ![[NUM0]] = !{}
4747
// CHECK: ![[NUM1]] = !{i32 1}
4848
// CHECK: ![[NUM42]] = !{i32 42}

‎clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp

Lines changed: 49 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -1,25 +1,58 @@
1-
// RUN: %clang_cc1 -fsycl-is-device -disable-llvm-passes -triple spir64-unknown-unknown-sycldevice -emit-llvm -o - %s | FileCheck %s
1+
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -sycl-std=2020 -emit-llvm -o - %s | FileCheck %s
22

3-
#include "Inputs/sycl.hpp"
4-
[[intel::scheduler_target_fmax_mhz(5)]] void
5-
func() {}
3+
#include "sycl.hpp"
4+
5+
using namespace cl::sycl;
6+
queue q;
7+
8+
class Foo {
9+
public:
10+
[[intel::scheduler_target_fmax_mhz(5)]] void operator()() const {}
11+
};
12+
13+
template <int N>
14+
class Functor {
15+
public:
16+
[[intel::scheduler_target_fmax_mhz(N)]] void operator()() const {}
17+
};
618

719
template <int N>
820
[[intel::scheduler_target_fmax_mhz(N)]] void zoo() {}
921

22+
[[intel::scheduler_target_fmax_mhz(2)]] void bar() {}
23+
1024
int main() {
11-
cl::sycl::kernel_single_task<class test_kernel1>(
12-
[]() [[intel::scheduler_target_fmax_mhz(2)]]{});
25+
q.submit([&](handler &h) {
26+
// Test attribute argument size.
27+
Foo boo;
28+
h.single_task<class kernel_name1>(boo);
29+
30+
// Test attribute is applied on lambda.
31+
h.single_task<class kernel_name2>(
32+
[]() [[intel::scheduler_target_fmax_mhz(42)]]{});
1333

14-
cl::sycl::kernel_single_task<class test_kernel2>(
15-
[]() { func(); });
34+
// Test class template argument.
35+
Functor<7> f;
36+
h.single_task<class kernel_name3>(f);
1637

17-
cl::sycl::kernel_single_task<class test_kernel3>(
18-
[]() { zoo<75>(); });
38+
// Test attribute is propagated.
39+
h.single_task<class kernel_name4>(
40+
[]() { bar(); });
41+
42+
// Test function template argument.
43+
h.single_task<class kernel_name5>(
44+
[]() { zoo<75>(); });
45+
});
46+
return 0;
1947
}
20-
// CHECK: define {{.*}}spir_kernel void @{{.*}}test_kernel1() {{.*}} !scheduler_target_fmax_mhz ![[PARAM1:[0-9]+]]
21-
// CHECK: define {{.*}}spir_kernel void @{{.*}}test_kernel2() {{.*}} !scheduler_target_fmax_mhz ![[PARAM2:[0-9]+]]
22-
// CHECK: define {{.*}}spir_kernel void @{{.*}}test_kernel3() {{.*}} !scheduler_target_fmax_mhz ![[PARAM3:[0-9]+]]
23-
// CHECK: ![[PARAM1]] = !{i32 2}
24-
// CHECK: ![[PARAM2]] = !{i32 5}
25-
// CHECK: ![[PARAM3]] = !{i32 75}
48+
49+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1"() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM5:[0-9]+]]
50+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM42:[0-9]+]]
51+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM7:[0-9]+]]
52+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM2:[0-9]+]]
53+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5"() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM75:[0-9]+]]
54+
// CHECK: ![[NUM5]] = !{i32 5}
55+
// CHECK: ![[NUM42]] = !{i32 42}
56+
// CHECK: ![[NUM7]] = !{i32 7}
57+
// CHECK: ![[NUM2]] = !{i32 2}
58+
// CHECK: ![[NUM75]] = !{i32 75}

‎clang/test/SemaSYCL/initiation_interval.cpp

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,8 +1,10 @@
11
// RUN: %clang_cc1 -fsycl-is-device -verify %s
22

3-
// Test that checks disable_loop_pipelining attribute support on Function.
3+
// Test that checks initiation_interval attribute support on function.
44

55
// Tests for incorrect argument values for Intel FPGA initiation_interval function attribute.
6+
[[intel::initiation_interval]] void one() {} // expected-error {{'initiation_interval' attribute takes one argument}}
7+
68
[[intel::initiation_interval(5)]] int a; // expected-error{{'initiation_interval' attribute only applies to 'for', 'while', 'do' statements, and functions}}
79

810
[[intel::initiation_interval("foo")]] void func() {} // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'const char [4]'}}
@@ -27,7 +29,7 @@
2729
[[intel::initiation_interval(1)]] void func6(); // expected-note {{previous attribute is here}}
2830
[[intel::initiation_interval(3)]] void func6(); // expected-warning {{attribute 'initiation_interval' is already applied with different arguments}}
2931

30-
// Tests for Intel FPGA loop fusion function attributes compatibility
32+
// Tests for Intel FPGA initiation_interval and disable_loop_pipelining attributes compatibility checks.
3133
// expected-error@+2 {{'initiation_interval' and 'disable_loop_pipelining' attributes are not compatible}}
3234
// expected-note@+1 {{conflicting attribute is here}}
3335
[[intel::disable_loop_pipelining]] [[intel::initiation_interval(2)]] void func7();

‎clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp

Lines changed: 69 additions & 50 deletions
Original file line numberDiff line numberDiff line change
@@ -1,60 +1,79 @@
1-
// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -triple spir64 -Wno-sycl-2017-compat -verify | FileCheck %s
1+
// RUN: %clang_cc1 -fsycl-is-device -verify %s
22

3-
#include "Inputs/sycl.hpp"
3+
// Test that checks scheduler_target_fmax_mhz attribute support on Function.
4+
5+
// Test for deprecated spelling of scheduler_target_fmax_mhz attribute.
46
// expected-warning@+2 {{attribute 'intelfpga::scheduler_target_fmax_mhz' is deprecated}}
57
// expected-note@+1 {{did you mean to use 'intel::scheduler_target_fmax_mhz' instead?}}
6-
[[intelfpga::scheduler_target_fmax_mhz(2)]] void
7-
func() {}
8+
[[intelfpga::scheduler_target_fmax_mhz(2)]] void deprecate() {}
9+
10+
// Tests for incorrect argument values for Intel FPGA scheduler_target_fmax_mhz function attribute.
11+
[[intel::scheduler_target_fmax_mhz(0)]] int Var = 0; // expected-error{{'scheduler_target_fmax_mhz' attribute only applies to functions}}
12+
13+
[[intel::scheduler_target_fmax_mhz(1048577)]] void correct() {} // OK
14+
15+
[[intel::scheduler_target_fmax_mhz("foo")]] void func() {} // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'const char [4]'}}
16+
17+
[[intel::scheduler_target_fmax_mhz(-1)]] void func1() {} // expected-error{{'scheduler_target_fmax_mhz' attribute requires a non-negative integral compile time constant expression}}
18+
19+
[[intel::scheduler_target_fmax_mhz(0, 1)]] void func2() {} // expected-error{{'scheduler_target_fmax_mhz' attribute takes one argument}}
20+
21+
// Tests for Intel FPGA scheduler_target_fmax_mhz function attribute duplication.
22+
// No diagnostic is emitted because the arguments match. Duplicate attribute is silently ignored.
23+
[[intel::scheduler_target_fmax_mhz(2)]]
24+
[[intel::scheduler_target_fmax_mhz(2)]] void
25+
func3() {}
826

927
// No diagnostic is emitted because the arguments match.
10-
[[intel::scheduler_target_fmax_mhz(12)]] void bar();
11-
[[intel::scheduler_target_fmax_mhz(12)]] void bar() {} // OK
28+
[[intel::scheduler_target_fmax_mhz(4)]] void func4();
29+
[[intel::scheduler_target_fmax_mhz(4)]] void func4(); // OK
1230

1331
// Diagnostic is emitted because the arguments mismatch.
14-
[[intel::scheduler_target_fmax_mhz(12)]] void baz(); // expected-note {{previous attribute is here}}
15-
[[intel::scheduler_target_fmax_mhz(100)]] void baz(); // expected-warning {{attribute 'scheduler_target_fmax_mhz' is already applied with different arguments}}
32+
[[intel::scheduler_target_fmax_mhz(2)]] // expected-note {{previous attribute is here}}
33+
[[intel::scheduler_target_fmax_mhz(4)]] void
34+
func5() {} // expected-warning@-1 {{attribute 'scheduler_target_fmax_mhz' is already applied with different arguments}}
35+
36+
[[intel::scheduler_target_fmax_mhz(1)]] void func6(); // expected-note {{previous attribute is here}}
37+
[[intel::scheduler_target_fmax_mhz(3)]] void func6(); // expected-warning {{attribute 'scheduler_target_fmax_mhz' is already applied with different arguments}}
1638

39+
// Tests that check template parameter support for Intel FPGA scheduler_target_fmax_mhz function attributes.
1740
template <int N>
18-
[[intel::scheduler_target_fmax_mhz(N)]] void zoo() {}
19-
20-
int main() {
21-
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel1 'void ()'
22-
// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
23-
// CHECK-NEXT: ConstantExpr {{.*}} 'int'
24-
// CHECK-NEXT: value: Int 5
25-
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 5
26-
// expected-warning@+3 {{attribute 'intelfpga::scheduler_target_fmax_mhz' is deprecated}}
27-
// expected-note@+2 {{did you mean to use 'intel::scheduler_target_fmax_mhz' instead?}}
28-
cl::sycl::kernel_single_task<class test_kernel1>(
29-
[]() [[intelfpga::scheduler_target_fmax_mhz(5)]]{});
30-
31-
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel2 'void ()'
32-
// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
33-
// CHECK-NEXT: ConstantExpr {{.*}} 'int'
34-
// CHECK-NEXT: value: Int 2
35-
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2
36-
cl::sycl::kernel_single_task<class test_kernel2>(
37-
[]() { func(); });
38-
39-
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 'void ()'
40-
// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
41-
// CHECK-NEXT: ConstantExpr {{.*}} 'int'
42-
// CHECK-NEXT: value: Int 75
43-
// CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}} 'int'
44-
// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} referenced 'int' depth 0 index 0 N
45-
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 75
46-
cl::sycl::kernel_single_task<class test_kernel3>(
47-
[]() { zoo<75>(); });
48-
49-
[[intel::scheduler_target_fmax_mhz(0)]] int Var = 0; // expected-error{{'scheduler_target_fmax_mhz' attribute only applies to functions}}
50-
51-
cl::sycl::kernel_single_task<class test_kernel4>(
52-
[]() [[intel::scheduler_target_fmax_mhz(1048577)]]{}); // OK
53-
54-
cl::sycl::kernel_single_task<class test_kernel5>(
55-
[]() [[intel::scheduler_target_fmax_mhz(-4)]]{}); // expected-error{{'scheduler_target_fmax_mhz' attribute requires a non-negative integral compile time constant expression}}
56-
57-
cl::sycl::kernel_single_task<class test_kernel6>(
58-
[]() [[intel::scheduler_target_fmax_mhz(1), // expected-note {{previous attribute is here}}
59-
intel::scheduler_target_fmax_mhz(2)]]{}); // expected-warning{{attribute 'scheduler_target_fmax_mhz' is already applied with different arguments}}
41+
[[intel::scheduler_target_fmax_mhz(N)]] void func7(); // expected-error {{'scheduler_target_fmax_mhz' attribute requires a non-negative integral compile time constant expression}}
42+
43+
template <int size>
44+
[[intel::scheduler_target_fmax_mhz(10)]] void func8(); // expected-note {{previous attribute is here}}
45+
template <int size>
46+
[[intel::scheduler_target_fmax_mhz(size)]] void func8() {} // expected-warning {{attribute 'scheduler_target_fmax_mhz' is already applied with different arguments}}
47+
48+
void checkTemplates() {
49+
func7<4>(); // OK
50+
func7<-1>(); // expected-note {{in instantiation of function template specialization 'func7<-1>' requested here}}
51+
func7<0>(); // OK
52+
func8<20>(); // expected-note {{in instantiation of function template specialization 'func8<20>' requested here}}
53+
}
54+
55+
// Test that checks expression is not a constant expression.
56+
// expected-note@+1{{declared here}}
57+
int baz();
58+
// expected-error@+2{{expression is not an integral constant expression}}
59+
// expected-note@+1{{non-constexpr function 'baz' cannot be used in a constant expression}}
60+
[[intel::scheduler_target_fmax_mhz(baz() + 1)]] void func9();
61+
62+
// Test that checks expression is a constant expression.
63+
constexpr int bar() { return 0; }
64+
[[intel::scheduler_target_fmax_mhz(bar() + 2)]] void func10(); // OK
65+
66+
// Test that checks wrong function template instantiation and ensures that the type
67+
// is checked properly when instantiating from the template definition.
68+
template <typename Ty>
69+
// expected-error@+2 {{integral constant expression must have integral or unscoped enumeration type, not 'S'}}
70+
// expected-error@+1 {{integral constant expression must have integral or unscoped enumeration type, not 'float'}}
71+
[[intel::scheduler_target_fmax_mhz(Ty{})]] void func11() {}
72+
73+
struct S {};
74+
void test() {
75+
//expected-note@+1{{in instantiation of function template specialization 'func11<S>' requested here}}
76+
func11<S>();
77+
//expected-note@+1{{in instantiation of function template specialization 'func11<float>' requested here}}
78+
func11<float>();
6079
}

‎clang/test/SemaSYCL/scheduler_target_fmax_mhz_ast.cpp

Lines changed: 111 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,111 @@
1+
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -ast-dump %s | FileCheck %s
2+
3+
// Tests for AST of Intel FPGA scheduler_target_fmax_mhz function attribute.
4+
#include "sycl.hpp"
5+
6+
sycl::queue deviceQueue;
7+
8+
// CHECK: FunctionDecl {{.*}} func1 'void ()'
9+
// CHECK-NEXT: CompoundStmt
10+
// CHECK-NEXT: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
11+
// CHECK-NEXT: ConstantExpr{{.*}}'int'
12+
// CHECK-NEXT: value: Int 4
13+
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 4
14+
[[intel::scheduler_target_fmax_mhz(4)]] void func1() {}
15+
16+
// Test that checks template parameter support on function.
17+
// CHECK: FunctionTemplateDecl {{.*}} func2
18+
// CHECK: FunctionDecl {{.*}} func2 'void ()'
19+
// CHECK-NEXT: CompoundStmt
20+
// CHECK_NEXT: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
21+
// CHECK_NEXT: DeclRefExpr {{.*}} 'int' NonTypeTemplateParm {{.*}} 'N' 'int'
22+
// CHECK: FunctionDecl {{.*}} func2 'void ()'
23+
// CHECK-NEXT: TemplateArgument integral 6
24+
// CHECK-NEXT: CompoundStmt
25+
// CHECK-NEXT: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
26+
// CHECK-NEXT: ConstantExpr{{.*}}'int'
27+
// CHECK-NEXT: value: Int 6
28+
// CHECK-NEXT: SubstNonTypeTemplateParmExpr
29+
// CHECK-NEXT: NonTypeTemplateParmDecl
30+
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 6
31+
template <int N>
32+
[[intel::scheduler_target_fmax_mhz(N)]] void func2() {}
33+
34+
template <int N>
35+
[[intel::scheduler_target_fmax_mhz(N)]] void func3() {}
36+
37+
// No diagnostic is emitted because the arguments match. Duplicate attribute is silently ignored.
38+
// CHECK: FunctionDecl {{.*}} {{.*}} func4 'void ()'
39+
// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
40+
// CHECK-NEXT: ConstantExpr {{.*}} 'int'
41+
// CHECK-NEXT: value: Int 10
42+
// CHECK-NEXT: IntegerLiteral{{.*}}10{{$}}
43+
[[intel::scheduler_target_fmax_mhz(10)]]
44+
[[intel::scheduler_target_fmax_mhz(10)]] void
45+
func4() {}
46+
47+
class KernelFunctor {
48+
public:
49+
void operator()() const {
50+
func1();
51+
}
52+
};
53+
54+
// Test that checks template parameter support on class member function.
55+
template <int N>
56+
class KernelFunctor2 {
57+
public:
58+
[[intel::scheduler_target_fmax_mhz(N)]] void operator()() const {
59+
}
60+
};
61+
62+
int main() {
63+
deviceQueue.submit([&](sycl::handler &h) {
64+
// CHECK-LABEL: FunctionDecl {{.*}}kernel_name_1
65+
// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr
66+
KernelFunctor f1;
67+
h.single_task<class kernel_name_1>(f1);
68+
69+
// CHECK-LABEL: FunctionDecl {{.*}}kernel_name_2
70+
// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
71+
// CHECK-NEXT: ConstantExpr{{.*}}'int'
72+
// CHECK-NEXT: value: Int 3
73+
// CHECK-NEXT: SubstNonTypeTemplateParmExpr
74+
// CHECK-NEXT: NonTypeTemplateParmDecl
75+
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 3
76+
KernelFunctor2<3> f2;
77+
h.single_task<class kernel_name_2>(f2);
78+
79+
// CHECK-LABEL: FunctionDecl {{.*}}kernel_name_3
80+
// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
81+
// CHECK-NEXT: ConstantExpr{{.*}}'int'
82+
// CHECK-NEXT: value: Int 4
83+
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 4
84+
h.single_task<class kernel_name_3>(
85+
[]() [[intel::scheduler_target_fmax_mhz(4)]]{});
86+
87+
// CHECK-LABEL: FunctionDecl {{.*}}kernel_name_4
88+
// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
89+
// CHECK-NEXT: ConstantExpr {{.*}} 'int'
90+
// CHECK-NEXT: value: Int 75
91+
// CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}} 'int'
92+
// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} referenced 'int' depth 0 index 0 N
93+
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 75
94+
h.single_task<class kernel_name_4>(
95+
[]() { func3<75>(); });
96+
97+
// Ignore duplicate attribute.
98+
h.single_task<class kernel_name_5>(
99+
// CHECK-LABEL: FunctionDecl {{.*}}kernel_name_5
100+
// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
101+
// CHECK-NEXT: ConstantExpr {{.*}} 'int'
102+
// CHECK-NEXT: value: Int 6
103+
// CHECK-NEXT: IntegerLiteral{{.*}}6{{$}}
104+
[]() [[intel::scheduler_target_fmax_mhz(6),
105+
intel::scheduler_target_fmax_mhz(6)]]{});
106+
});
107+
108+
func2<6>();
109+
110+
return 0;
111+
}

0 commit comments

Comments
 (0)