[FPGA][SYCL] Add support for memory attributes on device_global variables

clang/include/clang/Basic/Attr.td

@@ -2730,9 +2730,6 @@ def : MutualExclusions<[SYCLIntelRegister, SYCLIntelBankWidth]>;
 def SYCLIntelNumBanks : InheritableAttr {
  let Spellings = [CXX11<"intel", "numbanks">];
  let Args = [ExprArgument<"Value">];
- let Subjects = SubjectList<[SYCLIntelConstVar,
- SYCLIntelLocalStaticAgentMemVar,
- Field], ErrorDiag>;
  let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
  let Documentation = [SYCLIntelNumBanksAttrDocs];
 }

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -12215,6 +12215,10 @@ def err_sycl_attribute_internal_decl
  "in an anonymous namespace">;
 def err_sycl_attribute_not_device_global
  : Error<"%0 attribute can only be applied to 'device_global' variables">;
+def err_fpga_attribute_incorrrect_variable
+ : Error<"%0 attribute only applies to constant variables, local variables, "
+ "static variables, agent memory arguments, non-static data "
+ "members, and non-constant device_global variables">;
 def err_sycl_compiletime_property_duplication : Error<
  "can't apply %0 property twice to the same accessor">;
 def err_sycl_invalid_property_list_param_number : Error<

clang/lib/Sema/SemaDeclAttr.cpp

@@ -7745,6 +7745,16 @@ void Sema::AddSYCLIntelNumBanksAttr(Decl *D, const AttributeCommonInfo &CI,
  }
  }
 
+ if (auto *VD = dyn_cast<VarDecl>(D)) {
+ if (!(VD->getKind() != Decl::ImplicitParam &&
+ VD->getKind() != Decl::NonTypeTemplateParm &&
+ ((VD->getStorageClass() == SC_Static || VD->hasLocalStorage()) ||
+ (isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(VD->getType())
+ || VD->getType().isConstQualified())))){
+ Diag(CI.getLoc(), diag::err_fpga_attribute_incorrrect_variable) << CI;
+ }
+ }
+
  // Check to see if there's a duplicate attribute with different values
  // already applied to the declaration.
  if (const auto *DeclAttr = D->getAttr<SYCLIntelNumBanksAttr>()) {

clang/test/CodeGenSYCL/device_global.cpp

@@ -11,6 +11,9 @@ using namespace sycl;
 queue q;
 
 device_global<int> A;
+
+[[intel::numbanks(2)]] /*const*/ device_global<int> A1;
+
 #ifdef SYCL_EXTERNAL
 SYCL_EXTERNAL device_global<int> AExt;
 #endif
@@ -20,8 +23,10 @@ struct Foo {
  static device_global<int> C;
 };
 device_global<int> Foo::C;
+
 // CHECK-RDC: @AExt = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[AEXT_ATTRS:[0-9]+]]
 // CHECK: @A = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[A_ATTRS:[0-9]+]]
+// CHECK: @A1 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[A1_ATTRS:[0-9]+]]
 // CHECK: @_ZL1B = internal addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[B_ATTRS:[0-9]+]]
 // CHECK: @_ZN3Foo1CE = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[C_ATTRS:[0-9]+]]
 
@@ -70,6 +75,7 @@ void foo() {
  q.submit([&](handler &h) {
  h.single_task<class kernel_name_1>([=]() {
  (void)A;
+ (void)A1;
  (void)B;
  (void)Foo::C;
  (void)same_name;
@@ -103,6 +109,7 @@ void bar() {
 
 // CHECK-RDC: attributes #[[AEXT_ATTRS]] = { "sycl-unique-id"="_Z4AExt" }
 // CHECK: attributes #[[A_ATTRS]] = { "sycl-unique-id"="_Z1A" }
+// CHECK: attributes #[[A1_ATTRS]] = { "sycl-unique-id"="_Z2A1" }
 // CHECK: attributes #[[B_ATTRS]] = { "sycl-unique-id"="THE_PREFIX____ZL1B" }
 // CHECK: attributes #[[C_ATTRS]] = { "sycl-unique-id"="_ZN3Foo1CE" }
 // CHECK: attributes #[[SAME_NAME_ATTRS]] = { "sycl-unique-id"="_Z9same_name" }

clang/test/SemaSYCL/intel-fpga-device-global.cpp

@@ -0,0 +1,44 @@
+// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -sycl-std=2020 -verify %s
+#include "Inputs/sycl.hpp"
+
+// Tests that [[intel::numbanks()]] can be applied to device_global variables.
+using namespace sycl::ext::oneapi;
+
+[[intel::numbanks(2)]] device_global<int> dev_glob; // OK
+[[intel::numbanks(4)]] static device_global<float> static_dev_glob; // OK
+
+// expected-error@+1{{'numbanks' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members, and non-constant device_global variables}}
+[[intel::numbanks(2)]] int K;
+
+struct bar {
+ [[intel::numbanks(2)]] /*const*/ device_global<int> const_glob3; // OK
+ [[intel::numbanks(2)]] const device_global<int> const_glob4; // OK
+};
+
+void foo() {
+ [[intel::numbanks(2)]] int A1; // OK
+}
+
+struct [[__sycl_detail__::global_variable_allowed]] GlobAllowedVarOnly {
+};
+
+// expected-error@+1{{'numbanks' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members, and non-constant device_global variables}}
+[[intel::numbanks(2)]] GlobAllowedVarOnly GAVO;
+
+[[intel::numbanks(4)]] /*const*/ device_global<int> Good;
+[[intel::numbanks(4)]] extern device_global<int> Bad;
+
+int main() {
+ sycl::kernel_single_task<class KernelName1>([=]() {
+ Good.get();
+ // expected-error@+1 {{invalid reference to 'device_global' variable; external 'device_global' variable must be marked with SYCL_EXTERNAL macro}}
+ Bad.get();
+ (void)GAVO;
+ });
+ return 0;
+}
+
+//expected-error@+1{{'numbanks' attribute only applies to constant variables, local variables, static variables, agent memory arguments, non-static data members, and non-constant device_global variables}}
+[[intel::numbanks(2)]]
+__attribute__((opencl_global)) unsigned int ocl_glob_num_p2d[64] = {1, 2, 3};
+