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;