Skip to content

[FPGA][NFC] Refactor [[intel::max_work_group_size()]] attribute tests #6240

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 3 commits into from
Jun 3, 2022
Merged

[FPGA][NFC] Refactor [[intel::max_work_group_size()]] attribute tests #6240

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
be7cb69
[FPGA][NFC] improves tests for [[intel::max_work_group_size()]] attri…
smanna12 Jun 3, 2022
642cdea
fix clang format errors
smanna12 Jun 3, 2022
f0ab4c6
remove #if part
smanna12 Jun 3, 2022
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
131 changes: 131 additions & 0 deletions clang/test/SemaSYCL/intel-max-work-group-size-ast.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,131 @@
// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -triple spir64 | FileCheck %s

// The test checks support and functionality of [[intel:::max_work_group_size()]] attribute.
#include "sycl.hpp"

using namespace cl::sycl;
queue q;

[[intel::max_work_group_size(2, 2, 2)]] void func_do_not_ignore() {}

struct FuncObj {
[[intel::max_work_group_size(4, 4, 4)]] void operator()() const {}
};

// Test that checks template parameter support on member function of class template.
template <int SIZE, int SIZE1, int SIZE2>
class KernelFunctor {
public:
[[intel::max_work_group_size(SIZE, SIZE1, SIZE2)]] void operator()() {}
};

// Test that checks template parameter support on function.
template <int N, int N1, int N2>
[[intel::max_work_group_size(N, N1, N2)]] void func() {}

int check() {
// CHECK: ClassTemplateDecl {{.*}} {{.*}} KernelFunctor
// CHECK: ClassTemplateSpecializationDecl {{.*}} {{.*}} class KernelFunctor definition
// CHECK: CXXRecordDecl {{.*}} {{.*}} implicit class KernelFunctor
// CHECK: SYCLIntelMaxWorkGroupSizeAttr
// CHECK-NEXT: ConstantExpr{{.*}}'int'
// CHECK-NEXT: value: Int 4
// CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}}
// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}}
// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}}
// CHECK-NEXT: ConstantExpr{{.*}}'int'
// CHECK-NEXT: value: Int 4
// CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}}
// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}}
// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}}
// CHECK-NEXT: ConstantExpr{{.*}}'int'
// CHECK-NEXT: value: Int 4
// CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}}
// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}}
// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}}
KernelFunctor<4, 4, 4>();

// CHECK: FunctionTemplateDecl {{.*}} {{.*}} func
// CHECK: FunctionDecl {{.*}} {{.*}} used func 'void ()'
// CHECK: SYCLIntelMaxWorkGroupSizeAttr {{.*}}
// CHECK-NEXT: ConstantExpr{{.*}}'int'
// CHECK-NEXT: value: Int 8
// CHECK: SubstNonTypeTemplateParmExpr {{.*}}
// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}}
// CHECK-NEXT: IntegerLiteral{{.*}}8{{$}}
// CHECK-NEXT: ConstantExpr{{.*}}'int'
// CHECK-NEXT: value: Int 8
// CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}}
// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}}
// CHECK-NEXT: IntegerLiteral{{.*}}8{{$}}
// CHECK-NEXT: ConstantExpr{{.*}}'int'
// CHECK-NEXT: value: Int 8
// CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}}
// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}}
// CHECK-NEXT: IntegerLiteral{{.*}}8{{$}}
func<8, 8, 8>();
return 0;
}

int main() {
q.submit([&](handler &h) {
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel1
// CHECK: SYCLIntelMaxWorkGroupSizeAttr
// CHECK-NEXT: ConstantExpr{{.*}}'int'
// CHECK-NEXT: value: Int 4
// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}}
// CHECK-NEXT: ConstantExpr{{.*}}'int'
// CHECK-NEXT: value: Int 4
// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}}
// CHECK-NEXT: ConstantExpr{{.*}}'int'
// CHECK-NEXT: value: Int 4
// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}}
h.single_task<class test_kernel1>(FuncObj());

// CHECK-LABEL: FunctionDecl {{.*}}test_kernel2
// CHECK: SYCLIntelMaxWorkGroupSizeAttr
// CHECK-NEXT: ConstantExpr{{.*}}'int'
// CHECK-NEXT: value: Int 8
// CHECK-NEXT: IntegerLiteral{{.*}}8{{$}}
// CHECK-NEXT: ConstantExpr{{.*}}'int'
// CHECK-NEXT: value: Int 8
// CHECK-NEXT: IntegerLiteral{{.*}}8{{$}}
// CHECK-NEXT: ConstantExpr{{.*}}'int'
// CHECK-NEXT: value: Int 8
// CHECK-NEXT: IntegerLiteral{{.*}}8{{$}}
h.single_task<class test_kernel2>(
[]() [[intel::max_work_group_size(8, 8, 8)]] {});

// CHECK-LABEL: FunctionDecl {{.*}}test_kernel3
// CHECK: SYCLIntelMaxWorkGroupSizeAttr
// CHECK-NEXT: ConstantExpr{{.*}}'int'
// CHECK-NEXT: value: Int 2
// CHECK-NEXT: IntegerLiteral{{.*}}2{{$}}
// CHECK-NEXT: ConstantExpr{{.*}}'int'
// CHECK-NEXT: value: Int 2
// CHECK-NEXT: IntegerLiteral{{.*}}2{{$}}
// CHECK-NEXT: ConstantExpr{{.*}}'int'
// CHECK-NEXT: value: Int 2
// CHECK-NEXT: IntegerLiteral{{.*}}2{{$}}
h.single_task<class test_kernel3>(
[]() { func_do_not_ignore(); });

// Ignore duplicate attribute.
h.single_task<class test_kernel10>(
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel10
// CHECK: SYCLIntelMaxWorkGroupSizeAttr
// CHECK-NEXT: ConstantExpr{{.*}}'int'
// CHECK-NEXT: value: Int 2
// CHECK-NEXT: IntegerLiteral{{.*}}2{{$}}
// CHECK-NEXT: ConstantExpr{{.*}}'int'
// CHECK-NEXT: value: Int 2
// CHECK-NEXT: IntegerLiteral{{.*}}2{{$}}
// CHECK-NEXT: ConstantExpr{{.*}}'int'
// CHECK-NEXT: value: Int 2
// CHECK-NEXT: IntegerLiteral{{.*}}2{{$}}
// CHECK-NOT: SYCLIntelMaxWorkGroupSizeAttr
[]() [[intel::max_work_group_size(2, 2, 2),
intel::max_work_group_size(2, 2, 2)]] {});
});
return 0;
}
133 changes: 0 additions & 133 deletions clang/test/SemaSYCL/intel-max-work-group-size-device.cpp
View file
Open in desktop

This file was deleted.

39 changes: 39 additions & 0 deletions clang/test/SemaSYCL/intel-max-work-group-size.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,6 +5,7 @@
[[intel::max_work_group_size(12, 12, 12, 12)]] void f0(); // expected-error {{'max_work_group_size' attribute requires exactly 3 arguments}}
[[intel::max_work_group_size("derp", 1, 2)]] void f1(); // expected-error {{integral constant expression must have integral or unscoped enumeration type, not 'const char[5]'}}
[[intel::max_work_group_size(1, 1, 1)]] int i; // expected-error {{'max_work_group_size' attribute only applies to functions}}
[[intel::max_work_group_size(-8, 8, -8)]] void neg(); // expected-error 2{{'max_work_group_size' attribute requires a positive integral compile time constant expression}}

// Tests for Intel FPGA 'max_work_group_size' attribute duplication.
// No diagnostic is emitted because the arguments match. Duplicate attribute is silently ignored.
Expand Down Expand Up @@ -105,3 +106,41 @@ f18();

[[intel::max_work_group_size(16, 16, 1)]] void f19();
[[sycl::reqd_work_group_size(16, 16)]] void f19(); // 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, typename Ty1, typename Ty2>
// expected-error@+1 3{{integral constant expression must have integral or unscoped enumeration type, not 'S'}}
[[intel::max_work_group_size(Ty{}, Ty1{}, Ty2{})]] void f20() {}

struct S {};
void var() {
// expected-note@+1 {{in instantiation of function template specialization 'f20<S, S, S>' requested here}}
f20<S, S, S>();
}

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

// Test that checks expression is a constant expression.
constexpr int bar() { return 0; }
[[intel::max_work_group_size(bar() + 12, bar() + 12, bar() + 12)]] void f22(); // OK

struct DAFuncObj {
[[intel::max_work_group_size(4, 4, 4)]] // expected-note {{conflicting attribute is here}}
[[cl::reqd_work_group_size(8, 8, 4)]] // expected-error{{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}} \
// expected-warning{{attribute 'cl::reqd_work_group_size' is deprecated}} \
// expected-note{{did you mean to use 'sycl::reqd_work_group_size' instead?}}
void
operator()() const {}
};

struct Func {
// expected-warning@+1 {{unknown attribute 'max_work_group_size' ignored}}
[[intelfpga::max_work_group_size(1, 1, 1)]] void operator()() const {}
};
Loading