diff --git a/clang/lib/CodeGen/BackendUtil.cpp b/clang/lib/CodeGen/BackendUtil.cpp index d002aeb5b247b..3e91d1f53d7ed 100644 --- a/clang/lib/CodeGen/BackendUtil.cpp +++ b/clang/lib/CodeGen/BackendUtil.cpp @@ -1153,8 +1153,9 @@ void EmitAssemblyHelper::RunOptimizationPipeline( // Process properties and annotations MPM.addPass(CompileTimePropertiesPass()); - // Remove SYCL metadata added by the frontend, like sycl_aspects - // Note, this pass should be at the end of the pipeline + // Remove SYCL metadata added by the frontend, like + // sycl_types_that_use_aspects Note, this pass should be at the end of the + // pipeline MPM.addPass(CleanupSYCLMetadataPass()); } } diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 9a352247c3f27..fd43d36c3ba01 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -853,17 +853,6 @@ static llvm::MDNode *getAspectsMD(ASTContext &ASTContext, return llvm::MDNode::get(Ctx, AspectsMD); } -static llvm::MDNode *getAspectEnumValueMD(ASTContext &ASTContext, - llvm::LLVMContext &Ctx, - const EnumConstantDecl *ECD) { - SmallVector AspectEnumValMD; - AspectEnumValMD.push_back(llvm::MDString::get(Ctx, ECD->getName())); - AspectEnumValMD.push_back( - llvm::ConstantAsMetadata::get(llvm::ConstantInt::get( - llvm::Type::getInt32Ty(Ctx), ECD->getInitVal().getSExtValue()))); - return llvm::MDNode::get(Ctx, AspectEnumValMD); -} - static bool isStackProtectorOn(const LangOptions &LangOpts, const llvm::Triple &Triple, clang::LangOptions::StackProtectorMode Mode) { @@ -1337,15 +1326,6 @@ void CodeGenModule::Release() { RD->getAttr())); } } - - // Emit metadata for all aspects defined in the aspects enum. - if (AspectsEnumDecl) { - llvm::NamedMDNode *AspectEnumValsMD = - TheModule.getOrInsertNamedMetadata("sycl_aspects"); - for (const EnumConstantDecl *ECD : AspectsEnumDecl->enumerators()) - AspectEnumValsMD->addOperand( - getAspectEnumValueMD(Context, TheModule.getContext(), ECD)); - } } // HLSL related end of code gen work items. @@ -5720,16 +5700,6 @@ void CodeGenModule::maybeSetTrivialComdat(const Decl &D, GO.setComdat(TheModule.getOrInsertComdat(GO.getName())); } -void CodeGenModule::setAspectsEnumDecl(const EnumDecl *ED) { - if (AspectsEnumDecl && AspectsEnumDecl != ED) { - // Conflicting definitions of the aspect enum are not allowed. - Error(ED->getLocation(), "redefinition of aspect enum"); - getDiags().Report(AspectsEnumDecl->getLocation(), - diag::note_previous_definition); - } - AspectsEnumDecl = ED; -} - void CodeGenModule::generateIntelFPGAAnnotation( const Decl *D, llvm::SmallString<256> &AnnotStr) { llvm::raw_svector_ostream Out(AnnotStr); diff --git a/clang/lib/CodeGen/CodeGenModule.h b/clang/lib/CodeGen/CodeGenModule.h index e5f4949599135..fa45dc2bfd906 100644 --- a/clang/lib/CodeGen/CodeGenModule.h +++ b/clang/lib/CodeGen/CodeGenModule.h @@ -617,7 +617,6 @@ class CodeGenModule : public CodeGenTypeCache { MetadataTypeMap GeneralizedMetadataIdMap; llvm::DenseMap TypesWithAspects; - const EnumDecl *AspectsEnumDecl = nullptr; // Helps squashing blocks of TopLevelStmtDecl into a single llvm::Function // when used with -fincremental-extensions. std::pair, const TopLevelStmtDecl *> @@ -1122,8 +1121,6 @@ class CodeGenModule : public CodeGenTypeCache { TypesWithAspects[TypeName] = RD; } - void setAspectsEnumDecl(const EnumDecl *ED); - void generateIntelFPGAAnnotation(const Decl *D, llvm::SmallString<256> &AnnotStr); void addGlobalIntelFPGAAnnotation(const VarDecl *VD, llvm::GlobalValue *GV); diff --git a/clang/lib/CodeGen/CodeGenTypes.cpp b/clang/lib/CodeGen/CodeGenTypes.cpp index 88b6a3a19044a..e038be50e701e 100644 --- a/clang/lib/CodeGen/CodeGenTypes.cpp +++ b/clang/lib/CodeGen/CodeGenTypes.cpp @@ -178,10 +178,6 @@ void CodeGenTypes::UpdateCompletedType(const TagDecl *TD) { if (!ConvertType(ED->getIntegerType())->isIntegerTy(32)) TypeCache.clear(); } - // If this is the SYCL aspect enum it is saved for later processing. - if (const auto *Attr = ED->getAttr()) - if (Attr->getType() == SYCLTypeAttr::SYCLType::aspect) - CGM.setAspectsEnumDecl(ED); // If necessary, provide the full definition of a type only used with a // declaration so far. if (CGDebugInfo *DI = CGM.getModuleDebugInfo()) diff --git a/clang/test/CodeGenSYCL/aspect_enum.cpp b/clang/test/CodeGenSYCL/aspect_enum.cpp deleted file mode 100644 index 30a950b33cfa9..0000000000000 --- a/clang/test/CodeGenSYCL/aspect_enum.cpp +++ /dev/null @@ -1,14 +0,0 @@ -// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s - -// Tests for IR of [[__sycl_detail__::sycl_type(aspect)]] enum. -#include "sycl.hpp" - -// CHECK: !sycl_aspects = !{![[HOST:[0-9]+]], ![[CPU:[0-9]+]], ![[GPU:[0-9]+]], ![[ACC:[0-9]+]], ![[CUSTOM:[0-9]+]], ![[FP16:[0-9]+]], ![[FP64:[0-9]+]], ![[PRIVATE_ALLOCA:[0-9]+]]} -// CHECK: ![[HOST]] = !{!"host", i32 0} -// CHECK: ![[CPU]] = !{!"cpu", i32 1} -// CHECK: ![[GPU]] = !{!"gpu", i32 2} -// CHECK: ![[ACC]] = !{!"accelerator", i32 3} -// CHECK: ![[CUSTOM]] = !{!"custom", i32 4} -// CHECK: ![[FP16]] = !{!"fp16", i32 5} -// CHECK: ![[FP64]] = !{!"fp64", i32 6} -// CHECK: ![[PRIVATE_ALLOCA]] = !{!"ext_oneapi_private_alloca", i32 7} diff --git a/clang/test/CodeGenSYCL/multiple_aspect_enum.cpp b/clang/test/CodeGenSYCL/multiple_aspect_enum.cpp deleted file mode 100644 index 5ab9cc6db97d3..0000000000000 --- a/clang/test/CodeGenSYCL/multiple_aspect_enum.cpp +++ /dev/null @@ -1,12 +0,0 @@ -// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -triple spir64-unknown-unknown -verify -emit-llvm-only %s - -// Tests for error diagnostics when multiple definitions of -// [[__sycl_detail__::sycl_type(aspect)]] enums are present. -#include "sycl.hpp" - -// expected-note@#AspectEnum{{previous definition is here}} - -// expected-error@+1{{redefinition of aspect enum}} -enum class [[__sycl_detail__::sycl_type(aspect)]] aspect_redef { - imposter_value = 3 -}; diff --git a/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.hpp b/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.hpp index 6268ecc0b26f8..83cf041f03f3c 100644 --- a/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.hpp +++ b/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.hpp @@ -29,7 +29,13 @@ struct TargetInfo { }; using TargetTable_t = std::map; +// AspectTable maps aspect names to their corresponding enum value +// defined in the SYCL headers +using AspectTable_t = std::map; + #define GET_TargetTable_IMPL +#define GET_AspectTable_IMPL #include "llvm/SYCLLowerIR/DeviceConfigFile.inc" #undef GET_TargetTable_IMPL +#undef GET_AspectTable_IMPL } // namespace DeviceConfigFile diff --git a/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td b/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td index aaa55d3686ac5..394c53e42d1b3 100644 --- a/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td +++ b/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td @@ -8,81 +8,87 @@ include "llvm/TableGen/DynamicTable.td" -class Aspect { +class Aspect { string Name = name; + int Value = value; } -def AspectCpu : Aspect<"cpu">; -def AspectGpu : Aspect<"gpu">; -def AspectAccelerator : Aspect<"accelerator">; -def AspectCustom : Aspect<"custom">; -def AspectFp16 : Aspect<"fp16">; -def AspectFp64 : Aspect<"fp64">; -def AspectImage : Aspect<"image">; -def AspectOnline_compiler : Aspect<"online_compiler">; -def AspectOnline_linker : Aspect<"online_linker">; -def AspectQueue_profiling : Aspect<"queue_profiling">; -def AspectUsm_device_allocations : Aspect<"usm_device_allocations">; -def AspectUsm_host_allocations : Aspect<"usm_host_allocations">; -def AspectUsm_shared_allocations : Aspect<"usm_shared_allocations">; -def AspectUsm_system_allocations : Aspect<"usm_system_allocations">; -def AspectExt_intel_pci_address : Aspect<"ext_intel_pci_address">; -def AspectExt_intel_gpu_eu_count : Aspect<"ext_intel_gpu_eu_count">; -def AspectExt_intel_gpu_eu_simd_width : Aspect<"ext_intel_gpu_eu_simd_width">; -def AspectExt_intel_gpu_slices : Aspect<"ext_intel_gpu_slices">; -def AspectExt_intel_gpu_subslices_per_slice : Aspect<"ext_intel_gpu_subslices_per_slice">; -def AspectExt_intel_gpu_eu_count_per_subslice : Aspect<"ext_intel_gpu_eu_count_per_subslice">; -def AspectExt_intel_max_mem_bandwidth : Aspect<"ext_intel_max_mem_bandwidth">; -def AspectExt_intel_mem_channel : Aspect<"ext_intel_mem_channel">; -def AspectUsm_atomic_host_allocations : Aspect<"usm_atomic_host_allocations">; -def AspectUsm_atomic_shared_allocations : Aspect<"usm_atomic_shared_allocations">; -def AspectAtomic64 : Aspect<"atomic64">; -def AspectExt_intel_device_info_uuid : Aspect<"ext_intel_device_info_uuid">; -def AspectExt_oneapi_srgb : Aspect<"ext_oneapi_srgb">; -def AspectExt_oneapi_native_assert : Aspect<"ext_oneapi_native_assert">; -def AspectHost_debuggable : Aspect<"host_debuggable">; -def AspectExt_intel_gpu_hw_threads_per_eu : Aspect<"ext_intel_gpu_hw_threads_per_eu">; -def AspectExt_oneapi_cuda_async_barrier : Aspect<"ext_oneapi_cuda_async_barrier">; -def AspectExt_oneapi_bfloat16_math_functions : Aspect<"ext_oneapi_bfloat16_math_functions">; -def AspectExt_intel_free_memory : Aspect<"ext_intel_free_memory">; -def AspectExt_intel_device_id : Aspect<"ext_intel_device_id">; -def AspectExt_intel_memory_clock_rate : Aspect<"ext_intel_memory_clock_rate">; -def AspectExt_intel_memory_bus_width : Aspect<"ext_intel_memory_bus_width">; -def AspectEmulated : Aspect<"emulated">; -def AspectExt_intel_legacy_image : Aspect<"ext_intel_legacy_image">; -def AspectExt_oneapi_bindless_images : Aspect<"ext_oneapi_bindless_images">; -def AspectExt_oneapi_bindless_images_shared_usm : Aspect<"ext_oneapi_bindless_images_shared_usm">; -def AspectExt_oneapi_bindless_images_1d_usm : Aspect<"ext_oneapi_bindless_images_1d_usm">; -def AspectExt_oneapi_bindless_images_2d_usm : Aspect<"ext_oneapi_bindless_images_2d_usm">; -def AspectExt_oneapi_interop_memory_import : Aspect<"ext_oneapi_interop_memory_import">; -def AspectExt_oneapi_interop_memory_export : Aspect<"ext_oneapi_interop_memory_export">; -def AspectExt_oneapi_interop_semaphore_import : Aspect<"ext_oneapi_interop_semaphore_import">; -def AspectExt_oneapi_interop_semaphore_export : Aspect<"ext_oneapi_interop_semaphore_export">; -def AspectExt_oneapi_mipmap : Aspect<"ext_oneapi_mipmap">; -def AspectExt_oneapi_mipmap_anisotropy : Aspect<"ext_oneapi_mipmap_anisotropy">; -def AspectExt_oneapi_mipmap_level_reference : Aspect<"ext_oneapi_mipmap_level_reference">; -def AspectExt_intel_esimd : Aspect<"ext_intel_esimd">; -def AspectExt_oneapi_ballot_group : Aspect<"ext_oneapi_ballot_group">; -def AspectExt_oneapi_fixed_size_group : Aspect<"ext_oneapi_fixed_size_group">; -def AspectExt_oneapi_opportunistic_group : Aspect<"ext_oneapi_opportunistic_group">; -def AspectExt_oneapi_tangle_group : Aspect<"ext_oneapi_tangle_group">; -def AspectExt_intel_matrix : Aspect<"ext_intel_matrix">; -def AspectExt_oneapi_is_composite : Aspect<"ext_oneapi_is_composite">; -def AspectExt_oneapi_is_component : Aspect<"ext_oneapi_is_component">; -def AspectExt_oneapi_graph : Aspect<"ext_oneapi_graph">; -def AspectExt_intel_fpga_task_sequence : Aspect<"ext_intel_fpga_task_sequence">; -def AspectExt_oneapi_limited_graph : Aspect<"ext_oneapi_limited_graph">; -def AspectExt_oneapi_private_alloca : Aspect<"ext_oneapi_private_alloca">; +def AspectCpu : Aspect<"cpu", 1>; +def AspectGpu : Aspect<"gpu", 2>; +def AspectAccelerator : Aspect<"accelerator", 3>; +def AspectCustom : Aspect<"custom", 4>; +def AspectFp16 : Aspect<"fp16", 5>; +def AspectFp64 : Aspect<"fp64", 6>; +def AspectImage : Aspect<"image", 9>; +def AspectOnline_compiler : Aspect<"online_compiler", 10>; +def AspectOnline_linker : Aspect<"online_linker", 11>; +def AspectQueue_profiling : Aspect<"queue_profiling", 12>; +def AspectUsm_device_allocations : Aspect<"usm_device_allocations", 13>; +def AspectUsm_host_allocations : Aspect<"usm_host_allocations", 14>; +def AspectUsm_shared_allocations : Aspect<"usm_shared_allocations", 15>; +def AspectUsm_system_allocations : Aspect<"usm_system_allocations", 17>; +def AspectExt_intel_pci_address : Aspect<"ext_intel_pci_address", 18>; +def AspectExt_intel_gpu_eu_count : Aspect<"ext_intel_gpu_eu_count", 19>; +def AspectExt_intel_gpu_eu_simd_width : Aspect<"ext_intel_gpu_eu_simd_width", 20>; +def AspectExt_intel_gpu_slices : Aspect<"ext_intel_gpu_slices", 21>; +def AspectExt_intel_gpu_subslices_per_slice : Aspect<"ext_intel_gpu_subslices_per_slice", 22>; +def AspectExt_intel_gpu_eu_count_per_subslice : Aspect<"ext_intel_gpu_eu_count_per_subslice", 23>; +def AspectExt_intel_max_mem_bandwidth : Aspect<"ext_intel_max_mem_bandwidth", 24>; +def AspectExt_intel_mem_channel : Aspect<"ext_intel_mem_channel", 25>; +def AspectUsm_atomic_host_allocations : Aspect<"usm_atomic_host_allocations", 26>; +def AspectUsm_atomic_shared_allocations : Aspect<"usm_atomic_shared_allocations", 27>; +def AspectAtomic64 : Aspect<"atomic64", 28>; +def AspectExt_intel_device_info_uuid : Aspect<"ext_intel_device_info_uuid", 29>; +def AspectExt_oneapi_srgb : Aspect<"ext_oneapi_srgb", 30>; +def AspectExt_oneapi_native_assert : Aspect<"ext_oneapi_native_assert", 31>; +def AspectHost_debuggable : Aspect<"host_debuggable", 32>; +def AspectExt_intel_gpu_hw_threads_per_eu : Aspect<"ext_intel_gpu_hw_threads_per_eu", 33>; +def AspectExt_oneapi_cuda_async_barrier : Aspect<"ext_oneapi_cuda_async_barrier", 34>; +def AspectExt_oneapi_bfloat16_math_functions : Aspect<"ext_oneapi_bfloat16_math_functions", 35>; +def AspectExt_intel_free_memory : Aspect<"ext_intel_free_memory", 36>; +def AspectExt_intel_device_id : Aspect<"ext_intel_device_id", 37>; +def AspectExt_intel_memory_clock_rate : Aspect<"ext_intel_memory_clock_rate", 38>; +def AspectExt_intel_memory_bus_width : Aspect<"ext_intel_memory_bus_width", 39>; +def AspectEmulated : Aspect<"emulated", 40>; +def AspectExt_intel_legacy_image : Aspect<"ext_intel_legacy_image", 41>; +def AspectExt_oneapi_bindless_images : Aspect<"ext_oneapi_bindless_images", 42>; +def AspectExt_oneapi_bindless_images_shared_usm : Aspect<"ext_oneapi_bindless_images_shared_usm", 43>; +def AspectExt_oneapi_bindless_images_1d_usm : Aspect<"ext_oneapi_bindless_images_1d_usm", 44>; +def AspectExt_oneapi_bindless_images_2d_usm : Aspect<"ext_oneapi_bindless_images_2d_usm", 45>; +def AspectExt_oneapi_interop_memory_import : Aspect<"ext_oneapi_interop_memory_import", 46>; +def AspectExt_oneapi_interop_memory_export : Aspect<"ext_oneapi_interop_memory_export", 47>; +def AspectExt_oneapi_interop_semaphore_import : Aspect<"ext_oneapi_interop_semaphore_import", 48>; +def AspectExt_oneapi_interop_semaphore_export : Aspect<"ext_oneapi_interop_semaphore_export", 49>; +def AspectExt_oneapi_mipmap : Aspect<"ext_oneapi_mipmap", 50>; +def AspectExt_oneapi_mipmap_anisotropy : Aspect<"ext_oneapi_mipmap_anisotropy", 51>; +def AspectExt_oneapi_mipmap_level_reference : Aspect<"ext_oneapi_mipmap_level_reference", 52>; +def AspectExt_intel_esimd : Aspect<"ext_intel_esimd", 53>; +def AspectExt_oneapi_ballot_group : Aspect<"ext_oneapi_ballot_group", 54>; +def AspectExt_oneapi_fixed_size_group : Aspect<"ext_oneapi_fixed_size_group", 55>; +def AspectExt_oneapi_opportunistic_group : Aspect<"ext_oneapi_opportunistic_group", 56>; +def AspectExt_oneapi_tangle_group : Aspect<"ext_oneapi_tangle_group", 57>; +def AspectExt_intel_matrix : Aspect<"ext_intel_matrix", 58>; +def AspectExt_oneapi_is_composite : Aspect<"ext_oneapi_is_composite", 59>; +def AspectExt_oneapi_is_component : Aspect<"ext_oneapi_is_component", 60>; +def AspectExt_oneapi_graph : Aspect<"ext_oneapi_graph", 61>; +def AspectExt_intel_fpga_task_sequence : Aspect<"ext_intel_fpga_task_sequence", 62>; +def AspectExt_oneapi_limited_graph : Aspect<"ext_oneapi_limited_graph", 63>; +def AspectExt_oneapi_private_alloca : Aspect<"ext_oneapi_private_alloca", 64>; // Deprecated aspects -def AspectInt64_base_atomics : Aspect<"int64_base_atomics">; -def AspectInt64_extended_atomics : Aspect<"int64_extended_atomics">; -def AspectUsm_system_allocator : Aspect<"usm_system_allocator">; -def AspectUsm_restricted_shared_allocations : Aspect<"usm_restricted_shared_allocations">; -def AspectHost : Aspect<"host">; +def AspectInt64_base_atomics : Aspect<"int64_base_atomics", 7>; +def AspectInt64_extended_atomics : Aspect<"int64_extended_atomics", 8>; +def AspectUsm_system_allocator : Aspect<"usm_system_allocator", 17>; +def AspectUsm_restricted_shared_allocations : Aspect<"usm_restricted_shared_allocations", 16>; +def AspectHost : Aspect<"host", 0>; defvar AllUSMAspects = [AspectUsm_device_allocations, AspectUsm_host_allocations, AspectUsm_shared_allocations, AspectUsm_system_allocations, AspectUsm_atomic_host_allocations, AspectUsm_atomic_shared_allocations]; +def AspectTable : DynamicTable { + let FilterClass = "Aspect"; + let Fields = ["Name", "Value"]; + string CppTypeName = "int"; +} def TargetTable : DynamicTable { let FilterClass = "TargetInfo"; @@ -105,35 +111,6 @@ class TargetInfo aspectList, list subGroupS string aotToolchainOptions = options; } -// This definition serves the only purpose of testing whether the aspect list defined in here and in SYCL RT match. -def : TargetInfo<"__TestAspectList", - [AspectCpu, AspectGpu, AspectAccelerator, AspectCustom, AspectFp16, AspectFp64, AspectImage, AspectOnline_compiler, - AspectOnline_linker, AspectQueue_profiling, AspectUsm_device_allocations, AspectUsm_host_allocations, - AspectUsm_shared_allocations, AspectUsm_system_allocations, AspectExt_intel_pci_address, - AspectExt_intel_gpu_eu_count, AspectExt_intel_gpu_eu_simd_width, AspectExt_intel_gpu_slices, - AspectExt_intel_gpu_subslices_per_slice, AspectExt_intel_gpu_eu_count_per_subslice, - AspectExt_intel_max_mem_bandwidth, AspectExt_intel_mem_channel, AspectUsm_atomic_host_allocations, - AspectUsm_atomic_shared_allocations, AspectAtomic64, AspectExt_intel_device_info_uuid, AspectExt_oneapi_srgb, - AspectExt_oneapi_native_assert, AspectHost_debuggable, AspectExt_intel_gpu_hw_threads_per_eu, - AspectExt_oneapi_cuda_async_barrier, AspectExt_oneapi_bfloat16_math_functions, AspectExt_intel_free_memory, - AspectExt_intel_device_id, AspectExt_intel_memory_clock_rate, AspectExt_intel_memory_bus_width, AspectEmulated, - AspectExt_intel_legacy_image, AspectExt_oneapi_bindless_images, - AspectExt_oneapi_bindless_images_shared_usm, AspectExt_oneapi_bindless_images_1d_usm, AspectExt_oneapi_bindless_images_2d_usm, - AspectExt_oneapi_interop_memory_import, AspectExt_oneapi_interop_memory_export, - AspectExt_oneapi_interop_semaphore_import, AspectExt_oneapi_interop_semaphore_export, - AspectExt_oneapi_mipmap, AspectExt_oneapi_mipmap_anisotropy, AspectExt_oneapi_mipmap_level_reference, AspectExt_intel_esimd, - AspectExt_oneapi_ballot_group, AspectExt_oneapi_fixed_size_group, AspectExt_oneapi_opportunistic_group, - AspectExt_oneapi_tangle_group, AspectExt_intel_matrix, AspectExt_oneapi_is_composite, AspectExt_oneapi_is_component, - AspectExt_oneapi_graph, AspectExt_intel_fpga_task_sequence, AspectExt_oneapi_limited_graph, - AspectExt_oneapi_private_alloca], - []>; -// This definition serves the only purpose of testing whether the deprecated aspect list defined in here and in SYCL RT -// match. -def : TargetInfo<"__TestDeprecatedAspectList", - [AspectInt64_base_atomics, AspectInt64_extended_atomics, AspectUsm_system_allocator, - AspectUsm_restricted_shared_allocations, AspectHost], - []>; - def : TargetInfo<"spir64", [], [], "", "", 1>; def : TargetInfo<"spir64_gen", [], [], "", "", 1>; def : TargetInfo<"spir64_x86_64", [], [], "", "", 1>; diff --git a/llvm/lib/SYCLLowerIR/CleanupSYCLMetadata.cpp b/llvm/lib/SYCLLowerIR/CleanupSYCLMetadata.cpp index 799dd9188fbcf..5b0d4bc8fd487 100644 --- a/llvm/lib/SYCLLowerIR/CleanupSYCLMetadata.cpp +++ b/llvm/lib/SYCLLowerIR/CleanupSYCLMetadata.cpp @@ -36,8 +36,8 @@ PreservedAnalyses CleanupSYCLMetadataPass::run(Module &M, // Remove SYCL module-level metadata that will never be used again to avoid // duplication of their operands during llvm-link hence preventing // increase of the module size - llvm::SmallVector ModuleMDToRemove = { - "sycl_aspects", "sycl_types_that_use_aspects"}; + llvm::SmallVector ModuleMDToRemove = { + "sycl_types_that_use_aspects"}; for (const auto &MD : ModuleMDToRemove) cleanupSYCLCompilerMetadata(M, MD); diff --git a/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp b/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp index 87d3ede83ba18..31573f3685f12 100644 --- a/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp +++ b/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp @@ -31,6 +31,7 @@ #include "llvm/ADT/SetVector.h" #include "llvm/ADT/SmallPtrSet.h" #include "llvm/ADT/SmallSet.h" +#include "llvm/ADT/StringExtras.h" #include "llvm/IR/Constants.h" #include "llvm/IR/DiagnosticInfo.h" #include "llvm/IR/InstIterator.h" @@ -38,6 +39,7 @@ #include "llvm/IR/IntrinsicInst.h" #include "llvm/IR/Module.h" #include "llvm/Pass.h" +#include "llvm/SYCLLowerIR/DeviceConfigFile.hpp" #include "llvm/Support/CommandLine.h" #include "llvm/Support/Path.h" @@ -91,33 +93,7 @@ TypeToAspectsMapTy getTypesThatUseAspectsFromMetadata(const Module &M) { return Result; } -using AspectValueToNameMapTy = SmallMapVector; - -/// Retrieves from metadata (sycl_aspects) the mapping between SYCL aspect names -/// and their integral values. -AspectValueToNameMapTy getAspectsFromMetadata(const Module &M) { - const NamedMDNode *Node = M.getNamedMetadata("sycl_aspects"); - AspectValueToNameMapTy Result; - if (!Node) - return Result; - - for (const MDNode *N : Node->operands()) { - assert(N->getNumOperands() == 2 && - "Each operand of sycl_aspects must be a pair."); - - // The aspect's name is the first operand. - const auto *AspectName = cast(N->getOperand(0)); - - // The aspect's integral value is the second operand. - const auto *AspectCAM = cast(N->getOperand(1)); - const Constant *AspectC = AspectCAM->getValue(); - - Result[AspectName->getString()] = - cast(AspectC)->getSExtValue(); - } - - return Result; -} +using AspectValueToNameMapTy = DeviceConfigFile::AspectTable_t; using TypesEdgesTy = std::unordered_map>; @@ -163,8 +139,8 @@ void propagateAspectsThroughTypes(const TypesEdgesTy &Edges, const Type *Start, /// Time complexity: O((V + E) * T) where T is the number of input types /// containing aspects. void propagateAspectsToOtherTypesInModule( - const Module &M, TypeToAspectsMapTy &TypesWithAspects, - AspectValueToNameMapTy &AspectValues) { + const Module &M, TypeToAspectsMapTy &TypesWithAspects) { + const AspectValueToNameMapTy &AspectValues = DeviceConfigFile::AspectTable; std::unordered_set TypesToProcess; const Type *DoubleTy = Type::getDoubleTy(M.getContext()); @@ -365,9 +341,9 @@ void createUsedAspectsMetadataForFunctions( /// that function's sycl_declared_aspects metadata if present. A warning /// diagnostic is produced for each aspect this check fails for. void validateUsedAspectsForFunctions(const FunctionToAspectsMapTy &Map, - const AspectValueToNameMapTy &AspectValues, const std::vector &EntryPoints, const CallGraphTy &CG) { + const AspectValueToNameMapTy &AspectValues = DeviceConfigFile::AspectTable; for (auto &It : Map) { const AspectsSetTy &Aspects = It.second; if (Aspects.empty()) @@ -529,8 +505,8 @@ bool isEntryPoint(const Function &F) { } void setSyclFixedTargetsMD(const std::vector &EntryPoints, - const SmallVector &Targets, - AspectValueToNameMapTy &AspectValues) { + const SmallVector &Targets) { + const AspectValueToNameMapTy &AspectValues = DeviceConfigFile::AspectTable; if (EntryPoints.empty()) return; @@ -539,7 +515,7 @@ void setSyclFixedTargetsMD(const std::vector &EntryPoints, for (const auto &Target : Targets) { if (!Target.empty()) { - auto AspectIt = AspectValues.find(Target); + auto AspectIt = AspectValues.find(Target.str()); if (AspectIt != AspectValues.end()) { auto ConstIntTarget = ConstantInt::getSigned(Type::getInt32Ty(C), AspectIt->second); @@ -556,7 +532,6 @@ void setSyclFixedTargetsMD(const std::vector &EntryPoints, /// Returns a map of functions with corresponding used aspects. std::pair buildFunctionsToAspectsMap(Module &M, TypeToAspectsMapTy &TypesWithAspects, - const AspectValueToNameMapTy &AspectValues, const std::vector &EntryPoints, bool ValidateAspects) { FunctionToAspectsMapTy FunctionToUsedAspects; @@ -573,8 +548,7 @@ buildFunctionsToAspectsMap(Module &M, TypeToAspectsMapTy &TypesWithAspects, propagateAspectsThroughCG(F, CG, FunctionToUsedAspects, Visited); if (ValidateAspects) - validateUsedAspectsForFunctions(FunctionToUsedAspects, AspectValues, - EntryPoints, CG); + validateUsedAspectsForFunctions(FunctionToUsedAspects, EntryPoints, CG); // The set of aspects from FunctionToDeclaredAspects should be merged to the // set of FunctionToUsedAspects after validateUsedAspectsForFunctions call to @@ -592,17 +566,7 @@ buildFunctionsToAspectsMap(Module &M, TypeToAspectsMapTy &TypesWithAspects, PreservedAnalyses SYCLPropagateAspectsUsagePass::run(Module &M, ModuleAnalysisManager &MAM) { TypeToAspectsMapTy TypesWithAspects = getTypesThatUseAspectsFromMetadata(M); - AspectValueToNameMapTy AspectValues = getAspectsFromMetadata(M); - - // If there is no metadata for aspect values the source code must not have - // included the SYCL headers. In that case there should also not be any types - // that use aspects, so we can skip this pass. - if (AspectValues.empty()) { - assert(TypesWithAspects.empty() && - "sycl_aspects metadata is missing but " - "sycl_types_that_use_aspects is present."); - return PreservedAnalyses::all(); - } + const AspectValueToNameMapTy &AspectValues = DeviceConfigFile::AspectTable; if (ClSyclFixedTargets.getNumOccurrences() > 0) StringRef(ClSyclFixedTargets) @@ -621,25 +585,29 @@ SYCLPropagateAspectsUsagePass::run(Module &M, ModuleAnalysisManager &MAM) { if (isEntryPoint(F)) EntryPoints.push_back(&F); - propagateAspectsToOtherTypesInModule(M, TypesWithAspects, AspectValues); + propagateAspectsToOtherTypesInModule(M, TypesWithAspects); auto [FunctionToUsedAspects, FunctionToDeclaredAspects] = - buildFunctionsToAspectsMap(M, TypesWithAspects, AspectValues, EntryPoints, + buildFunctionsToAspectsMap(M, TypesWithAspects, EntryPoints, ValidateAspectUsage); // Create a set of excluded aspect values. AspectsSetTy ExcludedAspectVals; for (const StringRef &AspectName : ExcludedAspects) { - const auto AspectValIter = AspectValues.find(AspectName); - assert(AspectValIter != AspectValues.end() && - "Excluded aspect does not have a corresponding value."); - ExcludedAspectVals.insert(AspectValIter->second); + const auto AspectValIter = AspectValues.find(AspectName.str()); + if (AspectValIter != AspectValues.end()) + ExcludedAspectVals.insert(AspectValIter->second); + else { + int n; + assert(to_integer(AspectName, n, 10) && "Unrecognized excluded aspect!"); + ExcludedAspectVals.insert(n); + } } createUsedAspectsMetadataForFunctions( FunctionToUsedAspects, FunctionToDeclaredAspects, ExcludedAspectVals); - setSyclFixedTargetsMD(EntryPoints, TargetFixedAspects, AspectValues); + setSyclFixedTargetsMD(EntryPoints, TargetFixedAspects); return PreservedAnalyses::all(); } diff --git a/llvm/test/SYCLLowerIR/CleanupSYCLCompilerInternalMetadata/aspects-md-in-two-nodes.ll b/llvm/test/SYCLLowerIR/CleanupSYCLCompilerInternalMetadata/aspects-md-in-two-nodes.ll index 9f23f2b451c87..5caed8e8e0d6e 100644 --- a/llvm/test/SYCLLowerIR/CleanupSYCLCompilerInternalMetadata/aspects-md-in-two-nodes.ll +++ b/llvm/test/SYCLLowerIR/CleanupSYCLCompilerInternalMetadata/aspects-md-in-two-nodes.ll @@ -13,11 +13,9 @@ define spir_kernel void @kernel() !artificial !0 { } ; CHECK-NOT: sycl_types_that_use_aspects -; CHECK-NOT: sycl_aspects ; CHECK: !0 = !{!"A", i32 0} !sycl_types_that_use_aspects = !{!0} -!sycl_aspects = !{!1} !0 = !{!"A", i32 0} !1 = !{!"fp64", i32 6} diff --git a/llvm/test/SYCLLowerIR/CleanupSYCLCompilerInternalMetadata/multiple-aspects.ll b/llvm/test/SYCLLowerIR/CleanupSYCLCompilerInternalMetadata/multiple-aspects.ll index 576b76c98040c..e391727bc61c8 100644 --- a/llvm/test/SYCLLowerIR/CleanupSYCLCompilerInternalMetadata/multiple-aspects.ll +++ b/llvm/test/SYCLLowerIR/CleanupSYCLCompilerInternalMetadata/multiple-aspects.ll @@ -40,9 +40,7 @@ define spir_kernel void @kernel() !sycl_used_aspects !8 !sycl_fixed_targets !9 { } ; CHECK-NOT: sycl_types_that_use_aspects -; CHECK-NOT: sycl_aspects !sycl_types_that_use_aspects = !{!0, !1, !2, !3} -!sycl_aspects = !{!4} !0 = !{!"A", i32 0} !1 = !{!"B", i32 1} diff --git a/llvm/test/SYCLLowerIR/PropagateAspectsUsage/exclude-aspect.ll b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/exclude-aspect.ll index 223ef0e803287..0c889de0d1d90 100644 --- a/llvm/test/SYCLLowerIR/PropagateAspectsUsage/exclude-aspect.ll +++ b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/exclude-aspect.ll @@ -1,4 +1,4 @@ -; RUN: opt -passes=sycl-propagate-aspects-usage -sycl-propagate-aspects-usage-exclude-aspects=aspect4,aspect1 -S < %s | FileCheck %s +; RUN: opt -passes=sycl-propagate-aspects-usage -sycl-propagate-aspects-usage-exclude-aspects=4,1 -S < %s | FileCheck %s ; ; Test checks that the pass is able to collect all aspects used in a function diff --git a/llvm/utils/TableGen/DynamicTableEmitter.cpp b/llvm/utils/TableGen/DynamicTableEmitter.cpp index 35c29eff0d33a..bae17dfa91982 100644 --- a/llvm/utils/TableGen/DynamicTableEmitter.cpp +++ b/llvm/utils/TableGen/DynamicTableEmitter.cpp @@ -66,6 +66,8 @@ class DynamicTableEmitter { return SI->getAsString(); } + if (IntInit *II = dyn_cast(I)) + return Twine(getAsInt(II)).str(); if (BitsInit *BI = dyn_cast(I)) return "0x" + utohexstr(getAsInt(BI)); if (BitInit *BI = dyn_cast(I)) @@ -121,8 +123,8 @@ void DynamicTableEmitter::emitDynamicTable(const DynamicTable &Table, emitIfdef((Twine("GET_") + Table.PreprocessorGuard + "_IMPL"), OS); // The primary data table contains all the fields defined for this map. - OS << "std::map " << Table.Name - << " = {\n"; + OS << "static std::map " + << Table.Name << " = {\n"; // Iterate over the key-value pairs the dynamic table will contain. for (unsigned I = 0; I < Table.Entries.size(); ++I) { Record *Entry = Table.Entries[I]; diff --git a/sycl/test/basic_tests/device_config_file_aspects.cpp b/sycl/test/basic_tests/device_config_file_aspects.cpp index da08b87e60a2b..62911bf90e421 100644 --- a/sycl/test/basic_tests/device_config_file_aspects.cpp +++ b/sycl/test/basic_tests/device_config_file_aspects.cpp @@ -1,41 +1,48 @@ +// This test checks if DeviceConfigFile.td and aspects.def are in sync. // RUN: %clangxx -fsycl %s -o %t.out -I %llvm_main_include_dir // RUN: %t.out -// + +#include #include #include #include #include -#define __SYCL_ASPECT_DEPRECATED_ALIAS(ASPECT, ID, MESSAGE) \ - __SYCL_ASPECT_DEPRECATED(ASPECT, ID, MESSAGE) +void check(const char *aspect_name, int aspect_val, int &n_fail) { + const auto &aspectTable = DeviceConfigFile::AspectTable; + auto res = aspectTable.find(aspect_name); + if (res == aspectTable.end()) { + std::cout << "Aspect " << aspect_name + << " was not found in the device config file!\n"; + ++n_fail; + return; + } + if (res->second != aspect_val) { + std::cout << "Aspect " << aspect_name << " has value " << res->second + << " in the device config file but has value " << aspect_val + << " in aspects.def!\n"; + ++n_fail; + return; + } +} int main() { - auto testAspects = DeviceConfigFile::TargetTable.find("__TestAspectList"); - assert(testAspects != DeviceConfigFile::TargetTable.end()); - auto aspectsList = testAspects->second.aspects; - -#define __SYCL_ASPECT(ASPECT, ASPECT_VAL) \ - llvm::StringRef s##ASPECT(#ASPECT); \ - assert(std::find(aspectsList.begin(), aspectsList.end(), s##ASPECT) != \ - aspectsList.end()); - -#include - -#undef __SYCL_ASPECT + int n_fail = 0; - auto testDeprecatedAspects = - DeviceConfigFile::TargetTable.find("__TestDeprecatedAspectList"); - assert(testDeprecatedAspects != DeviceConfigFile::TargetTable.end()); - auto deprecatedAspectsList = testDeprecatedAspects->second.aspects; +#define __SYCL_ASPECT(ASPECT, ASPECT_VAL) check(#ASPECT, ASPECT_VAL, n_fail); #define __SYCL_ASPECT_DEPRECATED(ASPECT, ASPECT_VAL, MSG) \ - llvm::StringRef s##ASPECT(#ASPECT); \ - assert(std::find(deprecatedAspectsList.begin(), deprecatedAspectsList.end(), \ - s##ASPECT) != deprecatedAspectsList.end()); - + __SYCL_ASPECT(ASPECT, ASPECT_VAL) +#define __SYCL_ASPECT_DEPRECATED_ALIAS(ASPECT, ID, MESSAGE) +#include #include - +#undef __SYCL_ASPECT_DEPRECATED_ALIAS #undef __SYCL_ASPECT_DEPRECATED -} +#undef __SYCL_ASPECT -#undef __SYCL_ASPECT_DEPRECATED_ALIAS + if (n_fail > 0) { + std::cout << "Errors detected, DeviceConfigFile.td and aspects.def are out " + "of sync!\n"; + } + return n_fail; +} \ No newline at end of file diff --git a/sycl/test/extensions/properties/properties_kernel_device_has.cpp b/sycl/test/extensions/properties/properties_kernel_device_has.cpp index aa1988837d8c4..4fe0355b2f3ae 100644 --- a/sycl/test/extensions/properties/properties_kernel_device_has.cpp +++ b/sycl/test/extensions/properties/properties_kernel_device_has.cpp @@ -130,40 +130,5 @@ int main() { return 0; } -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_oneapi_cuda_async_barrier", i32 [[ext_oneapi_cuda_async_barrier_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_oneapi_bfloat16_math_functions", i32 [[ext_oneapi_bfloat16_math_functions_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"custom", i32 [[custom_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"fp16", i32 [[fp16_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"fp64", i32 [[fp64_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"image", i32 [[image_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"online_compiler", i32 [[online_compiler_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"online_linker", i32 [[online_linker_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"queue_profiling", i32 [[queue_profiling_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"usm_device_allocations", i32 [[usm_device_allocations_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"usm_system_allocations", i32 [[usm_system_allocations_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_pci_address", i32 [[ext_intel_pci_address_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"cpu", i32 [[cpu_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"gpu", i32 [[gpu_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"accelerator", i32 [[accelerator_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_gpu_eu_count", i32 [[ext_intel_gpu_eu_count_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_gpu_subslices_per_slice", i32 [[ext_intel_gpu_subslices_per_slice_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_gpu_eu_count_per_subslice", i32 [[ext_intel_gpu_eu_count_per_subslice_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_max_mem_bandwidth", i32 [[ext_intel_max_mem_bandwidth_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_mem_channel", i32 [[ext_intel_mem_channel_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"usm_atomic_host_allocations", i32 [[usm_atomic_host_allocations_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"usm_atomic_shared_allocations", i32 [[usm_atomic_shared_allocations_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"atomic64", i32 [[atomic64_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_device_info_uuid", i32 [[ext_intel_device_info_uuid_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_oneapi_srgb", i32 [[ext_oneapi_srgb_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_gpu_eu_simd_width", i32 [[ext_intel_gpu_eu_simd_width_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_gpu_slices", i32 [[ext_intel_gpu_slices_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_oneapi_native_assert", i32 [[ext_oneapi_native_assert_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"host_debuggable", i32 [[host_debuggable_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_gpu_hw_threads_per_eu", i32 [[ext_intel_gpu_hw_threads_per_eu_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"usm_host_allocations", i32 [[usm_host_allocations_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"usm_shared_allocations", i32 [[usm_shared_allocations_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_free_memory", i32 [[ext_intel_free_memory_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_device_id", i32 [[ext_intel_device_id_ASPECT_MD:[0-9]+]]} - -// CHECK-IR-DAG: attributes #[[DHAttr1]] = { {{.*}}"sycl-device-has"="[[ext_oneapi_cuda_async_barrier_ASPECT_MD]],[[ext_oneapi_bfloat16_math_functions_ASPECT_MD]],[[custom_ASPECT_MD]],[[fp16_ASPECT_MD]],[[fp64_ASPECT_MD]],[[image_ASPECT_MD]],[[online_compiler_ASPECT_MD]],[[online_linker_ASPECT_MD]],[[queue_profiling_ASPECT_MD]],[[usm_device_allocations_ASPECT_MD]],[[usm_system_allocations_ASPECT_MD]],[[ext_intel_pci_address_ASPECT_MD]],[[cpu_ASPECT_MD]],[[gpu_ASPECT_MD]],[[accelerator_ASPECT_MD]],[[ext_intel_gpu_eu_count_ASPECT_MD]],[[ext_intel_gpu_subslices_per_slice_ASPECT_MD]],[[ext_intel_gpu_eu_count_per_subslice_ASPECT_MD]],[[ext_intel_max_mem_bandwidth_ASPECT_MD]],[[ext_intel_mem_channel_ASPECT_MD]],[[usm_atomic_host_allocations_ASPECT_MD]],[[usm_atomic_shared_allocations_ASPECT_MD]],[[atomic64_ASPECT_MD]],[[ext_intel_device_info_uuid_ASPECT_MD]],[[ext_oneapi_srgb_ASPECT_MD]],[[ext_intel_gpu_eu_simd_width_ASPECT_MD]],[[ext_intel_gpu_slices_ASPECT_MD]],[[ext_oneapi_native_assert_ASPECT_MD]],[[host_debuggable_ASPECT_MD]],[[ext_intel_gpu_hw_threads_per_eu_ASPECT_MD]],[[usm_host_allocations_ASPECT_MD]],[[usm_shared_allocations_ASPECT_MD]],[[ext_intel_free_memory_ASPECT_MD]],[[ext_intel_device_id_ASPECT_MD]]" -// CHECK-IR-DAG: attributes #[[DHAttr2]] = { {{.*}}"sycl-device-has"="[[ext_oneapi_cuda_async_barrier_ASPECT_MD]],[[ext_oneapi_bfloat16_math_functions_ASPECT_MD]],[[custom_ASPECT_MD]],[[fp16_ASPECT_MD]],[[fp64_ASPECT_MD]],[[image_ASPECT_MD]],[[online_compiler_ASPECT_MD]],[[online_linker_ASPECT_MD]],[[queue_profiling_ASPECT_MD]],[[usm_device_allocations_ASPECT_MD]],[[usm_system_allocations_ASPECT_MD]],[[ext_intel_pci_address_ASPECT_MD]],[[cpu_ASPECT_MD]],[[gpu_ASPECT_MD]],[[accelerator_ASPECT_MD]],[[ext_intel_gpu_eu_count_ASPECT_MD]],[[ext_intel_gpu_subslices_per_slice_ASPECT_MD]],[[ext_intel_gpu_eu_count_per_subslice_ASPECT_MD]],[[ext_intel_max_mem_bandwidth_ASPECT_MD]],[[ext_intel_mem_channel_ASPECT_MD]],[[usm_atomic_host_allocations_ASPECT_MD]],[[usm_atomic_shared_allocations_ASPECT_MD]],[[atomic64_ASPECT_MD]],[[ext_intel_device_info_uuid_ASPECT_MD]],[[ext_oneapi_srgb_ASPECT_MD]],[[ext_intel_gpu_eu_simd_width_ASPECT_MD]],[[ext_intel_gpu_slices_ASPECT_MD]],[[ext_oneapi_native_assert_ASPECT_MD]],[[host_debuggable_ASPECT_MD]],[[ext_intel_gpu_hw_threads_per_eu_ASPECT_MD]],[[usm_host_allocations_ASPECT_MD]],[[usm_shared_allocations_ASPECT_MD]],[[ext_intel_free_memory_ASPECT_MD]],[[ext_intel_device_id_ASPECT_MD]]" +// CHECK-IR-DAG: attributes #[[DHAttr1]] = { {{.*}}"sycl-device-has"="34,35,4,5,6,9,10,11,12,13,17,18,1,2,3,19,22,23,24,25,26,27,28,29,30,20,21,31,32,33,14,15,36,37" +// CHECK-IR-DAG: attributes #[[DHAttr2]] = { {{.*}}"sycl-device-has"="34,35,4,5,6,9,10,11,12,13,17,18,1,2,3,19,22,23,24,25,26,27,28,29,30,20,21,31,32,33,14,15,36,37" diff --git a/sycl/test/extensions/properties/properties_kernel_device_has_macro.cpp b/sycl/test/extensions/properties/properties_kernel_device_has_macro.cpp index 3664e8f794ab4..d70644fdd5700 100644 --- a/sycl/test/extensions/properties/properties_kernel_device_has_macro.cpp +++ b/sycl/test/extensions/properties/properties_kernel_device_has_macro.cpp @@ -46,43 +46,6 @@ int main() { return 0; } -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_oneapi_cuda_async_barrier", i32 [[ext_oneapi_cuda_async_barrier_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_oneapi_bfloat16_math_functions", i32 [[ext_oneapi_bfloat16_math_functions_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"custom", i32 [[custom_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"fp16", i32 [[fp16_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"fp64", i32 [[fp64_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"image", i32 [[image_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"online_compiler", i32 [[online_compiler_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"online_linker", i32 [[online_linker_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"queue_profiling", i32 [[queue_profiling_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"usm_device_allocations", i32 [[usm_device_allocations_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"usm_system_allocations", i32 [[usm_system_allocations_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_pci_address", i32 [[ext_intel_pci_address_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"cpu", i32 [[cpu_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"gpu", i32 [[gpu_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"accelerator", i32 [[accelerator_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_gpu_eu_count", i32 [[ext_intel_gpu_eu_count_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_gpu_subslices_per_slice", i32 [[ext_intel_gpu_subslices_per_slice_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_gpu_eu_count_per_subslice", i32 [[ext_intel_gpu_eu_count_per_subslice_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_max_mem_bandwidth", i32 [[ext_intel_max_mem_bandwidth_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_mem_channel", i32 [[ext_intel_mem_channel_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"usm_atomic_host_allocations", i32 [[usm_atomic_host_allocations_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"usm_atomic_shared_allocations", i32 [[usm_atomic_shared_allocations_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"atomic64", i32 [[atomic64_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_device_info_uuid", i32 [[ext_intel_device_info_uuid_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_oneapi_srgb", i32 [[ext_oneapi_srgb_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_gpu_eu_simd_width", i32 [[ext_intel_gpu_eu_simd_width_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_gpu_slices", i32 [[ext_intel_gpu_slices_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_oneapi_native_assert", i32 [[ext_oneapi_native_assert_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"host_debuggable", i32 [[host_debuggable_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_gpu_hw_threads_per_eu", i32 [[ext_intel_gpu_hw_threads_per_eu_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"usm_host_allocations", i32 [[usm_host_allocations_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"usm_shared_allocations", i32 [[usm_shared_allocations_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_free_memory", i32 [[ext_intel_free_memory_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_device_id", i32 [[ext_intel_device_id_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_memory_clock_rate", i32 [[ext_intel_memory_clock_rate_ASPECT_MD:[0-9]+]]} -// CHECK-IR-DAG: !{{[0-9]+}} = !{!"ext_intel_memory_bus_width", i32 [[ext_intel_memory_bus_width_ASPECT_MD:[0-9]+]]} - -// CHECK-IR-DAG: attributes #[[DHAttr1]] = { {{.*}}"sycl-device-has"="[[ext_oneapi_cuda_async_barrier_ASPECT_MD]],[[ext_oneapi_bfloat16_math_functions_ASPECT_MD]],[[custom_ASPECT_MD]],[[fp16_ASPECT_MD]],[[fp64_ASPECT_MD]],[[image_ASPECT_MD]],[[online_compiler_ASPECT_MD]],[[online_linker_ASPECT_MD]],[[queue_profiling_ASPECT_MD]],[[usm_device_allocations_ASPECT_MD]],[[usm_system_allocations_ASPECT_MD]],[[ext_intel_pci_address_ASPECT_MD]],[[cpu_ASPECT_MD]],[[gpu_ASPECT_MD]],[[accelerator_ASPECT_MD]],[[ext_intel_gpu_eu_count_ASPECT_MD]],[[ext_intel_gpu_subslices_per_slice_ASPECT_MD]],[[ext_intel_gpu_eu_count_per_subslice_ASPECT_MD]],[[ext_intel_max_mem_bandwidth_ASPECT_MD]],[[ext_intel_mem_channel_ASPECT_MD]],[[usm_atomic_host_allocations_ASPECT_MD]],[[usm_atomic_shared_allocations_ASPECT_MD]],[[atomic64_ASPECT_MD]],[[ext_intel_device_info_uuid_ASPECT_MD]],[[ext_oneapi_srgb_ASPECT_MD]],[[ext_intel_gpu_eu_simd_width_ASPECT_MD]],[[ext_intel_gpu_slices_ASPECT_MD]],[[ext_oneapi_native_assert_ASPECT_MD]],[[host_debuggable_ASPECT_MD]],[[ext_intel_gpu_hw_threads_per_eu_ASPECT_MD]],[[usm_host_allocations_ASPECT_MD]],[[usm_shared_allocations_ASPECT_MD]],[[ext_intel_free_memory_ASPECT_MD]],[[ext_intel_device_id_ASPECT_MD]],[[ext_intel_memory_clock_rate_ASPECT_MD]],[[ext_intel_memory_bus_width_ASPECT_MD]]" +// CHECK-IR-DAG: attributes #[[DHAttr1]] = { {{.*}}"sycl-device-has"="34,35,4,5,6,9,10,11,12,13,17,18,1,2,3,19,22,23,24,25,26,27,28,29,30,20,21,31,32,33,14,15,36,37,38,39" // CHECK-IR-DAG: attributes #[[DHAttr2]] = { {{.*}}"sycl-device-has" {{.*}} -// CHECK-IR-DAG: attributes #[[DHAttr3]] = { {{.*}}"sycl-device-has"="[[fp16_ASPECT_MD]],[[atomic64_ASPECT_MD]]" +// CHECK-IR-DAG: attributes #[[DHAttr3]] = { {{.*}}"sycl-device-has"="5,28" diff --git a/sycl/test/optional_kernel_features/atomic_ref-atomic64-aspect.cpp b/sycl/test/optional_kernel_features/atomic_ref-atomic64-aspect.cpp index 2e11af832949f..e8340f6e7dbd7 100644 --- a/sycl/test/optional_kernel_features/atomic_ref-atomic64-aspect.cpp +++ b/sycl/test/optional_kernel_features/atomic_ref-atomic64-aspect.cpp @@ -3,9 +3,8 @@ // CHECK: !sycl_types_that_use_aspects = !{![[#MDNUM1:]], ![[#MDNUM2:]], ![[#MDNUM3:]]} // CHECK: ![[#MDNUM1]] = !{!"class.sycl::_V1::detail::atomic_ref_impl", i32 [[#ASPECT_NUM:]]} -// CHECK: ![[#MDNUM2]] = !{!"class.sycl::_V1::detail::atomic_ref_impl.2", i32 [[#ASPECT_NUM:]]} -// CHECK: ![[#MDNUM3]] = !{!"class.sycl::_V1::detail::atomic_ref_impl.7", i32 [[#ASPECT_NUM:]]} -// CHECK: !{{.*}} = !{!"atomic64", i32 [[#ASPECT_NUM]]} +// CHECK: ![[#MDNUM2]] = !{!"class.sycl::_V1::detail::atomic_ref_impl.2", i32 [[#ASPECT_NUM]]} +// CHECK: ![[#MDNUM3]] = !{!"class.sycl::_V1::detail::atomic_ref_impl.7", i32 [[#ASPECT_NUM]]} #include diff --git a/sycl/test/optional_kernel_features/esimd.cpp b/sycl/test/optional_kernel_features/esimd.cpp index fe787a86d6297..9be491a6c67a4 100644 --- a/sycl/test/optional_kernel_features/esimd.cpp +++ b/sycl/test/optional_kernel_features/esimd.cpp @@ -4,7 +4,6 @@ // CHECK: !sycl_types_that_use_aspects = !{![[#MDNUM1:]], ![[#MDNUM2:]]} // CHECK: ![[#MDNUM1]] = !{!"class.sycl::_V1::ext::intel::esimd::detail::simd_obj_impl", i32 [[#ASPECT_NUM:]]} // CHECK: ![[#MDNUM2]] = !{!"class.sycl::_V1::ext::intel::esimd::detail::simd_view_impl", i32 [[#ASPECT_NUM]]} -// CHECK: !{{.*}} = !{!"ext_intel_esimd", i32 [[#ASPECT_NUM]]} #include #include diff --git a/sycl/test/optional_kernel_features/half-aspect.cpp b/sycl/test/optional_kernel_features/half-aspect.cpp index 868d3c6a301c1..fd45c4beee5d6 100644 --- a/sycl/test/optional_kernel_features/half-aspect.cpp +++ b/sycl/test/optional_kernel_features/half-aspect.cpp @@ -3,7 +3,6 @@ // CHECK: !sycl_types_that_use_aspects = !{![[#MDNUM:]]} // CHECK: ![[#MDNUM]] = !{!"class.sycl::_V1::detail::half_impl::half", i32 [[#ASPECT_NUM:]]} -// CHECK: !{{.*}} = !{!"fp16", i32 [[#ASPECT_NUM]]} #include