Skip to content

[SYCL][NFC] Update tests #3632

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.

Sign up for GitHub

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 4 commits into from
Apr 28, 2021
Merged

[SYCL][NFC] Update tests #3632

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
1bec1d7
[SYCL][NFC] Update tests
smanna12 Apr 27, 2021
05fe9bb
fix clang format issues
smanna12 Apr 27, 2021
4f0d9e3
fix test
smanna12 Apr 27, 2021
609d517
fix test
smanna12 Apr 27, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 3 additions & 3 deletions clang/test/CodeGenSYCL/disable_loop_pipelining.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -29,8 +29,8 @@ int main() {
return 0;
}

// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel1"() #0 !kernel_arg_buffer_location ![[NUM4:[0-9]+]] !disable_loop_pipelining ![[NUM5:[0-9]+]]
// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel2"() #0 !kernel_arg_buffer_location ![[NUM4]]
// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel3"() #0 !kernel_arg_buffer_location ![[NUM4]] !disable_loop_pipelining ![[NUM5]]
// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel1"() #0 {{.*}} !disable_loop_pipelining ![[NUM5:[0-9]+]]
// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel2"() #0 {{.*}} ![[NUM4:[0-9]+]]
// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel3"() #0 {{.*}} !disable_loop_pipelining ![[NUM5]]
// CHECK: ![[NUM4]] = !{}
// CHECK: ![[NUM5]] = !{i32 1}
8 changes: 4 additions & 4 deletions clang/test/CodeGenSYCL/initiation_interval.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -39,10 +39,10 @@ int main() {
return 0;
}

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1"() #0 !kernel_arg_buffer_location ![[NUM0:[0-9]+]] !initiation_interval ![[NUM1:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 !kernel_arg_buffer_location ![[NUM0]] !initiation_interval ![[NUM42:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 !kernel_arg_buffer_location ![[NUM0]] !initiation_interval ![[NUM2:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 !kernel_arg_buffer_location ![[NUM0]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1"() #0 {{.*}} !initiation_interval ![[NUM1:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !initiation_interval ![[NUM42:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} !initiation_interval ![[NUM2:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} ![[NUM0:[0-9]+]]
// CHECK: ![[NUM0]] = !{}
// CHECK: ![[NUM1]] = !{i32 1}
// CHECK: ![[NUM42]] = !{i32 42}
Expand Down
65 changes: 49 additions & 16 deletions clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,25 +1,58 @@
// RUN: %clang_cc1 -fsycl-is-device -disable-llvm-passes -triple spir64-unknown-unknown-sycldevice -emit-llvm -o - %s | FileCheck %s
// 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

#include "Inputs/sycl.hpp"
[[intel::scheduler_target_fmax_mhz(5)]] void
func() {}
#include "sycl.hpp"

using namespace cl::sycl;
queue q;

class Foo {
public:
[[intel::scheduler_target_fmax_mhz(5)]] void operator()() const {}
};

template <int N>
class Functor {
public:
[[intel::scheduler_target_fmax_mhz(N)]] void operator()() const {}
};

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

[[intel::scheduler_target_fmax_mhz(2)]] void bar() {}

int main() {
cl::sycl::kernel_single_task<class test_kernel1>(
[]() [[intel::scheduler_target_fmax_mhz(2)]]{});
q.submit([&](handler &h) {
// Test attribute argument size.
Foo boo;
h.single_task<class kernel_name1>(boo);

// Test attribute is applied on lambda.
h.single_task<class kernel_name2>(
[]() [[intel::scheduler_target_fmax_mhz(42)]]{});

cl::sycl::kernel_single_task<class test_kernel2>(
[]() { func(); });
// Test class template argument.
Functor<7> f;
h.single_task<class kernel_name3>(f);

cl::sycl::kernel_single_task<class test_kernel3>(
[]() { zoo<75>(); });
// Test attribute is propagated.
h.single_task<class kernel_name4>(
[]() { bar(); });

// Test function template argument.
h.single_task<class kernel_name5>(
[]() { zoo<75>(); });
});
return 0;
}
// CHECK: define {{.*}}spir_kernel void @{{.*}}test_kernel1() {{.*}} !scheduler_target_fmax_mhz ![[PARAM1:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}test_kernel2() {{.*}} !scheduler_target_fmax_mhz ![[PARAM2:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}test_kernel3() {{.*}} !scheduler_target_fmax_mhz ![[PARAM3:[0-9]+]]
// CHECK: ![[PARAM1]] = !{i32 2}
// CHECK: ![[PARAM2]] = !{i32 5}
// CHECK: ![[PARAM3]] = !{i32 75}

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1"() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM5:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM42:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM7:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM2:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5"() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM75:[0-9]+]]
// CHECK: ![[NUM5]] = !{i32 5}
// CHECK: ![[NUM42]] = !{i32 42}
// CHECK: ![[NUM7]] = !{i32 7}
// CHECK: ![[NUM2]] = !{i32 2}
// CHECK: ![[NUM75]] = !{i32 75}
6 changes: 4 additions & 2 deletions clang/test/SemaSYCL/initiation_interval.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,8 +1,10 @@
// RUN: %clang_cc1 -fsycl-is-device -verify %s

// Test that checks disable_loop_pipelining attribute support on Function.
// Test that checks initiation_interval attribute support on function.

// Tests for incorrect argument values for Intel FPGA initiation_interval function attribute.
[[intel::initiation_interval]] void one() {} // expected-error {{'initiation_interval' attribute takes one argument}}

[[intel::initiation_interval(5)]] int a; // expected-error{{'initiation_interval' attribute only applies to 'for', 'while', 'do' statements, and functions}}

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

// Tests for Intel FPGA loop fusion function attributes compatibility
// Tests for Intel FPGA initiation_interval and disable_loop_pipelining attributes compatibility checks.
// expected-error@+2 {{'initiation_interval' and 'disable_loop_pipelining' attributes are not compatible}}
// expected-note@+1 {{conflicting attribute is here}}
[[intel::disable_loop_pipelining]] [[intel::initiation_interval(2)]] void func7();
Expand Down
119 changes: 69 additions & 50 deletions clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,60 +1,79 @@
// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -triple spir64 -Wno-sycl-2017-compat -verify | FileCheck %s
// RUN: %clang_cc1 -fsycl-is-device -verify %s

#include "Inputs/sycl.hpp"
// Test that checks scheduler_target_fmax_mhz attribute support on Function.

// Test for deprecated spelling of scheduler_target_fmax_mhz attribute.
// expected-warning@+2 {{attribute 'intelfpga::scheduler_target_fmax_mhz' is deprecated}}
// expected-note@+1 {{did you mean to use 'intel::scheduler_target_fmax_mhz' instead?}}
[[intelfpga::scheduler_target_fmax_mhz(2)]] void
func() {}
[[intelfpga::scheduler_target_fmax_mhz(2)]] void deprecate() {}

// Tests for incorrect argument values for Intel FPGA scheduler_target_fmax_mhz function attribute.
[[intel::scheduler_target_fmax_mhz(0)]] int Var = 0; // expected-error{{'scheduler_target_fmax_mhz' attribute only applies to functions}}

[[intel::scheduler_target_fmax_mhz(1048577)]] void correct() {} // OK

[[intel::scheduler_target_fmax_mhz("foo")]] void func() {} // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'const char [4]'}}

[[intel::scheduler_target_fmax_mhz(-1)]] void func1() {} // expected-error{{'scheduler_target_fmax_mhz' attribute requires a non-negative integral compile time constant expression}}

[[intel::scheduler_target_fmax_mhz(0, 1)]] void func2() {} // expected-error{{'scheduler_target_fmax_mhz' attribute takes one argument}}

// Tests for Intel FPGA scheduler_target_fmax_mhz function attribute duplication.
// No diagnostic is emitted because the arguments match. Duplicate attribute is silently ignored.
[[intel::scheduler_target_fmax_mhz(2)]]
[[intel::scheduler_target_fmax_mhz(2)]] void
func3() {}

// No diagnostic is emitted because the arguments match.
[[intel::scheduler_target_fmax_mhz(12)]] void bar();
[[intel::scheduler_target_fmax_mhz(12)]] void bar() {} // OK
[[intel::scheduler_target_fmax_mhz(4)]] void func4();
[[intel::scheduler_target_fmax_mhz(4)]] void func4(); // OK

// Diagnostic is emitted because the arguments mismatch.
[[intel::scheduler_target_fmax_mhz(12)]] void baz(); // expected-note {{previous attribute is here}}
[[intel::scheduler_target_fmax_mhz(100)]] void baz(); // expected-warning {{attribute 'scheduler_target_fmax_mhz' is already applied with different arguments}}
[[intel::scheduler_target_fmax_mhz(2)]] // expected-note {{previous attribute is here}}
[[intel::scheduler_target_fmax_mhz(4)]] void
func5() {} // expected-warning@-1 {{attribute 'scheduler_target_fmax_mhz' is already applied with different arguments}}

[[intel::scheduler_target_fmax_mhz(1)]] void func6(); // expected-note {{previous attribute is here}}
[[intel::scheduler_target_fmax_mhz(3)]] void func6(); // expected-warning {{attribute 'scheduler_target_fmax_mhz' is already applied with different arguments}}

// Tests that check template parameter support for Intel FPGA scheduler_target_fmax_mhz function attributes.
template <int N>
[[intel::scheduler_target_fmax_mhz(N)]] void zoo() {}

int main() {
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel1 'void ()'
// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
// CHECK-NEXT: ConstantExpr {{.*}} 'int'
// CHECK-NEXT: value: Int 5
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 5
// expected-warning@+3 {{attribute 'intelfpga::scheduler_target_fmax_mhz' is deprecated}}
// expected-note@+2 {{did you mean to use 'intel::scheduler_target_fmax_mhz' instead?}}
cl::sycl::kernel_single_task<class test_kernel1>(
[]() [[intelfpga::scheduler_target_fmax_mhz(5)]]{});

// CHECK-LABEL: FunctionDecl {{.*}}test_kernel2 'void ()'
// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
// CHECK-NEXT: ConstantExpr {{.*}} 'int'
// CHECK-NEXT: value: Int 2
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2
cl::sycl::kernel_single_task<class test_kernel2>(
[]() { func(); });

// CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 'void ()'
// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
// CHECK-NEXT: ConstantExpr {{.*}} 'int'
// CHECK-NEXT: value: Int 75
// CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}} 'int'
// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} referenced 'int' depth 0 index 0 N
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 75
cl::sycl::kernel_single_task<class test_kernel3>(
[]() { zoo<75>(); });

[[intel::scheduler_target_fmax_mhz(0)]] int Var = 0; // expected-error{{'scheduler_target_fmax_mhz' attribute only applies to functions}}

cl::sycl::kernel_single_task<class test_kernel4>(
[]() [[intel::scheduler_target_fmax_mhz(1048577)]]{}); // OK

cl::sycl::kernel_single_task<class test_kernel5>(
[]() [[intel::scheduler_target_fmax_mhz(-4)]]{}); // expected-error{{'scheduler_target_fmax_mhz' attribute requires a non-negative integral compile time constant expression}}

cl::sycl::kernel_single_task<class test_kernel6>(
[]() [[intel::scheduler_target_fmax_mhz(1), // expected-note {{previous attribute is here}}
intel::scheduler_target_fmax_mhz(2)]]{}); // expected-warning{{attribute 'scheduler_target_fmax_mhz' is already applied with different arguments}}
[[intel::scheduler_target_fmax_mhz(N)]] void func7(); // expected-error {{'scheduler_target_fmax_mhz' attribute requires a non-negative integral compile time constant expression}}

template <int size>
[[intel::scheduler_target_fmax_mhz(10)]] void func8(); // expected-note {{previous attribute is here}}
template <int size>
[[intel::scheduler_target_fmax_mhz(size)]] void func8() {} // expected-warning {{attribute 'scheduler_target_fmax_mhz' is already applied with different arguments}}

void checkTemplates() {
func7<4>(); // OK
func7<-1>(); // expected-note {{in instantiation of function template specialization 'func7<-1>' requested here}}
func7<0>(); // OK
func8<20>(); // expected-note {{in instantiation of function template specialization 'func8<20>' requested here}}
}

// Test that checks expression is not a constant expression.
// expected-note@+1{{declared here}}
int baz();
// expected-error@+2{{expression is not an integral constant expression}}
// expected-note@+1{{non-constexpr function 'baz' cannot be used in a constant expression}}
[[intel::scheduler_target_fmax_mhz(baz() + 1)]] void func9();

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

// 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@+2 {{integral constant expression must have integral or unscoped enumeration type, not 'S'}}
// expected-error@+1 {{integral constant expression must have integral or unscoped enumeration type, not 'float'}}
[[intel::scheduler_target_fmax_mhz(Ty{})]] void func11() {}

struct S {};
void test() {
//expected-note@+1{{in instantiation of function template specialization 'func11<S>' requested here}}
func11<S>();
//expected-note@+1{{in instantiation of function template specialization 'func11<float>' requested here}}
func11<float>();
}
111 changes: 111 additions & 0 deletions clang/test/SemaSYCL/scheduler_target_fmax_mhz_ast.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,111 @@
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -ast-dump %s | FileCheck %s

// Tests for AST of Intel FPGA scheduler_target_fmax_mhz function attribute.
#include "sycl.hpp"

sycl::queue deviceQueue;

// CHECK: FunctionDecl {{.*}} func1 'void ()'
// CHECK-NEXT: CompoundStmt
// CHECK-NEXT: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
// CHECK-NEXT: ConstantExpr{{.*}}'int'
// CHECK-NEXT: value: Int 4
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 4
[[intel::scheduler_target_fmax_mhz(4)]] void func1() {}

// Test that checks template parameter support on function.
// CHECK: FunctionTemplateDecl {{.*}} func2
// CHECK: FunctionDecl {{.*}} func2 'void ()'
// CHECK-NEXT: CompoundStmt
// CHECK_NEXT: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
// CHECK_NEXT: DeclRefExpr {{.*}} 'int' NonTypeTemplateParm {{.*}} 'N' 'int'
// CHECK: FunctionDecl {{.*}} func2 'void ()'
// CHECK-NEXT: TemplateArgument integral 6
// CHECK-NEXT: CompoundStmt
// CHECK-NEXT: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
// CHECK-NEXT: ConstantExpr{{.*}}'int'
// CHECK-NEXT: value: Int 6
// CHECK-NEXT: SubstNonTypeTemplateParmExpr
// CHECK-NEXT: NonTypeTemplateParmDecl
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 6
template <int N>
[[intel::scheduler_target_fmax_mhz(N)]] void func2() {}

template <int N>
[[intel::scheduler_target_fmax_mhz(N)]] void func3() {}

// No diagnostic is emitted because the arguments match. Duplicate attribute is silently ignored.
// CHECK: FunctionDecl {{.*}} {{.*}} func4 'void ()'
// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
// CHECK-NEXT: ConstantExpr {{.*}} 'int'
// CHECK-NEXT: value: Int 10
// CHECK-NEXT: IntegerLiteral{{.*}}10{{$}}
[[intel::scheduler_target_fmax_mhz(10)]]
[[intel::scheduler_target_fmax_mhz(10)]] void
func4() {}

class KernelFunctor {
public:
void operator()() const {
func1();
}
};

// Test that checks template parameter support on class member function.
template <int N>
class KernelFunctor2 {
public:
[[intel::scheduler_target_fmax_mhz(N)]] void operator()() const {
}
};

int main() {
deviceQueue.submit([&](sycl::handler &h) {
// CHECK-LABEL: FunctionDecl {{.*}}kernel_name_1
// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr
KernelFunctor f1;
h.single_task<class kernel_name_1>(f1);

// CHECK-LABEL: FunctionDecl {{.*}}kernel_name_2
// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
// CHECK-NEXT: ConstantExpr{{.*}}'int'
// CHECK-NEXT: value: Int 3
// CHECK-NEXT: SubstNonTypeTemplateParmExpr
// CHECK-NEXT: NonTypeTemplateParmDecl
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 3
KernelFunctor2<3> f2;
h.single_task<class kernel_name_2>(f2);

// CHECK-LABEL: FunctionDecl {{.*}}kernel_name_3
// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
// CHECK-NEXT: ConstantExpr{{.*}}'int'
// CHECK-NEXT: value: Int 4
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 4
h.single_task<class kernel_name_3>(
[]() [[intel::scheduler_target_fmax_mhz(4)]]{});

// CHECK-LABEL: FunctionDecl {{.*}}kernel_name_4
// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
// CHECK-NEXT: ConstantExpr {{.*}} 'int'
// CHECK-NEXT: value: Int 75
// CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}} 'int'
// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} referenced 'int' depth 0 index 0 N
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 75
h.single_task<class kernel_name_4>(
[]() { func3<75>(); });

// Ignore duplicate attribute.
h.single_task<class kernel_name_5>(
// CHECK-LABEL: FunctionDecl {{.*}}kernel_name_5
// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
// CHECK-NEXT: ConstantExpr {{.*}} 'int'
// CHECK-NEXT: value: Int 6
// CHECK-NEXT: IntegerLiteral{{.*}}6{{$}}
[]() [[intel::scheduler_target_fmax_mhz(6),
intel::scheduler_target_fmax_mhz(6)]]{});
});

func2<6>();

return 0;
}
Loading