[SYCL] Add template parameter support for no_global_work_offset attribute

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", /*default*/1>];
   let LangOpts = [SYCLIsDevice, SYCLIsHost];
   let Subjects = SubjectList<[Function], ErrorDiag>;
   let Documentation = [SYCLIntelNoGlobalWorkOffsetAttrDocs];

clang/include/clang/Sema/Sema.h

@@ -12971,7 +12971,16 @@ 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;

clang/lib/CodeGen/CodeGenFunction.cpp

@@ -693,8 +693,8 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD,
   }
 
   if (const SYCLIntelNoGlobalWorkOffsetAttr *A =
-          FD->getAttr<SYCLIntelNoGlobalWorkOffsetAttr>()) {
-    if (A->getEnabled())
+      FD->getAttr<SYCLIntelNoGlobalWorkOffsetAttr>()) {
+    if (A->getValue())
       Fn->setMetadata("no_global_work_offset", llvm::MDNode::get(Context, {}));
   }
 

clang/lib/Sema/SemaDeclAttr.cpp

@@ -5294,24 +5294,22 @@ static void handleNoGlobalWorkOffsetAttr(Sema &S, Decl *D,
 
   checkForDuplicateAttribute<SYCLIntelNoGlobalWorkOffsetAttr>(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;
-
   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));
+	<< "'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<SYCLIntelNoGlobalWorkOffsetAttr>(D, Attr,
+		                                                         E);
+  }
 }
 
 /// Handle the [[intelfpga::doublepump]] and [[intelfpga::singlepump]] attributes.

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<SYCLIntelNoGlobalWorkOffsetAttr>(TmplAttr)) {
+      instantiateIntelSYCLFunctionAttr<SYCLIntelNoGlobalWorkOffsetAttr>(
+          *this, TemplateArgs, SYCLIntelNoGlobalWorkOffset, New);
+      continue;
+    }
     // Existing DLL attribute on the instantiation takes precedence.
     if (TmplAttr->getKind() == attr::DLLExport ||
         TmplAttr->getKind() == attr::DLLImport) {

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 <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
-  kernelFunc();
-}
+template <int SIZE>
+class Functor {
+public:
+  [[intel::no_global_work_offset(SIZE)]] void operator()() const {}
+};
+
+int main() {
+  q.submit([&](handler &h) {
+    Foo boo;
+    h.single_task<class kernel_name1>(boo);
 
-void bar() {
-  Foo boo;
-  kernel<class kernel_name1>(boo);
+    h.single_task<class kernel_name2>(
+        []() [[intel::no_global_work_offset]]{});
 
-  kernel<class kernel_name2>(
-      []() [[intel::no_global_work_offset]]{});
+    h.single_task<class kernel_name3>(
+        []() [[intel::no_global_work_offset(0)]]{});
 
-  kernel<class kernel_name3>(
-      []() [[intel::no_global_work_offset(0)]]{});
+    Functor<1> f;
+    h.single_task<class kernel_name4>(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]] = !{}

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(1)]] void func_one() {
+[[intel::reqd_sub_group_size(2)]] 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<class KernelName>([]() {});
 #else
   parallel_for<class KernelName>([]() {}); // expected-error 2 {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}}

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]] void operator()() const {}
 };
 
-template <typename name, typename Func>
-void kernel(Func kernelFunc) {
-  kernelFunc();
-}
-
 int main() {
-  // CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}}Enabled
-  kernel<class test_kernel1>([]() {
-    FuncObj();
-  });
-
-  // CHECK: SYCLIntelNoGlobalWorkOffsetAttr
-  // CHECK-NOT: Enabled
-  kernel<class test_kernel2>(
-      []() [[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<class test_kernel3>(
-      []() [[intel::no_global_work_offset(42)]]{});
-
-  // expected-error@+2{{'no_global_work_offset' attribute requires a non-negative integral compile time constant expression}}
-  kernel<class test_kernel4>(
-      []() [[intel::no_global_work_offset(-1)]]{});
-
-  // expected-error@+2{{'no_global_work_offset' attribute requires parameter 0 to be an integer constant}}
-  kernel<class test_kernel5>(
-      []() [[intel::no_global_work_offset("foo")]]{});
-
-  kernel<class test_kernel6>([]() {
-    // expected-error@+1{{'no_global_work_offset' attribute only applies to functions}}
-    [[intel::no_global_work_offset(1)]] int a;
+  q.submit([&](handler &h) {
+    // CHECK:       SYCLIntelNoGlobalWorkOffsetAttr {{.*}}
+    // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
+    h.single_task<class test_kernel1>(FuncObj());
+
+    // CHECK: SYCLIntelNoGlobalWorkOffsetAttr {{.*}}
+    // CHECK-NEXT:  IntegerLiteral{{.*}}0{{$}}
+    h.single_task<class test_kernel2>(
+        []() [[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<class test_kernel3>(
+        []() [[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<class test_kernel4>(
+        []() [[intel::no_global_work_offset(-1)]]{});
+
+    // expected-error@+2{{'no_global_work_offset' attribute requires an integer constant}}
+    h.single_task<class test_kernel5>(
+        []() [[intel::no_global_work_offset("foo")]]{});
+
+    h.single_task<class test_kernel6>([]() {
+      // 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<class test_kernel7>(
+        []() [[intel::no_global_work_offset(0), intel::no_global_work_offset(1)]]{});
   });
-
-  // CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}}
-  // CHECK-NOT: Enabled
-  // CHECK: SYCLIntelNoGlobalWorkOffsetAttr{{.*}}Enabled
-  // expected-warning@+2{{attribute 'no_global_work_offset' is already applied}}
-  kernel<class test_kernel7>(
-      []() [[intel::no_global_work_offset(0), intel::no_global_work_offset(1)]]{});
-
   return 0;
 }

clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp

@@ -1,8 +1,11 @@
-// 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
@@ -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:  `-ReqdWorkGroupSizeAttr {{.*}} 2 2 2
-  cl::sycl::kernel_single_task<class test_kernel1>(
-      []() { 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<class test_kernel1>(
+        []() { func1(); });
 
 #else
-  cl::sycl::kernel_single_task<class test_kernel2>(
-      []() { func2(); }); // expected-error {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}}
+    h.single_task<class test_kernel2>(
+        []() { func2(); }); // expected-error {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}}
 
-  cl::sycl::kernel_single_task<class test_kernel3>(
-      []() { func3(); });
+    h.single_task<class test_kernel3>(
+        []() { func3(); });
 
-  cl::sycl::kernel_single_task<class test_kernel4>(
-      []() { func4(); });
+    h.single_task<class test_kernel4>(
+        []() { func4(); });
 #endif
+  });
+  return 0;
 }

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 <int SIZE>
+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{{$}}