From 3bb5f40b6584fe42c80ac9f1d8db83acbdfbc887 Mon Sep 17 00:00:00 2001 From: smanna12 Date: Wed, 6 Mar 2024 15:52:45 -0600 Subject: [PATCH] [SYCL][FPGA] Allow memory attributes on device_global variables for device compilation (#12930) https://github.com/intel/llvm/pull/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 --- clang/lib/Sema/SemaDeclAttr.cpp | 89 ++++++++++------- .../intel-fpga-device-global-host.cpp | 95 +++++++++++++++++++ .../intel-fpga-device-global-host.cpp | 84 ++++++++++++++++ 3 files changed, 232 insertions(+), 36 deletions(-) create mode 100644 clang/test/CodeGenSYCL/intel-fpga-device-global-host.cpp create mode 100644 clang/test/SemaSYCL/intel-fpga-device-global-host.cpp diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 813a6236cfa15..3f6d0be85f21b 100644 --- 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()) @@ -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()) 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(D)) { - if (!(isa(D) || + if (Context.getLangOpts().SYCLIsDevice && + (!(isa(D) || (VD->getKind() != Decl::ImplicitParam && VD->getKind() != Decl::NonTypeTemplateParm && VD->getKind() != Decl::ParmVar && (VD->hasLocalStorage() || isTypeDecoratedWithDeclAttribute( - 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 new file mode 100644 index 0000000000000..13ca7f669de40 --- /dev/null +++ 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 nonconst_glob; + [[intel::numbanks(4)]] const device_global const_glob; + [[intel::numbanks(8)]] unsigned int numbanks[64]; + + [[intel::max_replicates(2)]] device_global nonconst_glob1; + [[intel::max_replicates(4)]] const device_global const_glob1; + [[intel::max_replicates(8)]] unsigned int max_rep[64]; + + [[intel::force_pow2_depth(0)]] device_global nonconst_glob2; + [[intel::force_pow2_depth(0)]] const device_global const_glob2; + [[intel::force_pow2_depth(1)]] unsigned int force_dep[64]; + + [[intel::bankwidth(2)]] device_global nonconst_glob3; + [[intel::bankwidth(4)]] const device_global const_glob3; + [[intel::bankwidth(16)]] unsigned int bankw[64]; + + [[intel::simple_dual_port]] device_global nonconst_glob4; + [[intel::simple_dual_port]] const device_global const_glob4; + [[intel::simple_dual_port]] unsigned int simple[64]; + + [[intel::fpga_memory]] device_global nonconst_glob5; + [[intel::fpga_memory("MLAB")]] const device_global const_glob5; + [[intel::fpga_memory("BLOCK_RAM")]] unsigned int mem_block_ram[32]; + + [[intel::bank_bits(3, 4)]] device_global nonconst_glob6; + [[intel::bank_bits(4, 5)]] const device_global const_glob6; + [[intel::bank_bits(3, 4)]] unsigned int mem_block_bits[32]; + + [[intel::fpga_register]] device_global nonconst_glob7; + [[intel::fpga_register]] const device_global const_glob7; + [[intel::fpga_register]] unsigned int reg; + + [[intel::singlepump]] device_global nonconst_glob8; + [[intel::singlepump]] const device_global const_glob8; + [[intel::singlepump]] unsigned int spump; + + [[intel::doublepump]] device_global nonconst_glob9; + [[intel::doublepump]] const device_global const_glob9; + [[intel::doublepump]] unsigned int dpump; + + [[intel::merge("mrg6", "depth")]] device_global nonconst_glob10; + [[intel::merge("mrg6", "depth")]] const device_global const_glob10; + [[intel::merge("mrg6", "width")]] unsigned int mergewidth; + + [[intel::private_copies(32)]] device_global nonconst_glob11; + [[intel::private_copies(8)]] const device_global const_glob11; + [[intel::private_copies(8)]] unsigned int pc; +}; + +[[intel::numbanks(4)]] device_global nonconst_ignore; + +[[intel::max_replicates(8)]] device_global nonconst_ignore1; + +[[intel::force_pow2_depth(0)]] device_global nonconst_ignore2; + +[[intel::bankwidth(2)]] device_global nonconst_ignore3; + +[[intel::simple_dual_port]] device_global nonconst_ignore4; + +[[intel::fpga_memory("MLAB")]] device_global nonconst_ignore5; + +[[intel::bank_bits(6, 7)]] device_global nonconst_ignore6; + +[[intel::fpga_register]] device_global nonconst_ignore7; + +[[intel::doublepump]] device_global nonconst_ignore8; + +[[intel::singlepump]] device_global nonconst_ignore9; + +[[intel::merge("mrg1", "depth")]] device_global nonconst_ignore10; + +[[intel::private_copies(16)]] device_global 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 new file mode 100644 index 0000000000000..5500206bab2fe --- /dev/null +++ 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 nonconst_glob; + [[intel::numbanks(4)]] const device_global const_glob; + [[intel::numbanks(8)]] unsigned int numbanks[64]; + + [[intel::max_replicates(2)]] device_global nonconst_glob1; + [[intel::max_replicates(4)]] const device_global const_glob1; + [[intel::max_replicates(8)]] unsigned int max_rep[64]; + + [[intel::force_pow2_depth(0)]] device_global nonconst_glob2; + [[intel::force_pow2_depth(0)]] const device_global const_glob2; + [[intel::force_pow2_depth(1)]] unsigned int force_dep[64]; + + [[intel::bankwidth(2)]] device_global nonconst_glob3; + [[intel::bankwidth(4)]] const device_global const_glob3; + [[intel::bankwidth(16)]] unsigned int bankw[64]; + + [[intel::simple_dual_port]] device_global nonconst_glob4; + [[intel::simple_dual_port]] const device_global const_glob4; + [[intel::simple_dual_port]] unsigned int simple[64]; + + [[intel::fpga_memory]] device_global nonconst_glob5; + [[intel::fpga_memory("MLAB")]] const device_global const_glob5; + [[intel::fpga_memory("BLOCK_RAM")]] unsigned int mem_block_ram[32]; + + [[intel::bank_bits(3, 4)]] device_global nonconst_glob6; + [[intel::bank_bits(4, 5)]] const device_global const_glob6; + [[intel::bank_bits(3, 4)]] unsigned int mem_block_bits[32]; + + [[intel::fpga_register]] device_global nonconst_glob7; + [[intel::fpga_register]] const device_global const_glob7; + [[intel::fpga_register]] unsigned int reg; + + [[intel::singlepump]] device_global nonconst_glob8; + [[intel::singlepump]] const device_global const_glob8; + [[intel::singlepump]] unsigned int spump; + + [[intel::doublepump]] device_global nonconst_glob9; + [[intel::doublepump]] const device_global const_glob9; + [[intel::doublepump]] unsigned int dpump; + + [[intel::merge("mrg6", "depth")]] device_global nonconst_glob10; + [[intel::merge("mrg6", "depth")]] const device_global const_glob10; + [[intel::merge("mrg6", "width")]] unsigned int mergewidth; + + [[intel::private_copies(32)]] device_global nonconst_glob11; + [[intel::private_copies(8)]] const device_global const_glob11; + [[intel::private_copies(8)]] unsigned int pc; +}; + +[[intel::numbanks(4)]] device_global nonconst_ignore; + +[[intel::max_replicates(8)]] device_global nonconst_ignore1; + +[[intel::force_pow2_depth(0)]] device_global nonconst_ignore2; + +[[intel::bankwidth(2)]] device_global nonconst_ignore3; + +[[intel::simple_dual_port]] device_global nonconst_ignore4; + +[[intel::fpga_memory("MLAB")]] device_global nonconst_ignore5; + +[[intel::bank_bits(6, 7)]] device_global nonconst_ignore6; + +[[intel::fpga_register]] device_global nonconst_ignore7; + +[[intel::doublepump]] device_global nonconst_ignore8; + +[[intel::singlepump]] device_global nonconst_ignore9; + +[[intel::merge("mrg1", "depth")]] device_global nonconst_ignore10; + +[[intel::private_copies(16)]] device_global nonconst_ignore11;