From 5918ff310cedd6759912c3023870b94698e5573b Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 30 Nov 2020 22:29:43 -0800 Subject: [PATCH 01/13] [SYCL] Add template parameter support for no_global_work_offset attribute 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 --- clang/include/clang/Basic/Attr.td | 2 +- clang/include/clang/Sema/Sema.h | 10 ++- clang/lib/CodeGen/CodeGenFunction.cpp | 10 ++- clang/lib/Sema/SemaDeclAttr.cpp | 15 +--- .../lib/Sema/SemaTemplateInstantiateDecl.cpp | 6 ++ .../intel-fpga-no-global-work-offset.cpp | 45 ++++++---- .../check-notdirect-attribute-propagation.cpp | 4 +- .../intel-fpga-no-global-work-offset.cpp | 83 +++++++++---------- .../redeclaration-attribute-propagation.cpp | 26 +++--- ...evice-intel-fpga-no-global-work-offset.cpp | 25 ++++++ 10 files changed, 140 insertions(+), 86 deletions(-) create mode 100644 clang/test/SemaSYCL/sycl-device-intel-fpga-no-global-work-offset.cpp diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 53a95c1c8c365..fb171859a680f 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1303,7 +1303,7 @@ def SYCLIntelMaxGlobalWorkDim : InheritableAttr { def SYCLIntelNoGlobalWorkOffset : InheritableAttr { let Spellings = [CXX11<"intelfpga","no_global_work_offset">, CXX11<"intel","no_global_work_offset">]; - let Args = [BoolArgument<"Enabled", 1>]; + let Args = [ExprArgument<"Value">, BoolArgument<"Enabled", 1>]; let LangOpts = [SYCLIsDevice, SYCLIsHost]; let Subjects = SubjectList<[Function], ErrorDiag>; let Documentation = [SYCLIntelNoGlobalWorkOffsetAttrDocs]; diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index cc49a9994b21b..51983d8276378 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -12971,7 +12971,15 @@ void Sema::addIntelSYCLSingleArgFunctionAttr(Decl *D, return; } } - if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelMaxGlobalWorkDim) { + if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset) { + if (ArgInt > 1) { + Diag(E->getExprLoc(), diag::warn_boolean_attribute_argument_is_not_valid) + << CI.getAttrName(); + return; + } + } + if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelMaxGlobalWorkDim || + CI.getParsedKind() == ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset) { if (ArgInt < 0) { Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) << CI.getAttrName() << /*non-negative*/ 1; diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index f94a6a6973b6a..3af461be136da 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -694,8 +694,14 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, if (const SYCLIntelNoGlobalWorkOffsetAttr *A = FD->getAttr()) { - if (A->getEnabled()) - Fn->setMetadata("no_global_work_offset", llvm::MDNode::get(Context, {})); + llvm::LLVMContext &Context = getLLVMContext(); + Optional ArgVal = + A->getValue()->getIntegerConstantExpr(FD->getASTContext()); + assert(ArgVal.hasValue() && "Not an integer constant expression"); + llvm::Metadata *AttrMDArgs[] = {llvm::ConstantAsMetadata::get( + Builder.getInt32(ArgVal->getSExtValue()))}; + Fn->setMetadata("no_global_work_offset", + llvm::MDNode::get(Context, AttrMDArgs)); } if (FD->hasAttr()) { diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index d1e5855ff2e0b..e4db90b8b8ec6 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -5294,24 +5294,15 @@ static void handleNoGlobalWorkOffsetAttr(Sema &S, Decl *D, checkForDuplicateAttribute(S, D, Attr); - uint32_t Enabled = 1; - if (Attr.getNumArgs()) { - const Expr *E = Attr.getArgAsExpr(0); - if (!checkUInt32Argument(S, Attr, E, Enabled, 0, - /*StrictlyUnsigned=*/true)) - return; - } - if (Enabled > 1) - S.Diag(Attr.getLoc(), diag::warn_boolean_attribute_argument_is_not_valid) - << Attr; + Expr *E = Attr.getArgAsExpr(0); if (Attr.getKind() == ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset && checkDeprecatedSYCLAttributeSpelling(S, Attr)) S.Diag(Attr.getLoc(), diag::note_spelling_suggestion) << "'intel::no_global_work_offset'"; - D->addAttr(::new (S.Context) - SYCLIntelNoGlobalWorkOffsetAttr(S.Context, Attr, Enabled)); + S.addIntelSYCLSingleArgFunctionAttr(D, Attr, + E); } /// Handle the [[intelfpga::doublepump]] and [[intelfpga::singlepump]] attributes. diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index 5130bad1f7b5c..85fd57c459bf3 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -775,6 +775,12 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs, *this, TemplateArgs, SYCLIntelMaxGlobalWorkDim, New); continue; } + if (const auto *SYCLIntelNoGlobalWorkOffset = + dyn_cast(TmplAttr)) { + instantiateIntelSYCLFunctionAttr( + *this, TemplateArgs, SYCLIntelNoGlobalWorkOffset, New); + continue; + } // Existing DLL attribute on the instantiation takes precedence. if (TmplAttr->getKind() == attr::DLLExport || TmplAttr->getKind() == attr::DLLImport) { diff --git a/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp b/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp index a4d93478134a1..22d34fa88a51e 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp @@ -1,28 +1,41 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s +// 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 + +#include "sycl.hpp" + +using namespace cl::sycl; +queue q; class Foo { public: [[intel::no_global_work_offset(1)]] void operator()() const {} }; -template -__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { - kernelFunc(); -} +template +class Functor { +public: + [[intel::no_global_work_offset(SIZE)]] void operator()() const {} +}; + +int main() { + q.submit([&](handler &h) { + Foo boo; + h.single_task(boo); -void bar() { - Foo boo; - kernel(boo); + h.single_task( + []() [[intel::no_global_work_offset(1)]]{}); - kernel( - []() [[intel::no_global_work_offset]]{}); + h.single_task( + []() [[intel::no_global_work_offset(0)]]{}); - kernel( - []() [[intel::no_global_work_offset(0)]]{}); + Functor<1> f; + h.single_task(f); + }); + return 0; } -// CHECK: define spir_kernel void @{{.*}}kernel_name1() {{.*}} !no_global_work_offset ![[NUM5:[0-9]+]] -// CHECK: define spir_kernel void @{{.*}}kernel_name2() {{.*}} !no_global_work_offset ![[NUM5]] -// CHECK: define spir_kernel void @{{.*}}kernel_name3() {{.*}} ![[NUM4:[0-9]+]] +// CHECK: define spir_kernel void @{{.*}}kernel_name1"() #0 {{.*}} !no_global_work_offset ![[NUM5:[0-9]+]] +// CHECK: define spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !no_global_work_offset ![[NUM5]] +// CHECK: define spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} ![[NUM4:[0-9]+]] +// CHECK: define spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} !no_global_work_offset ![[NUM5]] // CHECK-NOT: ![[NUM4]] = !{i32 0} -// CHECK: ![[NUM5]] = !{} +// CHECK: ![[NUM5]] = !{i32 1} diff --git a/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp b/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp index 777ef128123eb..fb4193ee95db2 100644 --- a/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp +++ b/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp @@ -3,7 +3,7 @@ // RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -triple spir64 | FileCheck %s #ifndef TRIGGER_ERROR -[[intel::no_global_work_offset]] void not_direct_one() {} // expected-no-diagnostics +[[intel::no_global_work_offset(1)]] void not_direct_one() {} // expected-no-diagnostics [[intel::reqd_sub_group_size(1)]] void func_one() { not_direct_one(); @@ -46,7 +46,7 @@ void invoke_foo2() { // CHECK-LABEL: FunctionDecl {{.*}} invoke_foo2 'void ()' // CHECK: `-FunctionDecl {{.*}}KernelName 'void ()' // CHECK: -IntelReqdSubGroupSizeAttr {{.*}} - // CHECK: `-SYCLIntelNoGlobalWorkOffsetAttr {{.*}} Enabled + // CHECK: `-SYCLIntelNoGlobalWorkOffsetAttr {{.*}} parallel_for([]() {}); #else parallel_for([]() {}); // expected-error 2 {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}} diff --git a/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp b/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp index ff816237d6fb1..938f7814ff0e1 100644 --- a/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp +++ b/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp @@ -1,51 +1,50 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-return-type -fcxx-exceptions -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s +// 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 + +#include "sycl.hpp" + +using namespace cl::sycl; +queue q; struct FuncObj { //expected-warning@+2 {{attribute 'intelfpga::no_global_work_offset' is deprecated}} //expected-note@+1 {{did you mean to use 'intel::no_global_work_offset' instead?}} - [[intelfpga::no_global_work_offset]] void operator()() {} + [[intelfpga::no_global_work_offset(1)]] void operator()() const {} }; -template -void kernel(Func kernelFunc) { - kernelFunc(); -} - int main() { - // CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}}Enabled - kernel([]() { - FuncObj(); - }); - - // CHECK: SYCLIntelNoGlobalWorkOffsetAttr - // CHECK-NOT: Enabled - kernel( - []() [[intel::no_global_work_offset(0)]]{}); - - // CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}}Enabled - // expected-warning@+2{{'no_global_work_offset' attribute should be 0 or 1. Adjusted to 1}} - kernel( - []() [[intel::no_global_work_offset(42)]]{}); - - // expected-error@+2{{'no_global_work_offset' attribute requires a non-negative integral compile time constant expression}} - kernel( - []() [[intel::no_global_work_offset(-1)]]{}); - - // expected-error@+2{{'no_global_work_offset' attribute requires parameter 0 to be an integer constant}} - kernel( - []() [[intel::no_global_work_offset("foo")]]{}); - - kernel([]() { - // expected-error@+1{{'no_global_work_offset' attribute only applies to functions}} - [[intel::no_global_work_offset(1)]] int a; - }); - - // CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}} - // CHECK-NOT: Enabled - // CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}}Enabled - // expected-warning@+2{{attribute 'no_global_work_offset' is already applied}} - kernel( - []() [[intel::no_global_work_offset(0), intel::no_global_work_offset(1)]]{}); - + q.submit([&](handler &h) { + // CHECK: SYCLIntelNoGlobalWorkOffsetAttr {{.*}} + // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + h.single_task(FuncObj()); + + // CHECK: SYCLIntelNoGlobalWorkOffsetAttr {{.*}} + // CHECK-NEXT: IntegerLiteral{{.*}}0{{$}} + h.single_task( + []() [[intel::no_global_work_offset(0)]]{}); + + // CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}} + // expected-warning@+2{{'no_global_work_offset' attribute should be 0 or 1. Adjusted to 1}} + h.single_task( + []() [[intel::no_global_work_offset(42)]]{}); + + // expected-error@+2{{'no_global_work_offset' attribute requires a non-negative integral compile time constant expression}} + h.single_task( + []() [[intel::no_global_work_offset(-1)]]{}); + + // expected-error@+2{{'no_global_work_offset' attribute requires an integer constant}} + h.single_task( + []() [[intel::no_global_work_offset("foo")]]{}); + + h.single_task([]() { + // expected-error@+1{{'no_global_work_offset' attribute only applies to functions}} + [[intel::no_global_work_offset(1)]] int a; + + }); + // CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}} + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 1 + // expected-warning@+2{{attribute 'no_global_work_offset' is already applied}} + h.single_task( + []() [[intel::no_global_work_offset(0), intel::no_global_work_offset(1)]]{}); + }); return 0; } diff --git a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp index 4b2777c1bb9dd..942a708883de9 100644 --- a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp +++ b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp @@ -1,12 +1,15 @@ -// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -triple spir64 -Wno-sycl-2017-compat -verify -// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -triple spir64 -DTRIGGER_ERROR -Wno-sycl-2017-compat -verify -// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -triple spir64 -Wno-sycl-2017-compat | FileCheck %s +// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -Wno-sycl-2017-compat -verify +// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -DTRIGGER_ERROR -Wno-sycl-2017-compat -verify +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -Wno-sycl-2017-compat | FileCheck %s -#include "Inputs/sycl.hpp" +#include "sycl.hpp" + +using namespace cl::sycl; +queue q; #ifndef TRIGGER_ERROR //first case - good case -[[intel::no_global_work_offset]] // expected-no-diagnostics +[[intel::no_global_work_offset(1)]] // expected-no-diagnostics void func1(); @@ -46,23 +49,26 @@ func4() {} // expected-error {{'max_work_group_size' attribute conflicts with '' #endif int main() { + q.submit([&](handler &h) { #ifndef TRIGGER_ERROR // CHECK-LABEL: FunctionDecl {{.*}} main 'int ()' // CHECK: `-FunctionDecl {{.*}}test_kernel1 'void ()' // CHECK: -SYCLIntelMaxWorkGroupSizeAttr {{.*}} Inherited 4 4 4 - // CHECK: -SYCLIntelNoGlobalWorkOffsetAttr {{.*}} Inherited Enabled + // CHECK: -SYCLIntelNoGlobalWorkOffsetAttr {{.*}} // CHECK: `-ReqdWorkGroupSizeAttr {{.*}} 2 2 2 - cl::sycl::kernel_single_task( + h.single_task( []() { func1(); }); #else - cl::sycl::kernel_single_task( + h.single_task( []() { func2(); }); // expected-error {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}} - cl::sycl::kernel_single_task( + h.single_task( []() { func3(); }); - cl::sycl::kernel_single_task( + h.single_task( []() { func4(); }); #endif + }); + return 0; } diff --git a/clang/test/SemaSYCL/sycl-device-intel-fpga-no-global-work-offset.cpp b/clang/test/SemaSYCL/sycl-device-intel-fpga-no-global-work-offset.cpp new file mode 100644 index 0000000000000..0a05e94dc5fac --- /dev/null +++ b/clang/test/SemaSYCL/sycl-device-intel-fpga-no-global-work-offset.cpp @@ -0,0 +1,25 @@ +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s + +// Test that checkes template parameter support for 'no_global_work_offset' attribute on sycl device. + +template +class KernelFunctor { +public: + // expected-error@+1{{'no_global_work_offset' attribute requires a non-negative integral compile time constant expression}} + [[intel::no_global_work_offset(SIZE)]] void operator()() {} +}; + +int main() { + //expected-note@+1{{in instantiation of template class 'KernelFunctor<-1>' requested here}} + KernelFunctor<-1>(); + // no error expected + KernelFunctor<1>(); +} + +// CHECK: ClassTemplateDecl {{.*}} {{.*}} KernelFunctor +// CHECK: ClassTemplateSpecializationDecl {{.*}} {{.*}} class KernelFunctor definition +// CHECK: CXXRecordDecl {{.*}} {{.*}} implicit class KernelFunctor +// CHECK: SYCLIntelNoGlobalWorkOffsetAttr {{.*}} +// CHECK: SubstNonTypeTemplateParmExpr {{.*}} +// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} +// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} From b54e4012d75c99380443a7e9c11c55dd60c54f82 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 30 Nov 2020 22:55:30 -0800 Subject: [PATCH 02/13] Fix Clang format issues Signed-off-by: Soumi Manna --- clang/include/clang/Sema/Sema.h | 5 ++-- clang/lib/CodeGen/CodeGenFunction.cpp | 2 +- .../intel-fpga-no-global-work-offset.cpp | 8 +++--- .../redeclaration-attribute-propagation.cpp | 26 +++++++++---------- 4 files changed, 21 insertions(+), 20 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 51983d8276378..add7e1e945e31 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -12973,13 +12973,14 @@ void Sema::addIntelSYCLSingleArgFunctionAttr(Decl *D, } if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset) { if (ArgInt > 1) { - Diag(E->getExprLoc(), diag::warn_boolean_attribute_argument_is_not_valid) + Diag(E->getExprLoc(), + diag::warn_boolean_attribute_argument_is_not_valid) << CI.getAttrName(); return; } } if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelMaxGlobalWorkDim || - CI.getParsedKind() == ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset) { + CI.getParsedKind() == ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset) { if (ArgInt < 0) { Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) << CI.getAttrName() << /*non-negative*/ 1; diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index 3af461be136da..099dcad1e7822 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -701,7 +701,7 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, llvm::Metadata *AttrMDArgs[] = {llvm::ConstantAsMetadata::get( Builder.getInt32(ArgVal->getSExtValue()))}; Fn->setMetadata("no_global_work_offset", - llvm::MDNode::get(Context, AttrMDArgs)); + llvm::MDNode::get(Context, AttrMDArgs)); } if (FD->hasAttr()) { diff --git a/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp b/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp index 938f7814ff0e1..237a4bc50970a 100644 --- a/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp +++ b/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp @@ -36,15 +36,15 @@ int main() { []() [[intel::no_global_work_offset("foo")]]{}); h.single_task([]() { - // expected-error@+1{{'no_global_work_offset' attribute only applies to functions}} - [[intel::no_global_work_offset(1)]] int a; + // expected-error@+1{{'no_global_work_offset' attribute only applies to functions}} + [[intel::no_global_work_offset(1)]] int a; + }); - }); // CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}} // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 1 // expected-warning@+2{{attribute 'no_global_work_offset' is already applied}} h.single_task( []() [[intel::no_global_work_offset(0), intel::no_global_work_offset(1)]]{}); - }); + }); return 0; } diff --git a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp index 942a708883de9..c0a09931561f7 100644 --- a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp +++ b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp @@ -51,23 +51,23 @@ func4() {} // expected-error {{'max_work_group_size' attribute conflicts with '' int main() { q.submit([&](handler &h) { #ifndef TRIGGER_ERROR - // CHECK-LABEL: FunctionDecl {{.*}} main 'int ()' - // CHECK: `-FunctionDecl {{.*}}test_kernel1 'void ()' - // CHECK: -SYCLIntelMaxWorkGroupSizeAttr {{.*}} Inherited 4 4 4 - // CHECK: -SYCLIntelNoGlobalWorkOffsetAttr {{.*}} - // CHECK: `-ReqdWorkGroupSizeAttr {{.*}} 2 2 2 - h.single_task( - []() { func1(); }); + // CHECK-LABEL: FunctionDecl {{.*}} main 'int ()' + // CHECK: `-FunctionDecl {{.*}}test_kernel1 'void ()' + // CHECK: -SYCLIntelMaxWorkGroupSizeAttr {{.*}} Inherited 4 4 4 + // CHECK: -SYCLIntelNoGlobalWorkOffsetAttr {{.*}} + // CHECK: `-ReqdWorkGroupSizeAttr {{.*}} 2 2 2 + h.single_task( + []() { func1(); }); #else - h.single_task( - []() { func2(); }); // expected-error {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}} + h.single_task( + []() { func2(); }); // expected-error {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}} - h.single_task( - []() { func3(); }); + h.single_task( + []() { func3(); }); - h.single_task( - []() { func4(); }); + h.single_task( + []() { func4(); }); #endif }); return 0; From 7791b8ad91218bee66fea8b2e82345aa71d071ff Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 30 Nov 2020 23:13:37 -0800 Subject: [PATCH 03/13] Fix Clang format issues Signed-off-by: Soumi Manna --- clang/include/clang/Sema/Sema.h | 2 +- clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index add7e1e945e31..2827d1aceb316 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -12974,7 +12974,7 @@ void Sema::addIntelSYCLSingleArgFunctionAttr(Decl *D, if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset) { if (ArgInt > 1) { Diag(E->getExprLoc(), - diag::warn_boolean_attribute_argument_is_not_valid) + diag::warn_boolean_attribute_argument_is_not_valid) << CI.getAttrName(); return; } diff --git a/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp b/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp index 237a4bc50970a..19e2845ba424c 100644 --- a/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp +++ b/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp @@ -38,7 +38,7 @@ int main() { h.single_task([]() { // expected-error@+1{{'no_global_work_offset' attribute only applies to functions}} [[intel::no_global_work_offset(1)]] int a; - }); + }); // CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}} // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 1 From 7ceb0241bc5e4188aba2b6e73ac59544e1542fcc Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Wed, 9 Dec 2020 22:33:46 -0800 Subject: [PATCH 04/13] Add support for optional attribute parameter Signed-off-by: Soumi Manna --- clang/include/clang/Basic/Attr.td | 2 +- clang/lib/CodeGen/CodeGenFunction.cpp | 12 +++--------- clang/lib/Sema/SemaDeclAttr.cpp | 19 +++++++++++++------ .../intel-fpga-no-global-work-offset.cpp | 4 ++-- .../check-notdirect-attribute-propagation.cpp | 4 ++-- .../intel-fpga-no-global-work-offset.cpp | 2 +- .../redeclaration-attribute-propagation.cpp | 2 +- 7 files changed, 23 insertions(+), 22 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index fb171859a680f..c2866cc854d93 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1303,7 +1303,7 @@ def SYCLIntelMaxGlobalWorkDim : InheritableAttr { def SYCLIntelNoGlobalWorkOffset : InheritableAttr { let Spellings = [CXX11<"intelfpga","no_global_work_offset">, CXX11<"intel","no_global_work_offset">]; - let Args = [ExprArgument<"Value">, BoolArgument<"Enabled", 1>]; + let Args = [ExprArgument<"Value", /*default*/1>]; let LangOpts = [SYCLIsDevice, SYCLIsHost]; let Subjects = SubjectList<[Function], ErrorDiag>; let Documentation = [SYCLIntelNoGlobalWorkOffsetAttrDocs]; diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index 099dcad1e7822..1002e53449f23 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -693,15 +693,9 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, } if (const SYCLIntelNoGlobalWorkOffsetAttr *A = - FD->getAttr()) { - llvm::LLVMContext &Context = getLLVMContext(); - Optional ArgVal = - A->getValue()->getIntegerConstantExpr(FD->getASTContext()); - assert(ArgVal.hasValue() && "Not an integer constant expression"); - llvm::Metadata *AttrMDArgs[] = {llvm::ConstantAsMetadata::get( - Builder.getInt32(ArgVal->getSExtValue()))}; - Fn->setMetadata("no_global_work_offset", - llvm::MDNode::get(Context, AttrMDArgs)); + FD->getAttr()) { + if (A->getValue()) + Fn->setMetadata("no_global_work_offset", llvm::MDNode::get(Context, {})); } if (FD->hasAttr()) { diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index e4db90b8b8ec6..1b19aefdbfedc 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -5294,15 +5294,22 @@ static void handleNoGlobalWorkOffsetAttr(Sema &S, Decl *D, checkForDuplicateAttribute(S, D, Attr); - Expr *E = Attr.getArgAsExpr(0); - if (Attr.getKind() == ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset && checkDeprecatedSYCLAttributeSpelling(S, Attr)) S.Diag(Attr.getLoc(), diag::note_spelling_suggestion) - << "'intel::no_global_work_offset'"; - - S.addIntelSYCLSingleArgFunctionAttr(D, Attr, - E); + << "'intel::no_global_work_offset'"; + + // If no attribute argument is specified, set to default value '1'. + if (!Attr.isArgExpr(0)) { + Expr *E = IntegerLiteral::Create(S.Context, llvm::APInt(32, 1), + S.Context.IntTy, Attr.getLoc()); + D->addAttr(::new (S.Context) SYCLIntelNoGlobalWorkOffsetAttr(S.Context, + Attr, E)); + } else { + Expr *E = Attr.getArgAsExpr(0); + S.addIntelSYCLSingleArgFunctionAttr(D, Attr, + E); + } } /// Handle the [[intelfpga::doublepump]] and [[intelfpga::singlepump]] attributes. diff --git a/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp b/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp index 22d34fa88a51e..ece6a5f6f6d64 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp @@ -22,7 +22,7 @@ int main() { h.single_task(boo); h.single_task( - []() [[intel::no_global_work_offset(1)]]{}); + []() [[intel::no_global_work_offset]]{}); h.single_task( []() [[intel::no_global_work_offset(0)]]{}); @@ -38,4 +38,4 @@ int main() { // CHECK: define spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} ![[NUM4:[0-9]+]] // CHECK: define spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} !no_global_work_offset ![[NUM5]] // CHECK-NOT: ![[NUM4]] = !{i32 0} -// CHECK: ![[NUM5]] = !{i32 1} +// CHECK: ![[NUM5]] = !{} diff --git a/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp b/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp index fb4193ee95db2..3967f4322a937 100644 --- a/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp +++ b/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp @@ -3,9 +3,9 @@ // RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -triple spir64 | FileCheck %s #ifndef TRIGGER_ERROR -[[intel::no_global_work_offset(1)]] void not_direct_one() {} // expected-no-diagnostics +[[intel::no_global_work_offset]] void not_direct_one() {} // expected-no-diagnostics -[[intel::reqd_sub_group_size(1)]] void func_one() { +[[intel::reqd_sub_group_size(2)]] void func_one() { not_direct_one(); } diff --git a/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp b/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp index 19e2845ba424c..a6a4bffe401ef 100644 --- a/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp +++ b/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp @@ -8,7 +8,7 @@ queue q; struct FuncObj { //expected-warning@+2 {{attribute 'intelfpga::no_global_work_offset' is deprecated}} //expected-note@+1 {{did you mean to use 'intel::no_global_work_offset' instead?}} - [[intelfpga::no_global_work_offset(1)]] void operator()() const {} + [[intelfpga::no_global_work_offset]] void operator()() const {} }; int main() { diff --git a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp index c0a09931561f7..436f9b3186de2 100644 --- a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp +++ b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp @@ -9,7 +9,7 @@ queue q; #ifndef TRIGGER_ERROR //first case - good case -[[intel::no_global_work_offset(1)]] // expected-no-diagnostics +[[intel::no_global_work_offset]] // expected-no-diagnostics void func1(); From 77cb8e951a57bbd16eed0567902893774d5d71ee Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Wed, 9 Dec 2020 23:07:21 -0800 Subject: [PATCH 05/13] Fix calng format and address review comments Signed-off-by: Soumi Manna --- clang/lib/CodeGen/CodeGenFunction.cpp | 2 +- clang/lib/Sema/SemaDeclAttr.cpp | 10 +++++----- .../SemaSYCL/check-notdirect-attribute-propagation.cpp | 2 +- .../test/SemaSYCL/intel-fpga-no-global-work-offset.cpp | 2 ++ .../sycl-device-intel-fpga-no-global-work-offset.cpp | 2 +- 5 files changed, 10 insertions(+), 8 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index 1002e53449f23..7a99043f848d8 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -693,7 +693,7 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, } if (const SYCLIntelNoGlobalWorkOffsetAttr *A = - FD->getAttr()) { + FD->getAttr()) { if (A->getValue()) Fn->setMetadata("no_global_work_offset", llvm::MDNode::get(Context, {})); } diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 1b19aefdbfedc..d35872991b754 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -5297,18 +5297,18 @@ static void handleNoGlobalWorkOffsetAttr(Sema &S, Decl *D, if (Attr.getKind() == ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset && checkDeprecatedSYCLAttributeSpelling(S, Attr)) S.Diag(Attr.getLoc(), diag::note_spelling_suggestion) - << "'intel::no_global_work_offset'"; + << "'intel::no_global_work_offset'"; // If no attribute argument is specified, set to default value '1'. if (!Attr.isArgExpr(0)) { Expr *E = IntegerLiteral::Create(S.Context, llvm::APInt(32, 1), S.Context.IntTy, Attr.getLoc()); - D->addAttr(::new (S.Context) SYCLIntelNoGlobalWorkOffsetAttr(S.Context, - Attr, E)); + D->addAttr(::new (S.Context) + SYCLIntelNoGlobalWorkOffsetAttr(S.Context, Attr, E)); } else { Expr *E = Attr.getArgAsExpr(0); - S.addIntelSYCLSingleArgFunctionAttr(D, Attr, - E); + S.addIntelSYCLSingleArgFunctionAttr( + D, Attr, E); } } diff --git a/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp b/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp index 3967f4322a937..d2ade080f86cc 100644 --- a/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp +++ b/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp @@ -5,7 +5,7 @@ #ifndef TRIGGER_ERROR [[intel::no_global_work_offset]] void not_direct_one() {} // expected-no-diagnostics -[[intel::reqd_sub_group_size(2)]] void func_one() { +[[intel::reqd_sub_group_size(1)]] void func_one() { not_direct_one(); } diff --git a/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp b/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp index a6a4bffe401ef..8c30f15812cad 100644 --- a/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp +++ b/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp @@ -42,6 +42,8 @@ int main() { // CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}} // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 1 + // CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}} + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 0 // expected-warning@+2{{attribute 'no_global_work_offset' is already applied}} h.single_task( []() [[intel::no_global_work_offset(0), intel::no_global_work_offset(1)]]{}); diff --git a/clang/test/SemaSYCL/sycl-device-intel-fpga-no-global-work-offset.cpp b/clang/test/SemaSYCL/sycl-device-intel-fpga-no-global-work-offset.cpp index 0a05e94dc5fac..b1dad09d86585 100644 --- a/clang/test/SemaSYCL/sycl-device-intel-fpga-no-global-work-offset.cpp +++ b/clang/test/SemaSYCL/sycl-device-intel-fpga-no-global-work-offset.cpp @@ -1,6 +1,6 @@ // RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s -// Test that checkes template parameter support for 'no_global_work_offset' attribute on sycl device. +// Test that checks template parameter support for 'no_global_work_offset' attribute on sycl device. template class KernelFunctor { From e51dfff9c959a0f8bd0ee6a33330d78559dbd11f Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 10 Dec 2020 07:04:20 -0800 Subject: [PATCH 06/13] Address review comments Signed-off-by: Soumi Manna --- clang/include/clang/Basic/Attr.td | 2 +- clang/include/clang/Sema/Sema.h | 3 +-- clang/lib/Sema/SemaDeclAttr.cpp | 15 +++++---------- 3 files changed, 7 insertions(+), 13 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index c2866cc854d93..b5b97dfea0c0c 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1303,7 +1303,7 @@ def SYCLIntelMaxGlobalWorkDim : InheritableAttr { def SYCLIntelNoGlobalWorkOffset : InheritableAttr { let Spellings = [CXX11<"intelfpga","no_global_work_offset">, CXX11<"intel","no_global_work_offset">]; - let Args = [ExprArgument<"Value", /*default*/1>]; + let Args = [ExprArgument<"Value", /*optional*/1>]; let LangOpts = [SYCLIsDevice, SYCLIsHost]; let Subjects = SubjectList<[Function], ErrorDiag>; let Documentation = [SYCLIntelNoGlobalWorkOffsetAttrDocs]; diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 2827d1aceb316..70198b62206e4 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -12974,8 +12974,7 @@ void Sema::addIntelSYCLSingleArgFunctionAttr(Decl *D, if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset) { if (ArgInt > 1) { Diag(E->getExprLoc(), - diag::warn_boolean_attribute_argument_is_not_valid) - << CI.getAttrName(); + diag::warn_boolean_attribute_argument_is_not_valid) << CI; return; } } diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index d35872991b754..28213a6f7da54 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -5300,16 +5300,11 @@ static void handleNoGlobalWorkOffsetAttr(Sema &S, Decl *D, << "'intel::no_global_work_offset'"; // If no attribute argument is specified, set to default value '1'. - if (!Attr.isArgExpr(0)) { - Expr *E = IntegerLiteral::Create(S.Context, llvm::APInt(32, 1), - S.Context.IntTy, Attr.getLoc()); - D->addAttr(::new (S.Context) - SYCLIntelNoGlobalWorkOffsetAttr(S.Context, Attr, E)); - } else { - Expr *E = Attr.getArgAsExpr(0); - S.addIntelSYCLSingleArgFunctionAttr( - D, Attr, E); - } + Expr *E = Attr.isArgExpr(0) ? Attr.getArgAsExpr(0) + : IntegerLiteral::Create(S.Context, + llvm::APInt(32, 1), S.Context.IntTy, Attr.getLoc()); + S.addIntelSYCLSingleArgFunctionAttr(D, Attr, + E); } /// Handle the [[intelfpga::doublepump]] and [[intelfpga::singlepump]] attributes. From 19685c37b094f214deefb5ed95de1b2dfbd3dd84 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 10 Dec 2020 07:12:03 -0800 Subject: [PATCH 07/13] Fix Clang format issues Signed-off-by: Soumi Manna --- clang/include/clang/Sema/Sema.h | 5 +++-- clang/lib/Sema/SemaDeclAttr.cpp | 7 ++++--- 2 files changed, 7 insertions(+), 5 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 70198b62206e4..3e0504d087923 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -12973,8 +12973,9 @@ void Sema::addIntelSYCLSingleArgFunctionAttr(Decl *D, } if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset) { if (ArgInt > 1) { - Diag(E->getExprLoc(), - diag::warn_boolean_attribute_argument_is_not_valid) << CI; + Diag(E->getExprLoc(), + diag::warn_boolean_attribute_argument_is_not_valid) + << CI; return; } } diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 28213a6f7da54..cc219422c963b 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -5300,9 +5300,10 @@ static void handleNoGlobalWorkOffsetAttr(Sema &S, Decl *D, << "'intel::no_global_work_offset'"; // If no attribute argument is specified, set to default value '1'. - Expr *E = Attr.isArgExpr(0) ? Attr.getArgAsExpr(0) - : IntegerLiteral::Create(S.Context, - llvm::APInt(32, 1), S.Context.IntTy, Attr.getLoc()); + Expr *E = Attr.isArgExpr(0) + ? Attr.getArgAsExpr(0) + : IntegerLiteral::Create(S.Context, llvm::APInt(32, 1), + S.Context.IntTy, Attr.getLoc()); S.addIntelSYCLSingleArgFunctionAttr(D, Attr, E); } From bdd752a73dcc0b77a0cad045bd92e0174ffa1638 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 10 Dec 2020 07:16:15 -0800 Subject: [PATCH 08/13] Fix error Signed-off-by: Soumi Manna --- clang/include/clang/Sema/Sema.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 3e0504d087923..d06c2080eca99 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -12973,7 +12973,7 @@ void Sema::addIntelSYCLSingleArgFunctionAttr(Decl *D, } if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset) { if (ArgInt > 1) { - Diag(E->getExprLoc(), + Diag(E->getExprLoc(), diag::warn_boolean_attribute_argument_is_not_valid) << CI; return; From 7810812f8ff6334c0514f851e1ca4b3a18700101 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Tue, 15 Dec 2020 07:52:28 -0800 Subject: [PATCH 09/13] Remove diagnostic and update patch Signed-off-by: Soumi Manna --- clang/include/clang/Basic/DiagnosticGroups.td | 4 +--- clang/include/clang/Basic/DiagnosticSemaKinds.td | 3 --- clang/include/clang/Sema/Sema.h | 10 ++-------- .../test/SemaSYCL/intel-fpga-no-global-work-offset.cpp | 5 ----- 4 files changed, 3 insertions(+), 19 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticGroups.td b/clang/include/clang/Basic/DiagnosticGroups.td index e2894e764dba9..4b42a37fd33aa 100644 --- a/clang/include/clang/Basic/DiagnosticGroups.td +++ b/clang/include/clang/Basic/DiagnosticGroups.td @@ -667,10 +667,8 @@ def NSReturnsMismatch : DiagGroup<"nsreturns-mismatch">; def IndependentClassAttribute : DiagGroup<"IndependentClass-attribute">; def UnknownAttributes : DiagGroup<"unknown-attributes">; def IgnoredAttributes : DiagGroup<"ignored-attributes">; -def AdjustedAttributes : DiagGroup<"adjusted-attributes">; def Attributes : DiagGroup<"attributes", [UnknownAttributes, - IgnoredAttributes, - AdjustedAttributes]>; + IgnoredAttributes]>; def UnknownSanitizers : DiagGroup<"unknown-sanitizers">; def UnnamedTypeTemplateArgs : DiagGroup<"unnamed-type-template-args", [CXX98CompatUnnamedTypeTemplateArgs]>; diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index a0eaf8d7bdb3c..f4511886eb212 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -11134,9 +11134,6 @@ def err_sycl_function_attribute_mismatch : Error< "SYCL kernel without %0 attribute can't call a function with this attribute">; def err_sycl_x_y_z_arguments_must_be_one : Error< "%0 X-, Y- and Z- sizes must be 1 when %1 attribute is used with value 0">; -def warn_boolean_attribute_argument_is_not_valid: Warning< - "The value of %0 attribute should be 0 or 1. Adjusted to 1">, - InGroup; def err_sycl_attibute_cannot_be_applied_here : Error<"%0 attribute cannot be applied to a " "static function or function in an anonymous namespace">; diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index d06c2080eca99..d6303fdb14feb 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -12971,14 +12971,6 @@ void Sema::addIntelSYCLSingleArgFunctionAttr(Decl *D, return; } } - if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset) { - if (ArgInt > 1) { - Diag(E->getExprLoc(), - diag::warn_boolean_attribute_argument_is_not_valid) - << CI; - return; - } - } if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelMaxGlobalWorkDim || CI.getParsedKind() == ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset) { if (ArgInt < 0) { @@ -12986,6 +12978,8 @@ void Sema::addIntelSYCLSingleArgFunctionAttr(Decl *D, << CI.getAttrName() << /*non-negative*/ 1; return; } + } + if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelMaxGlobalWorkDim) { if (ArgInt > 3) { Diag(E->getBeginLoc(), diag::err_attribute_argument_out_of_range) << CI.getAttrName() << 0 << 3 << E->getSourceRange(); diff --git a/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp b/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp index 8c30f15812cad..fdad6e5f4a1cb 100644 --- a/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp +++ b/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp @@ -23,7 +23,6 @@ int main() { []() [[intel::no_global_work_offset(0)]]{}); // CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}} - // expected-warning@+2{{'no_global_work_offset' attribute should be 0 or 1. Adjusted to 1}} h.single_task( []() [[intel::no_global_work_offset(42)]]{}); @@ -40,10 +39,6 @@ int main() { [[intel::no_global_work_offset(1)]] int a; }); - // CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}} - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 1 - // CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}} - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 0 // expected-warning@+2{{attribute 'no_global_work_offset' is already applied}} h.single_task( []() [[intel::no_global_work_offset(0), intel::no_global_work_offset(1)]]{}); From 2a0df8909a57a4bea269f36d0b71d45e8e7d9ae8 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Tue, 15 Dec 2020 19:38:41 -0800 Subject: [PATCH 10/13] Address review comments and add new test cases Signed-off-by: Soumi Manna --- clang/lib/CodeGen/CodeGenFunction.cpp | 9 +++-- .../intel-fpga-no-global-work-offset.cpp | 8 +++++ ...evice-intel-fpga-no-global-work-offset.cpp | 36 +++++++++++++++++++ 3 files changed, 51 insertions(+), 2 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index 7a99043f848d8..364acd535181c 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -693,8 +693,13 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, } if (const SYCLIntelNoGlobalWorkOffsetAttr *A = - FD->getAttr()) { - if (A->getValue()) + FD->getAttr()) { + const Expr *Arg = A->getValue(); + assert(Arg && "Got an unexpected null argument"); + Optional ArgVal = + Arg->getIntegerConstantExpr(FD->getASTContext()); + assert(ArgVal.hasValue() && "Not an integer constant expression"); + if (ArgVal) Fn->setMetadata("no_global_work_offset", llvm::MDNode::get(Context, {})); } diff --git a/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp b/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp index ece6a5f6f6d64..a2d33c2ac2932 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp @@ -16,6 +16,9 @@ class Functor { [[intel::no_global_work_offset(SIZE)]] void operator()() const {} }; +template +[[intel::no_global_work_offset(N)]] void func() {} + int main() { q.submit([&](handler &h) { Foo boo; @@ -29,6 +32,10 @@ int main() { Functor<1> f; h.single_task(f); + + h.single_task([]() { + func<1>(); + }); }); return 0; } @@ -37,5 +44,6 @@ int main() { // CHECK: define spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !no_global_work_offset ![[NUM5]] // CHECK: define spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} ![[NUM4:[0-9]+]] // CHECK: define spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} !no_global_work_offset ![[NUM5]] +// CHECK: define spir_kernel void @{{.*}}kernel_name5"() #0 {{.*}} !no_global_work_offset ![[NUM5]] // CHECK-NOT: ![[NUM4]] = !{i32 0} // CHECK: ![[NUM5]] = !{} diff --git a/clang/test/SemaSYCL/sycl-device-intel-fpga-no-global-work-offset.cpp b/clang/test/SemaSYCL/sycl-device-intel-fpga-no-global-work-offset.cpp index b1dad09d86585..76786c8a09bf6 100644 --- a/clang/test/SemaSYCL/sycl-device-intel-fpga-no-global-work-offset.cpp +++ b/clang/test/SemaSYCL/sycl-device-intel-fpga-no-global-work-offset.cpp @@ -2,6 +2,25 @@ // Test that checks template parameter support for 'no_global_work_offset' attribute on sycl device. +// Test that checks wrong template instantiation. +template +[[intel::no_global_work_offset(Ty{})]] void func() {} + +struct S {}; + // expected-error@+2{{template specialization requires 'template<>'}} + // expected-error@+1{{C++ requires a type specifier for all declarations}} + func(); + +// Test that checks expression is not a constant expression. +int foo(); +// expected-error@+1{{'no_global_work_offset' attribute requires an integer constant}} +[[intel::no_global_work_offset(foo() + 12)]] void func1(); + +// Test that checks expression is a constant expression. +constexpr int bar() { return 0; } +[[intel::no_global_work_offset(bar() + 12)]] void func2(); // OK + +// Test that checks template parameter suppport on member function of class template. template class KernelFunctor { public: @@ -23,3 +42,20 @@ int main() { // CHECK: SubstNonTypeTemplateParmExpr {{.*}} // CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + +// Test that checks template parameter suppport on function. +template +[[intel::no_global_work_offset(N)]] void func3() {} + +int check() { + func3<1>(); + return 0; +} + +// CHECK: FunctionTemplateDecl {{.*}} {{.*}} func3 +// CHECK: NonTypeTemplateParmDecl {{.*}} {{.*}} referenced 'int' depth 0 index 0 N +// CHECK: FunctionDecl {{.*}} {{.*}} func3 'void ()' +// CHECK: SYCLIntelNoGlobalWorkOffsetAttr {{.*}} +// CHECK: SubstNonTypeTemplateParmExpr {{.*}} +// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} +// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} From fca61c5aca813443c3cdbd07044bea2e55185d17 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Tue, 15 Dec 2020 19:46:00 -0800 Subject: [PATCH 11/13] Fix Clang format errors Signed-off-by: Soumi Manna --- clang/lib/CodeGen/CodeGenFunction.cpp | 4 ++-- .../sycl-device-intel-fpga-no-global-work-offset.cpp | 6 +++--- 2 files changed, 5 insertions(+), 5 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index 364acd535181c..a59583bf80e8b 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -693,11 +693,11 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, } if (const SYCLIntelNoGlobalWorkOffsetAttr *A = - FD->getAttr()) { + FD->getAttr()) { const Expr *Arg = A->getValue(); assert(Arg && "Got an unexpected null argument"); Optional ArgVal = - Arg->getIntegerConstantExpr(FD->getASTContext()); + Arg->getIntegerConstantExpr(FD->getASTContext()); assert(ArgVal.hasValue() && "Not an integer constant expression"); if (ArgVal) Fn->setMetadata("no_global_work_offset", llvm::MDNode::get(Context, {})); diff --git a/clang/test/SemaSYCL/sycl-device-intel-fpga-no-global-work-offset.cpp b/clang/test/SemaSYCL/sycl-device-intel-fpga-no-global-work-offset.cpp index 76786c8a09bf6..3391d20b0dd85 100644 --- a/clang/test/SemaSYCL/sycl-device-intel-fpga-no-global-work-offset.cpp +++ b/clang/test/SemaSYCL/sycl-device-intel-fpga-no-global-work-offset.cpp @@ -7,9 +7,9 @@ template [[intel::no_global_work_offset(Ty{})]] void func() {} struct S {}; - // expected-error@+2{{template specialization requires 'template<>'}} - // expected-error@+1{{C++ requires a type specifier for all declarations}} - func(); +// expected-error@+2{{template specialization requires 'template<>'}} +// expected-error@+1{{C++ requires a type specifier for all declarations}} +func(); // Test that checks expression is not a constant expression. int foo(); From 87cb866400ac63cb43dfaf3368b5eeb963b80a86 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Wed, 16 Dec 2020 07:56:17 -0800 Subject: [PATCH 12/13] address review comments Signed-off-by: Soumi Manna --- clang/include/clang/Sema/Sema.h | 5 +---- clang/lib/CodeGen/CodeGenFunction.cpp | 2 +- .../SemaSYCL/intel-fpga-no-global-work-offset.cpp | 2 +- ...cl-device-intel-fpga-no-global-work-offset.cpp | 15 +++++++-------- 4 files changed, 10 insertions(+), 14 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index d6303fdb14feb..cc49a9994b21b 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -12971,15 +12971,12 @@ void Sema::addIntelSYCLSingleArgFunctionAttr(Decl *D, return; } } - if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelMaxGlobalWorkDim || - CI.getParsedKind() == ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset) { + if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelMaxGlobalWorkDim) { if (ArgInt < 0) { Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) << CI.getAttrName() << /*non-negative*/ 1; return; } - } - if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelMaxGlobalWorkDim) { if (ArgInt > 3) { Diag(E->getBeginLoc(), diag::err_attribute_argument_out_of_range) << CI.getAttrName() << 0 << 3 << E->getSourceRange(); diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index a59583bf80e8b..5c947cc66b42d 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -699,7 +699,7 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, Optional ArgVal = Arg->getIntegerConstantExpr(FD->getASTContext()); assert(ArgVal.hasValue() && "Not an integer constant expression"); - if (ArgVal) + if (ArgVal->getBoolValue()) Fn->setMetadata("no_global_work_offset", llvm::MDNode::get(Context, {})); } diff --git a/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp b/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp index fdad6e5f4a1cb..a2134ecaf644b 100644 --- a/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp +++ b/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp @@ -26,7 +26,7 @@ int main() { h.single_task( []() [[intel::no_global_work_offset(42)]]{}); - // expected-error@+2{{'no_global_work_offset' attribute requires a non-negative integral compile time constant expression}} + // CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}} h.single_task( []() [[intel::no_global_work_offset(-1)]]{}); diff --git a/clang/test/SemaSYCL/sycl-device-intel-fpga-no-global-work-offset.cpp b/clang/test/SemaSYCL/sycl-device-intel-fpga-no-global-work-offset.cpp index 3391d20b0dd85..c6e2bb0475d28 100644 --- a/clang/test/SemaSYCL/sycl-device-intel-fpga-no-global-work-offset.cpp +++ b/clang/test/SemaSYCL/sycl-device-intel-fpga-no-global-work-offset.cpp @@ -2,14 +2,17 @@ // Test that checks template parameter support for 'no_global_work_offset' attribute on sycl device. -// Test that checks wrong template instantiation. +// Test that checks wrong function template instantiation and ensures that the type +// is checked properly when instantiating from the template definition. template +// expected-error@+1{{'no_global_work_offset' attribute requires an integer constant}} [[intel::no_global_work_offset(Ty{})]] void func() {} struct S {}; -// expected-error@+2{{template specialization requires 'template<>'}} -// expected-error@+1{{C++ requires a type specifier for all declarations}} -func(); +void var() { + //expected-note@+1{{in instantiation of function template specialization 'func' requested here}} + func(); +} // Test that checks expression is not a constant expression. int foo(); @@ -24,14 +27,10 @@ constexpr int bar() { return 0; } template class KernelFunctor { public: - // expected-error@+1{{'no_global_work_offset' attribute requires a non-negative integral compile time constant expression}} [[intel::no_global_work_offset(SIZE)]] void operator()() {} }; int main() { - //expected-note@+1{{in instantiation of template class 'KernelFunctor<-1>' requested here}} - KernelFunctor<-1>(); - // no error expected KernelFunctor<1>(); } From 6471266853d7cf9c520c3cfed02452a431a05235 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Wed, 16 Dec 2020 09:43:37 -0800 Subject: [PATCH 13/13] Add CHECKs for Integer literals Signed-off-by: Soumi Manna --- clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp b/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp index a2134ecaf644b..c8c54a6913587 100644 --- a/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp +++ b/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp @@ -23,10 +23,13 @@ int main() { []() [[intel::no_global_work_offset(0)]]{}); // CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}} + // CHECK-NEXT: IntegerLiteral{{.*}}42{{$}} h.single_task( []() [[intel::no_global_work_offset(42)]]{}); // CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}} + // CHECK-NEXT: UnaryOperator{{.*}} 'int' prefix '-' + // CHECK-NEXT-NEXT: IntegerLiteral{{.*}}1{{$}} h.single_task( []() [[intel::no_global_work_offset(-1)]]{});