diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 13f6e863b541e..e1e9f3c6b17be 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -5907,10 +5907,50 @@ void CodeGenModule::generateIntelFPGAAnnotation( } } +/** + * Adds global Intel FPGA annotations for a given variable declaration. + * This function handles both simple global variables and fields within + * structs that are annotated with Intel FPGA attributes. For structs, + * it recursively visits all fields and base classes to collect annotations. + * + * @param VD The variable declaration to annotate. + * @param GV The LLVM GlobalValue corresponding to the variable declaration. + */ void CodeGenModule::addGlobalIntelFPGAAnnotation(const VarDecl *VD, llvm::GlobalValue *GV) { SmallString<256> AnnotStr; + + // Handle annotations for fields within a device_global struct. + if (getLangOpts().IntelFPGA && VD->getType()->isRecordType()) { + auto RT = VD->getType()->getAs(); + + auto Gen = [&AnnotStr, this](const RecordType *Ty, auto &&Gen) -> void { + const CXXRecordDecl *RD = cast(Ty->getDecl()); + + // Iterate over the fields of the struct. + for (const auto *Field : RD->fields()) { + generateIntelFPGAAnnotation(Field, AnnotStr); + + if (const auto *FT = + Field->getType() + ->getPointeeOrArrayElementType() // Strip pointers/arrays + ->getAs()) + Gen(FT, Gen); + } + + // Iterate over the base classes of the struct. + for (const auto Base : RD->bases()) { + QualType BaseTy = Base.getType(); + + if (const auto *BRT = BaseTy->getAs()) + Gen(BRT, Gen); + } + }; + Gen(RT, Gen); + } + generateIntelFPGAAnnotation(VD, AnnotStr); + if (!AnnotStr.empty()) { // Get the globals for file name, annotation, and the line number. llvm::Constant *AnnoGV = EmitAnnotationString(AnnotStr), diff --git a/clang/test/CodeGenSYCL/device_global.cpp b/clang/test/CodeGenSYCL/device_global.cpp index 372853ff1a185..3b70256bab767 100644 --- a/clang/test/CodeGenSYCL/device_global.cpp +++ b/clang/test/CodeGenSYCL/device_global.cpp @@ -1,5 +1,8 @@ // RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -fsycl-unique-prefix=THE_PREFIX -std=c++17 -fgpu-rdc -emit-llvm %s -o - | FileCheck %s --check-prefixes=CHECK,CHECK-RDC // RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -fsycl-unique-prefix=THE_PREFIX -std=c++17 -fno-gpu-rdc -emit-llvm %s -o - | FileCheck %s --check-prefixes=CHECK,CHECK-NORDC +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -fsycl-unique-prefix=THE_PREFIX -std=c++17 -emit-llvm %s -o - | FileCheck %s --check-prefixes=CHECK,CHECK-NOINTELFPGA +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -fsycl-unique-prefix=THE_PREFIX -std=c++17 -fintelfpga -emit-llvm %s -o - | FileCheck %s --check-prefixes=CHECK,CHECK-INTELFPGA + #include "sycl.hpp" // Test cases below show that 'sycl-unique-id' LLVM IR attribute is attached to the @@ -124,8 +127,318 @@ void foo() { namespace { device_global same_name; } + +// CHECK: [[ANN_numbanks:@.str]] = {{.*}}{memory:DEFAULT}{sizeinfo:8}{numbanks:2}{{.*}} +// CHECK: [[ANN_max_replicates:@.str.[0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:8}{max_replicates:2}{{.*}} +// CHECK: [[ANN_force_pow2_depth:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:8}{force_pow2_depth:1}{{.*}} +// CHECK: [[ANN_bankwidth:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:8}{bankwidth:2}{{.*}} +// CHECK: [[ANN_simple_dual_port:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:8}{simple_dual_port:1}{{.*}} +// CHECK: [[ANN_memory_default:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:8}{{.*}} +// CHECK: [[ANN_bank_bits:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:8}{numbanks:4}{bank_bits:3,4}{{.*}} +// CHECK: [[ANN_register:@.str[.0-9]*]] = {{.*}}{register:1}{{.*}} +// CHECK: [[ANN_doublepump:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:8}{pump:2}{{.*}} +// CHECK: [[ANN_singlepump:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:8}{pump:1}{{.*}} +// CHECK: [[ANN_merge:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:8}{merge:mrg5:width}{{.*}} +// CHECK: [[ANN_private_copies:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:8}{private_copies:8}{{.*}} +// CHECK: @counter = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global.3" zeroinitializer, align 8 #[[DEV_GLOB_FPGA_ATTRS:[0-9]+]] +// CHECK: @counter1 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global.4" zeroinitializer, align 8 #[[DEV_GLOB_FPGA_ATTRS1:[0-9]+]] +// CHECK: [[ANN_max_replicates1:@.str.[0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:4,155}{max_replicates:2}{{.*}} +// CHECK: @counter2 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global.5" zeroinitializer, align 8 #[[DEV_GLOB_FPGA_ATTRS2:[0-9]+]] +// CHECK: [[ANN_bankwidth1:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:4,155}{bankwidth:2}{{.*}} +// CHECK: @counter3 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global.6" zeroinitializer, align 8 #[[DEV_GLOB_FPGA_ATTRS3:[0-9]+]] +// CHECK: [[ANN_memory_default1:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:4,155}{{.*}} +// CHECK: @counter4 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global.7" zeroinitializer, align 8 #[[DEV_GLOB_FPGA_ATTRS4:[0-9]+]] +// CHECK: [[ANN_numbanks1:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:4,155}{numbanks:2}{{.*}} +// CHECK: @counter5 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global.8" zeroinitializer, align 8 #[[DEV_GLOB_FPGA_ATTRS5:[0-9]+]] +// CHECK: [[ANN_force_pow2_depth1:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:4,155}{force_pow2_depth:1}{{.*}} +// CHECK: @counter6 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global.9" zeroinitializer, align 8 #[[DEV_GLOB_FPGA_ATTRS6:[0-9]+]] +// CHECK: [[ANN_bank_bits1:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:4,155}{numbanks:4}{bank_bits:3,4}{{.*}} +// CHECK: @counter7 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global.10" zeroinitializer, align 8 #[[DEV_GLOB_FPGA_ATTRS7:[0-9]+]] +// CHECK: [[ANN_doublepump1:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:4,155}{pump:2}{{.*}} +// CHECK: @counter8 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global.11" zeroinitializer, align 8 #[[DEV_GLOB_FPGA_ATTRS8:[0-9]+]] +// CHECK: [[ANN_singlepump1:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:4,155}{pump:1}{{.*}} +// CHECK: @counter9 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global.12" zeroinitializer, align 8 #[[DEV_GLOB_FPGA_ATTRS9:[0-9]+]] +// CHECK: [[ANN_merge1:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:4,155}{merge:mrg5:width}{{.*}} +// CHECK: @counter10 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global.13" zeroinitializer, align 8 #[[DEV_GLOB_FPGA_ATTRS10:[0-9]+]] +// CHECK: [[ANN_private_copies1:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:4,155}{private_copies:8}{{.*}} +// CHECK: @counter11 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global.14" zeroinitializer, align 8 #[[DEV_GLOB_FPGA_ATTRS11:[0-9]+]] +// CHECK: [[ANN_simple_dual_port1:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:4,155}{simple_dual_port:1}{{.*}} +// CHECK: @counter12 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global.15" zeroinitializer, align 8 #[[DEV_GLOB_FPGA_ATTRS12:[0-9]+]] +// CHECK: [[ANN_simple_dual_port2:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:4}{simple_dual_port:1}{{.*}} +// CHECK: [[ANN_private_copies2:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:4}{private_copies:16}{{.*}} +// CHECK: @counter13 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global.16" zeroinitializer, align 8 #[[DEV_GLOB_FPGA_ATTRS13:[0-9]+]] +// CHECK: [[ANN_numbanks2:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:4}{numbanks:2}{{.*}} +// CHECK: [[ANN_merge2:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:4,155}{merge:foo:depth}{{.*}} +// CHECK: @counter14 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global.17" zeroinitializer, align 8 #[[DEV_GLOB_FPGA_ATTRS14:[0-9]+]] +// CHECK: [[ANN_force_pow2_depth2:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:4}{force_pow2_depth:0}{{.*}} +// CHECK: [[ANN_doublepump2:@.str[.0-9]*]] = {{.*}}"{memory:DEFAULT}{sizeinfo:4}{pump:2}{{.*}} +// CHECK: [[ANN_bank_bits2:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:4,155}{numbanks:4}{bank_bits:2,3}{{.*}} +// CHECK: @counter15 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global.18" zeroinitializer, align 8 #[[DEV_GLOB_FPGA_ATTRS15:[0-9]+]] +// CHECK-INTELFPGA: @global_wrapper = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global.19" zeroinitializer, align 8 #[[DEV_GLOB_FPGA_ATTRS16:[0-9]+]] +// CHECK-INTELFPGA: [[ANN_memory3:@.str[.0-9]*]] = {{.*}}{memory:MLAB}{sizeinfo:48}{{.*}} +// CHECK-INTELFPGA: [[ANN_numbanks4:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:4,10}{numbanks:4}{{.*}} +// CHECK-INTELFPGA: @global_wrapper1 = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global.20" zeroinitializer, align 8 #[[DEV_GLOB_FPGA_ATTRS17:[0-9]+]] +// CHECK-INTELFPGA: [[ANN_private_copies4:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:48}{private_copies:16}{{.*}} +// CHECK-INTELFPGA: [[ANN_simple_dual_port4:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:4,10}{simple_dual_port:1}{{.*}} +// CHECK-INTELFPGA: [[ANN_private_copies_imple_dual_port2:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:4}{simple_dual_port:1}{memory:DEFAULT}{sizeinfo:4}{private_copies:16}{{.*}} +// CHECK-INTELFPGA: [[ANN_numbanks3_merge3:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:4}{numbanks:2}{memory:DEFAULT}{sizeinfo:4,155}{merge:foo:depth}{{.*}} +// CHECK-INTELFPGA: [[ANN_force_pow2_depth3_doublepump3_numbanks3:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:4}{force_pow2_depth:0}{memory:DEFAULT}{sizeinfo:4}{pump:2}{memory:DEFAULT}{sizeinfo:4,155}{numbanks:4}{bank_bits:2,3}{{.*}} +// CHECK-INTELFPGA: [[ANN_memory4_numbanks4_register4:@.str[.0-9]*]] = {{.*}}{memory:MLAB}{sizeinfo:48}{memory:DEFAULT}{sizeinfo:4,10}{numbanks:4}{register:1}{register:1}{{.*}} +// CHECK-INTELFPGA: [[ANN_private_copies5_simple_dual_port5_register5:@.str[.0-9]*]] = {{.*}}{memory:DEFAULT}{sizeinfo:48}{private_copies:16}{memory:DEFAULT}{sizeinfo:4,10}{simple_dual_port:1}{register:1}{{.*}} // CHECK: @_ZN12_GLOBAL__N_19same_nameE = internal addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[SAME_NAME_ANON_NS_ATTRS:[0-9]+]] +struct bar { + int t1; + int t2; + [[intel::fpga_register]] int arr[155]; +}; + +device_global counter; + +struct bar1 { + int t3; + int t4; + [[intel::max_replicates(2)]] int arr1[155]; +}; + +device_global counter1; + +struct bar2 { + int t5; + int t6; + [[intel::bankwidth(2)]] int arr2[155]; +}; + +device_global counter2; + +struct bar3 { + int t7; + int t8; + [[intel::fpga_memory]] int arr3[155]; +}; + +device_global counter3; + +struct bar4 { + int t9; + int t10; + [[intel::numbanks(2)]] int arr4[155]; +}; + +device_global counter4; + +struct bar5 { + int t11; + int t12; + [[intel::force_pow2_depth(1)]] int arr5[155]; +}; + +device_global counter5; + +struct bar6 { + int t13; + int t14; + [[intel::bank_bits(3, 4)]] int arr6[155]; +}; + +device_global counter6; + +struct bar7 { + int t15; + int t16; + [[intel::doublepump]] int arr7[155]; +}; + +device_global counter7; + +struct bar8 { + int t17; + int t18; + [[intel::singlepump]] int arr8[155]; +}; + +device_global counter8; + +struct bar9 { + int t19; + int t20; + [[intel::merge("mrg5", "width")]] int arr9[155]; +}; + +device_global counter9; + +struct bar10 { + int t21; + int t22; + [[intel::private_copies(8)]] int arr10[155]; +}; + +device_global counter10; + +struct bar11 { + int t23; + int t24; + [[intel::simple_dual_port]] int arr11[155]; +}; + +device_global counter11; + +struct bar12 { + [[intel::simple_dual_port]] int t25; + [[intel::private_copies(16)]] int t26; +}; + +device_global counter12; + +struct bar13 { + [[intel::numbanks(2)]] int t27; + [[intel::merge("foo", "depth")]] int arr12[155]; +}; + +device_global counter13; + +struct bar14 { + [[intel::force_pow2_depth(0)]] int t28; + [[intel::doublepump]] int t29; + [[intel::bank_bits(2, 3)]] int arr13[155]; +}; + +device_global counter14; + +struct bar15 { + int t30; + int t31; + int arr14[155]; +}; + +[[intel::fpga_register]] device_global counter15; + +// Base class with different attributes +class Base { +public: + [[intel::fpga_register]] int reg_attr; + int no_attr; +}; + +// Derived class with additional attributes +class Derived : public Base { +public: + [[intel::numbanks(4)]] int arr_attr[10]; +}; + +// Class with class type member with attributes +class Wrapper { +public: + [[intel::fpga_memory("MLAB")]] Derived derived_attr; +}; + +// Global instance with FPGA attributes +[[intel::fpga_register]] device_global global_wrapper; + +// Base class with different attributes +class Base1 { +public: + [[intel::fpga_register]] int reg_attr1; + int no_attr1; +}; + +// Derived class with additional attributes +class Derived1 : public Base1 { +public: + [[intel::simple_dual_port]] int arr_attr1[10]; +}; + +// Class with class type member with attributes +class Wrapper1 { +public: + [[intel::private_copies(16)]] Derived1 derived_attr1; +}; + +// Global instance with FPGA attributes +device_global global_wrapper1; + + +int main() { + queue q; + + q.submit([&](handler &h) { + h.single_task([=] { + + auto& non_const_counter = const_cast(counter.get()); + non_const_counter.arr[0]++; + + auto& non_const_counter1 = const_cast(counter1.get()); + non_const_counter1.arr1[0]++; + + auto& non_const_counter2 = const_cast(counter2.get()); + non_const_counter2.arr2[0]++; + + auto& non_const_counter3 = const_cast(counter3.get()); + non_const_counter3.arr3[0]++; + + auto& non_const_counter4 = const_cast(counter4.get()); + non_const_counter4.arr4[0]++; + + auto& non_const_counter5 = const_cast(counter5.get()); + non_const_counter5.arr5[0]++; + + auto& non_const_counter6 = const_cast(counter6.get()); + non_const_counter6.arr6[0]++; + + auto& non_const_counter7 = const_cast(counter7.get()); + non_const_counter7.arr7[0]++; + + auto& non_const_counter8 = const_cast(counter8.get()); + non_const_counter8.arr8[0]++; + + auto& non_const_counter9 = const_cast(counter9.get()); + non_const_counter9.arr9[0]++; + + auto& non_const_counter10 = const_cast(counter10.get()); + non_const_counter10.arr10[0]++; + + auto& non_const_counter11 = const_cast(counter11.get()); + non_const_counter11.arr11[0]++; + + auto& non_const_counter12 = const_cast(counter12.get()); + non_const_counter12.t25 = 5; + non_const_counter12.t26 = 20; + + auto& non_const_counter13 = const_cast(counter13.get()); + non_const_counter13.t27 = 30; + non_const_counter13.arr12[0]++; + + auto& non_const_counter14 = const_cast(counter14.get()); + non_const_counter14.t28 = 35; + non_const_counter14.t29 = 36; + non_const_counter14.arr13[0]++; + + auto& non_const_counter15 = const_cast(counter15.get()); + non_const_counter15.t30 = 5; + non_const_counter15.t31 = 9; + non_const_counter15.arr14[0]++; + + auto& non_const_global_wrapper = const_cast(global_wrapper.get()); + non_const_global_wrapper.derived_attr.reg_attr = 5; + non_const_global_wrapper.derived_attr.no_attr = 10; + for (int i = 0; i < 10; ++i) { + non_const_global_wrapper.derived_attr.arr_attr[i] = i; + } + + auto& non_const_global_wrapper1 = const_cast(global_wrapper1.get()); + non_const_global_wrapper1.derived_attr1.reg_attr1 = 3; + non_const_global_wrapper1.derived_attr1.no_attr1 = 20; + for (int i = 0; i < 10; ++i) { + non_const_global_wrapper1.derived_attr1.arr_attr1[i] = i; + } + + }); + }); + + q.wait(); + return 0; +} + namespace { void bar() { q.submit([&](handler &h) { @@ -135,6 +448,38 @@ void bar() { } // namespace // CHECK: @llvm.global_ctors = appending global [2 x { i32, ptr, ptr addrspace(4) }] [{ i32, ptr, ptr addrspace(4) } { i32 65535, ptr @__cxx_global_var_init{{.*}}, ptr addrspace(4) addrspacecast (ptr addrspace(1) @[[TEMPL_DEV_GLOB]] to ptr addrspace(4)) }, { i32, ptr, ptr addrspace(4) } { i32 65535, ptr @_GLOBAL__sub_I_device_global.cpp, ptr addrspace(4) null }] + +// CHECK: @llvm.global.annotations +// CHECK-NOINTELFPGA-SAME: ptr addrspace(1) @Nonconst_glob, ptr addrspace(1) [[ANN_numbanks]]{{.*}} i32 18, ptr addrspace(1) null +// CHECK-NOINTELFPGA-SAME: ptr addrspace(1) @Nonconst_glob1, ptr addrspace(1) [[ANN_max_replicates]]{{.*}} i32 19, ptr addrspace(1) null +// CHECK-NOINTELFPGA-SAME: ptr addrspace(1) @Nonconst_glob2, ptr addrspace(1) [[ANN_force_pow2_depth]]{{.*}} i32 20, ptr addrspace(1) null +// CHECK-NOINTELFPGA-SAME: ptr addrspace(1) @Nonconst_glob3, ptr addrspace(1) [[ANN_bankwidth]]{{.*}} i32 21, ptr addrspace(1) null +// CHECK-NOINTELFPGA-SAME: ptr addrspace(1) @Nonconst_glob4, ptr addrspace(1) [[ANN_simple_dual_port]]{{.*}} i32 22, ptr addrspace(1) null +// CHECK-NOINTELFPGA-SAME: ptr addrspace(1) @Nonconst_glob5, ptr addrspace(1) [[ANN_memory_default]]{{.*}} i32 23, ptr addrspace(1) null +// CHECK-NOINTELFPGA-SAME: ptr addrspace(1) @Nonconst_glob6, ptr addrspace(1) [[ANN_bank_bits]]{{.*}} i32 24, ptr addrspace(1) null +// CHECK-NOINTELFPGA-SAME: ptr addrspace(1) @Nonconst_glob7, ptr addrspace(1) [[ANN_register]]{{.*}} i32 25, ptr addrspace(1) null +// CHECK-NOINTELFPGA-SAME: ptr addrspace(1) @Nonconst_glob8, ptr addrspace(1) [[ANN_doublepump]]{{.*}} i32 26, ptr addrspace(1) null +// CHECK-NOINTELFPGA-SAME: ptr addrspace(1) @Nonconst_glob9, ptr addrspace(1) [[ANN_singlepump]]{{.*}} i32 27, ptr addrspace(1) null +// CHECK-NOINTELFPGA-SAME: ptr addrspace(1) @Nonconst_glob10, ptr addrspace(1) [[ANN_merge]]{{.*}} i32 28, ptr addrspace(1) null +// CHECK-NOINTELFPGA-SAME: ptr addrspace(1) @Nonconst_glob11, ptr addrspace(1) [[ANN_private_copies]]{{.*}} i32 29, ptr addrspace(1) null +// CHECK-INTELFPGA-SAME: ptr addrspace(1) @counter, ptr addrspace(1) [[ANN_register]]{{.*}}, i32 196, ptr addrspace(1) null +// CHECK-INTELFPGA-SAME: ptr addrspace(1) @counter1, ptr addrspace(1) [[ANN_max_replicates1]]{{.*}}, i32 204, ptr addrspace(1) null +// CHECK-INTELFPGA-SAME: ptr addrspace(1) @counter2, ptr addrspace(1) [[ANN_bankwidth1]]{{.*}}, i32 212, ptr addrspace(1) null +// CHECK-INTELFPGA-SAME: ptr addrspace(1) @counter3, ptr addrspace(1) [[ANN_memory_default1]]{{.*}}, i32 220, ptr addrspace(1) null +// CHECK-INTELFPGA-SAME: ptr addrspace(1) @counter4, ptr addrspace(1) [[ANN_numbanks1]]{{.*}}, i32 228, ptr addrspace(1) null +// CHECK-INTELFPGA-SAME: ptr addrspace(1) @counter5, ptr addrspace(1) [[ANN_force_pow2_depth1]]{{.*}}, i32 236, ptr addrspace(1) null +// CHECK-INTELFPGA-SAME: ptr addrspace(1) @counter6, ptr addrspace(1) [[ANN_bank_bits1]]{{.*}}, i32 244, ptr addrspace(1) null +// CHECK-INTELFPGA-SAME: ptr addrspace(1) @counter7, ptr addrspace(1) [[ANN_doublepump1]]{{.*}}, i32 252, ptr addrspace(1) null +// CHECK-INTELFPGA-SAME: ptr addrspace(1) @counter8, ptr addrspace(1) [[ANN_singlepump1]]{{.*}}, i32 260, ptr addrspace(1) null +// CHECK-INTELFPGA-SAME: ptr addrspace(1) @counter9, ptr addrspace(1) [[ANN_merge1]]{{.*}}, i32 268, ptr addrspace(1) null +// CHECK-INTELFPGA-SAME: ptr addrspace(1) @counter10, ptr addrspace(1) [[ANN_private_copies1]]{{.*}}, i32 276, ptr addrspace(1) null +// CHECK-INTELFPGA-SAME: ptr addrspace(1) @counter11, ptr addrspace(1) [[ANN_simple_dual_port1]]{{.*}}, i32 284, ptr addrspace(1) null +// CHECK-INTELFPGA-SAME: ptr addrspace(1) @counter12, ptr addrspace(1) [[ANN_private_copies_imple_dual_port2]]{{.*}}, i32 291, ptr addrspace(1) null +// CHECK-INTELFPGA-SAME: ptr addrspace(1) @counter13, ptr addrspace(1) [[ANN_numbanks3_merge3]]{{.*}}, i32 298, ptr addrspace(1) null +// CHECK-INTELFPGA-SAME: ptr addrspace(1) @counter14, ptr addrspace(1) [[ANN_force_pow2_depth3_doublepump3_numbanks3]]{{.*}}, i32 306, ptr addrspace(1) null +// CHECK-INTELFPGA-SAME: ptr addrspace(1) @counter15, ptr addrspace(1) [[ANN_register]]{{.*}}, i32 314, ptr addrspace(1) null +// CHECK-INTELFPGA-SAME: ptr addrspace(1) @global_wrapper, ptr addrspace(1) [[ANN_memory4_numbanks4_register4]]{{.*}}i32 336, ptr addrspace(1) null +// CHECK-INTELFPGA-SAME: ptr addrspace(1) @global_wrapper1, ptr addrspace(1) [[ANN_private_copies5_simple_dual_port5_register5]]{{.*}}i32 358, ptr addrspace(1) null // CHECK: @llvm.used = appending addrspace(1) global [1 x ptr addrspace(4)] [ptr addrspace(4) addrspacecast (ptr addrspace(1) @[[TEMPL_DEV_GLOB]] to ptr addrspace(4))], section "llvm.metadata" // CHECK: @llvm.compiler.used = appending addrspace(1) global [2 x ptr addrspace(4)] // CHECK-SAME: @_ZL1B @@ -161,4 +506,23 @@ void bar() { // CHECK: attributes #[[BAR_INLINENS_ATTRS]] = { "sycl-unique-id"="_ZN3Bar8InlineNSE" } // CHECK: attributes #[[TEMPLATED_WRAPPER_ATTRS]] = { "sycl-unique-id"="_ZN2TSIiE1dE" } // CHECK: attributes #[[TEMPL_DEV_GLOB_ATTRS]] = { "sycl-unique-id"="_Z16templ_dev_globalIiE" } +// CHECK: attributes #[[DEV_GLOB_FPGA_ATTRS]] = { "sycl-unique-id"="_Z7counter" } +// CHECK: attributes #[[DEV_GLOB_FPGA_ATTRS1]] = { "sycl-unique-id"="_Z8counter1" } +// CHECK: attributes #[[DEV_GLOB_FPGA_ATTRS2]] = { "sycl-unique-id"="_Z8counter2" } +// CHECK: attributes #[[DEV_GLOB_FPGA_ATTRS3]] = { "sycl-unique-id"="_Z8counter3" } +// CHECK: attributes #[[DEV_GLOB_FPGA_ATTRS4]] = { "sycl-unique-id"="_Z8counter4" } +// CHECK: attributes #[[DEV_GLOB_FPGA_ATTRS5]] = { "sycl-unique-id"="_Z8counter5" } +// CHECK: attributes #[[DEV_GLOB_FPGA_ATTRS6]] = { "sycl-unique-id"="_Z8counter6" } +// CHECK: attributes #[[DEV_GLOB_FPGA_ATTRS7]] = { "sycl-unique-id"="_Z8counter7" } +// CHECK: attributes #[[DEV_GLOB_FPGA_ATTRS8]] = { "sycl-unique-id"="_Z8counter8" } +// CHECK: attributes #[[DEV_GLOB_FPGA_ATTRS9]] = { "sycl-unique-id"="_Z8counter9" } +// CHECK: attributes #[[DEV_GLOB_FPGA_ATTRS10]] = { "sycl-unique-id"="_Z9counter10" } +// CHECK: attributes #[[DEV_GLOB_FPGA_ATTRS11]] = { "sycl-unique-id"="_Z9counter11" } +// CHECK: attributes #[[DEV_GLOB_FPGA_ATTRS12]] = { "sycl-unique-id"="_Z9counter12" } +// CHECK: attributes #[[DEV_GLOB_FPGA_ATTRS13]] = { "sycl-unique-id"="_Z9counter13" } +// CHECK: attributes #[[DEV_GLOB_FPGA_ATTRS14]] = { "sycl-unique-id"="_Z9counter14" } +// CHECK: attributes #[[DEV_GLOB_FPGA_ATTRS15]] = { "sycl-unique-id"="_Z9counter15" } +// CHECK-INTELFPGA: attributes #[[DEV_GLOB_FPGA_ATTRS16]] = { "sycl-unique-id"="_Z14global_wrapper" } +// CHECK-INTELFPGA: attributes #[[DEV_GLOB_FPGA_ATTRS17]] = { "sycl-unique-id"="_Z15global_wrapper1" } + // CHECK: attributes #[[SAME_NAME_ANON_NS_ATTRS]] = { "sycl-unique-id"="THE_PREFIX____ZN12_GLOBAL__N_19same_nameE" }