[SYCL][FPGA] Allow memory attributes on device_global variables for d…

…evice compilation (#12930) #12785 added support for non-const device_global variables on all FPGA memory attributes. FPGA memory attributes should only work for device codes. They should have no effect on host code. This patch fixes bug where we emitted error when attributes are applied to device_globals variables for host compilation. --------- Signed-off-by: Soumi Manna <soumi.manna@intel.com>
intel · Mar 6, 2024 · 3bb5f40 · 3bb5f40
1 parent 984c88c
commit 3bb5f40
Show file tree

Hide file tree

Showing 3 changed files with 232 additions and 36 deletions.
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -7540,13 +7540,15 @@ static void handleSYCLIntelSinglePumpAttr(Sema &S, Decl *D,
   }
 
   // Check attribute applies to field, constant variables, local variables,
-  // static variables, non-static data members, and device_global variables.
-  if ((D->getKind() == Decl::ParmVar) ||
-      CheckValidFPGAMemoryAttributesVar(S, D)) {
+  // static variables, non-static data members, and device_global variables
+  // for the device compilation.
+  if (S.Context.getLangOpts().SYCLIsDevice &&
+      ((D->getKind() == Decl::ParmVar) ||
+       CheckValidFPGAMemoryAttributesVar(S, D))) {
     S.Diag(AL.getLoc(), diag::err_fpga_attribute_incorrect_variable)
         << AL << /*agent memory arguments*/ 0;
     return;
-   }
+  }
 
   // If the declaration does not have an [[intel::fpga_memory]]
   // attribute, this creates one as an implicit attribute.
@@ -7572,13 +7574,15 @@ static void handleSYCLIntelDoublePumpAttr(Sema &S, Decl *D,
   }
 
   // Check attribute applies to field, constant variables, local variables,
-  // static variables, non-static data members, and device_global variables.
-  if ((D->getKind() == Decl::ParmVar) ||
-      CheckValidFPGAMemoryAttributesVar(S, D)) {
+  // static variables, non-static data members, and device_global variables
+  // for the device compilation.
+  if (S.Context.getLangOpts().SYCLIsDevice &&
+      ((D->getKind() == Decl::ParmVar) ||
+       CheckValidFPGAMemoryAttributesVar(S, D))) {
     S.Diag(AL.getLoc(), diag::err_fpga_attribute_incorrect_variable)
         << AL << /*agent memory arguments*/ 0;
     return;
-   }
+  }
 
   // If the declaration does not have an [[intel::fpga_memory]]
   // attribute, this creates one as an implicit attribute.
@@ -7629,12 +7633,13 @@ static void handleSYCLIntelMemoryAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
 
   // Check attribute applies to field, constant variables, local variables,
   // static variables, agent memory arguments, non-static data members,
-  // and device_global variables.
-  if (CheckValidFPGAMemoryAttributesVar(S, D)) {
+  // and device_global variables for the device compilation.
+  if (S.Context.getLangOpts().SYCLIsDevice &&
+      CheckValidFPGAMemoryAttributesVar(S, D)) {
     S.Diag(AL.getLoc(), diag::err_fpga_attribute_incorrect_variable)
         << AL << /*agent memory arguments*/ 1;
     return;
-   }
+  }
 
   D->addAttr(::new (S.Context) SYCLIntelMemoryAttr(S.Context, AL, Kind));
 }
@@ -7669,13 +7674,15 @@ static void handleSYCLIntelRegisterAttr(Sema &S, Decl *D,
   }
 
   // Check attribute applies to field, constant variables, local variables,
-  // static variables, non-static data members, and device_global variables.
-  if ((D->getKind() == Decl::ParmVar) ||
-      CheckValidFPGAMemoryAttributesVar(S, D)) {
+  // static variables, non-static data members, and device_global variables
+  // for the device compilation.
+  if (S.Context.getLangOpts().SYCLIsDevice &&
+      ((D->getKind() == Decl::ParmVar) ||
+       CheckValidFPGAMemoryAttributesVar(S, D))) {
     S.Diag(A.getLoc(), diag::err_fpga_attribute_incorrect_variable)
         << A << /*agent memory arguments*/ 0;
     return;
-   }
+  }
 
   if (checkIntelFPGARegisterAttrCompatibility(S, D, A))
     return;
@@ -7717,8 +7724,9 @@ void Sema::AddSYCLIntelBankWidthAttr(Decl *D, const AttributeCommonInfo &CI,
 
     // Check attribute applies to field, constant variables, local variables,
     // static variables, agent memory arguments, non-static data members,
-    // and device_global variables.
-    if (CheckValidFPGAMemoryAttributesVar(*this, D)) {
+    // and device_global variables for the device compilation.
+    if (Context.getLangOpts().SYCLIsDevice &&
+        CheckValidFPGAMemoryAttributesVar(*this, D)) {
       Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable)
           << CI << /*agent memory arguments*/ 1;
       return;
@@ -7810,8 +7818,9 @@ void Sema::AddSYCLIntelNumBanksAttr(Decl *D, const AttributeCommonInfo &CI,
 
     // Check attribute applies to constant variables, local variables,
     // static variables, agent memory arguments, non-static data members,
-    // and device_global variables.
-    if (CheckValidFPGAMemoryAttributesVar(*this, D)) {
+    // and device_global variables for the device compilation.
+    if (Context.getLangOpts().SYCLIsDevice &&
+        CheckValidFPGAMemoryAttributesVar(*this, D)) {
       Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable)
           << CI << /*agent memory arguments*/ 1;
       return;
@@ -7886,11 +7895,12 @@ static void handleIntelSimpleDualPortAttr(Sema &S, Decl *D,
 
   // Check attribute applies to field, constant variables, local variables,
   // static variables, agent memory arguments, non-static data members,
-  // and device_global variables.
-  if (CheckValidFPGAMemoryAttributesVar(S, D)) {
-    S.Diag(AL.getLoc(), diag::err_fpga_attribute_incorrect_variable)
-        << AL << /*agent memory arguments*/ 1;
-    return;
+  // and device_global variables for the device compilation.
+  if (S.Context.getLangOpts().SYCLIsDevice &&
+      CheckValidFPGAMemoryAttributesVar(S, D)) {
+      S.Diag(AL.getLoc(), diag::err_fpga_attribute_incorrect_variable)
+          << AL << /*agent memory arguments*/ 1;
+      return;
   }
 
   if (!D->hasAttr<SYCLIntelMemoryAttr>())
@@ -7921,8 +7931,9 @@ void Sema::AddSYCLIntelMaxReplicatesAttr(Decl *D, const AttributeCommonInfo &CI,
 
     // Check attribute applies to field, constant variables, local variables,
     // static variables, agent memory arguments, non-static data members,
-    // and device_global variables.
-    if (CheckValidFPGAMemoryAttributesVar(*this, D)) {
+    // and device_global variables for the device compilation.
+    if (Context.getLangOpts().SYCLIsDevice &&
+        CheckValidFPGAMemoryAttributesVar(*this, D)) {
       Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable)
           << CI << /*agent memory arguments*/ 1;
       return;
@@ -8012,13 +8023,15 @@ static void handleSYCLIntelMergeAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
   }
 
   // Check attribute applies to field, constant variables, local variables,
-  // static variables, non-static data members, and device_global variables.
-  if ((D->getKind() == Decl::ParmVar) ||
-      CheckValidFPGAMemoryAttributesVar(S, D)) {
+  // static variables, non-static data members, and device_global variables
+  // for the device compilation.
+  if (S.Context.getLangOpts().SYCLIsDevice &&
+      ((D->getKind() == Decl::ParmVar) ||
+       CheckValidFPGAMemoryAttributesVar(S, D))) {
     S.Diag(AL.getLoc(), diag::err_fpga_attribute_incorrect_variable)
         << AL << /*agent memory arguments*/ 0;
     return;
-   }
+  }
 
   if (!D->hasAttr<SYCLIntelMemoryAttr>())
     D->addAttr(SYCLIntelMemoryAttr::CreateImplicit(
@@ -8107,8 +8120,9 @@ void Sema::AddSYCLIntelBankBitsAttr(Decl *D, const AttributeCommonInfo &CI,
 
   // Check attribute applies to field, constant variables, local variables,
   // static variables, agent memory arguments, non-static data members,
-  // and device_global variables.
-  if (CheckValidFPGAMemoryAttributesVar(*this, D)) {
+  // and device_global variables for the device compilation.
+  if (Context.getLangOpts().SYCLIsDevice &&
+      CheckValidFPGAMemoryAttributesVar(*this, D)) {
     Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable)
         << CI << /*agent memory arguments*/ 1;
     return;
@@ -8142,14 +8156,16 @@ void Sema::AddSYCLIntelPrivateCopiesAttr(Decl *D, const AttributeCommonInfo &CI,
 
     // Check attribute applies to field as well as const variables, non-static
     // local variables, non-static data members, and device_global variables.
+    // for the device compilation.
     if (const auto *VD = dyn_cast<VarDecl>(D)) {
-      if (!(isa<FieldDecl>(D) ||
+      if (Context.getLangOpts().SYCLIsDevice &&
+	  (!(isa<FieldDecl>(D) ||
             (VD->getKind() != Decl::ImplicitParam &&
              VD->getKind() != Decl::NonTypeTemplateParm &&
              VD->getKind() != Decl::ParmVar &&
              (VD->hasLocalStorage() ||
               isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
-                  VD->getType()))))) {
+                  VD->getType())))))) {
         Diag(CI.getLoc(), diag::err_fpga_attribute_invalid_decl) << CI;
         return;
       }
@@ -8207,8 +8223,9 @@ void Sema::AddSYCLIntelForcePow2DepthAttr(Decl *D,
 
     // Check attribute applies to field, constant variables, local variables,
     // static variables, agent memory arguments, non-static data members,
-    // and device_global variables.
-    if (CheckValidFPGAMemoryAttributesVar(*this, D)) {
+    // and device_global variables for the device compilation.
+    if (Context.getLangOpts().SYCLIsDevice &&
+        CheckValidFPGAMemoryAttributesVar(*this, D)) {
       Diag(CI.getLoc(), diag::err_fpga_attribute_incorrect_variable)
           << CI << /*agent memory arguments*/ 1;
       return;

diff --git a/clang/test/CodeGenSYCL/intel-fpga-device-global-host.cpp b/clang/test/CodeGenSYCL/intel-fpga-device-global-host.cpp
@@ -0,0 +1,95 @@
+// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-host -triple x86_64-unknown-linux-gnu -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s
+
+#include "sycl.hpp"
+
+// Tests that [[intel::numbanks()]], [[intel::fpga_register]], [[intel::private_copies()]], [[intel::doublepump]], [[intel::singlepump]], [[intel::merge()]], [[intel::fpga_memory()]], [[intel::bank_bits()]], [[intel::force_pow2_depth()]], [[intel::max_replicates()]], [[intel::bankwidth()]], [[intel::simple_dual_port]] attributes are ignored on host code.
+
+using namespace sycl::ext::oneapi;
+using namespace sycl;
+
+struct bar {
+  [[intel::numbanks(2)]] device_global<int> nonconst_glob;
+  [[intel::numbanks(4)]] const device_global<int> const_glob;
+  [[intel::numbanks(8)]] unsigned int numbanks[64];
+
+  [[intel::max_replicates(2)]] device_global<int> nonconst_glob1;
+  [[intel::max_replicates(4)]] const device_global<int> const_glob1;
+  [[intel::max_replicates(8)]] unsigned int max_rep[64];
+
+  [[intel::force_pow2_depth(0)]] device_global<int> nonconst_glob2;
+  [[intel::force_pow2_depth(0)]] const device_global<int> const_glob2;
+  [[intel::force_pow2_depth(1)]] unsigned int force_dep[64];
+
+  [[intel::bankwidth(2)]] device_global<int> nonconst_glob3;
+  [[intel::bankwidth(4)]] const device_global<int> const_glob3;
+  [[intel::bankwidth(16)]] unsigned int bankw[64];
+
+  [[intel::simple_dual_port]] device_global<int> nonconst_glob4;
+  [[intel::simple_dual_port]] const device_global<int> const_glob4;
+  [[intel::simple_dual_port]] unsigned int simple[64];
+
+  [[intel::fpga_memory]] device_global<int> nonconst_glob5;
+  [[intel::fpga_memory("MLAB")]] const device_global<int> const_glob5;
+  [[intel::fpga_memory("BLOCK_RAM")]] unsigned int mem_block_ram[32];
+
+  [[intel::bank_bits(3, 4)]] device_global<int> nonconst_glob6;
+  [[intel::bank_bits(4, 5)]] const device_global<int> const_glob6;
+  [[intel::bank_bits(3, 4)]] unsigned int mem_block_bits[32];
+
+  [[intel::fpga_register]] device_global<int> nonconst_glob7;
+  [[intel::fpga_register]] const device_global<int> const_glob7;
+  [[intel::fpga_register]] unsigned int reg;
+
+  [[intel::singlepump]] device_global<int> nonconst_glob8;
+  [[intel::singlepump]] const device_global<int> const_glob8;
+  [[intel::singlepump]] unsigned int spump;
+
+  [[intel::doublepump]] device_global<int> nonconst_glob9;
+  [[intel::doublepump]] const device_global<int> const_glob9;
+  [[intel::doublepump]] unsigned int dpump;
+
+  [[intel::merge("mrg6", "depth")]] device_global<int> nonconst_glob10;
+  [[intel::merge("mrg6", "depth")]] const device_global<int> const_glob10;
+  [[intel::merge("mrg6", "width")]] unsigned int mergewidth;
+
+  [[intel::private_copies(32)]] device_global<int> nonconst_glob11;
+  [[intel::private_copies(8)]] const device_global<int> const_glob11;
+  [[intel::private_copies(8)]] unsigned int pc;
+};
+
+[[intel::numbanks(4)]] device_global<int> nonconst_ignore;
+
+[[intel::max_replicates(8)]] device_global<int> nonconst_ignore1;
+
+[[intel::force_pow2_depth(0)]] device_global<int> nonconst_ignore2;
+
+[[intel::bankwidth(2)]] device_global<int> nonconst_ignore3;
+
+[[intel::simple_dual_port]] device_global<int> nonconst_ignore4;
+
+[[intel::fpga_memory("MLAB")]] device_global<int> nonconst_ignore5;
+
+[[intel::bank_bits(6, 7)]] device_global<int> nonconst_ignore6;
+
+[[intel::fpga_register]] device_global<int> nonconst_ignore7;
+
+[[intel::doublepump]] device_global<int> nonconst_ignore8;
+
+[[intel::singlepump]] device_global<int> nonconst_ignore9;
+
+[[intel::merge("mrg1", "depth")]] device_global<int> nonconst_ignore10;
+
+[[intel::private_copies(16)]] device_global<int> nonconst_ignore11;
+
+// CHECK-NOT: !private_copies
+// CHECK-NOT: !singlepump
+// CHECK-NOT: !doublepump
+// CHECK-NOT: !force_pow2_depth
+// CHECK-NOT: !max_replicates
+// CHECK-NOT: !numbanks
+// CHECK-NOT: !bank_bits
+// CHECK-NOT: !bankwidth
+// CHECK-NOT: !simple_dual_port
+// CHECK-NOT: !merge
+// CHECK-NOT: !fpga_memory
+// CHECK-NOT: !fpga_register
diff --git a/clang/test/SemaSYCL/intel-fpga-device-global-host.cpp b/clang/test/SemaSYCL/intel-fpga-device-global-host.cpp
@@ -0,0 +1,84 @@
+// RUN: %clang_cc1 -fsycl-is-host -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s
+
+#include "Inputs/sycl.hpp"
+
+// Tests that [[intel::numbanks()]], [[intel::fpga_register]], [[intel::private_copies()]], [[intel::doublepump]], [[intel::singlepump]], [[intel::merge()]], [[intel::fpga_memory()]], [[intel::bank_bits()]], [[intel::force_pow2_depth()]], [[intel::max_replicates()]], [[intel::bankwidth()]], [[intel::simple_dual_port]] attributes are ignored during host compilation where -fsycl-is-host is passed on cc1.
+
+// expected-no-diagnostics
+
+using namespace sycl::ext::oneapi;
+using namespace sycl;
+
+struct bar {
+  [[intel::numbanks(2)]] device_global<int> nonconst_glob;
+  [[intel::numbanks(4)]] const device_global<int> const_glob;
+  [[intel::numbanks(8)]] unsigned int numbanks[64];
+
+  [[intel::max_replicates(2)]] device_global<int> nonconst_glob1;
+  [[intel::max_replicates(4)]] const device_global<int> const_glob1;
+  [[intel::max_replicates(8)]] unsigned int max_rep[64];
+
+  [[intel::force_pow2_depth(0)]] device_global<int> nonconst_glob2;
+  [[intel::force_pow2_depth(0)]] const device_global<int> const_glob2;
+  [[intel::force_pow2_depth(1)]] unsigned int force_dep[64];
+
+  [[intel::bankwidth(2)]] device_global<int> nonconst_glob3;
+  [[intel::bankwidth(4)]] const device_global<int> const_glob3;
+  [[intel::bankwidth(16)]] unsigned int bankw[64];
+
+  [[intel::simple_dual_port]] device_global<int> nonconst_glob4;
+  [[intel::simple_dual_port]] const device_global<int> const_glob4;
+  [[intel::simple_dual_port]] unsigned int simple[64];
+
+  [[intel::fpga_memory]] device_global<int> nonconst_glob5;
+  [[intel::fpga_memory("MLAB")]] const device_global<int> const_glob5;
+  [[intel::fpga_memory("BLOCK_RAM")]] unsigned int mem_block_ram[32];
+
+  [[intel::bank_bits(3, 4)]] device_global<int> nonconst_glob6;
+  [[intel::bank_bits(4, 5)]] const device_global<int> const_glob6;
+  [[intel::bank_bits(3, 4)]] unsigned int mem_block_bits[32];
+
+  [[intel::fpga_register]] device_global<int> nonconst_glob7;
+  [[intel::fpga_register]] const device_global<int> const_glob7;
+  [[intel::fpga_register]] unsigned int reg;
+
+  [[intel::singlepump]] device_global<int> nonconst_glob8;
+  [[intel::singlepump]] const device_global<int> const_glob8;
+  [[intel::singlepump]] unsigned int spump;
+
+  [[intel::doublepump]] device_global<int> nonconst_glob9;
+  [[intel::doublepump]] const device_global<int> const_glob9;
+  [[intel::doublepump]] unsigned int dpump;
+
+  [[intel::merge("mrg6", "depth")]] device_global<int> nonconst_glob10;
+  [[intel::merge("mrg6", "depth")]] const device_global<int> const_glob10;
+  [[intel::merge("mrg6", "width")]] unsigned int mergewidth;
+
+  [[intel::private_copies(32)]] device_global<int> nonconst_glob11;
+  [[intel::private_copies(8)]] const device_global<int> const_glob11;
+  [[intel::private_copies(8)]] unsigned int pc;
+};
+
+[[intel::numbanks(4)]] device_global<int> nonconst_ignore;
+
+[[intel::max_replicates(8)]] device_global<int> nonconst_ignore1;
+
+[[intel::force_pow2_depth(0)]] device_global<int> nonconst_ignore2;
+
+[[intel::bankwidth(2)]] device_global<int> nonconst_ignore3;
+
+[[intel::simple_dual_port]] device_global<int> nonconst_ignore4;
+
+[[intel::fpga_memory("MLAB")]] device_global<int> nonconst_ignore5;
+
+[[intel::bank_bits(6, 7)]] device_global<int> nonconst_ignore6;
+
+[[intel::fpga_register]] device_global<int> nonconst_ignore7;
+
+[[intel::doublepump]] device_global<int> nonconst_ignore8;
+
+[[intel::singlepump]] device_global<int> nonconst_ignore9;
+
+[[intel::merge("mrg1", "depth")]] device_global<int> nonconst_ignore10;
+
+[[intel::private_copies(16)]] device_global<int> nonconst_ignore11;