[SYCL] Add device config file consistency test (#16369)

Co-authored-by: Kornev, Nikita <[email protected]>

Author: jzc

buildbot/configure.py

@@ -66,6 +66,7 @@ def do_configure(args, passthrough_args):
     xpti_enable_werror = "OFF"
     llvm_enable_zstd = "ON"
     spirv_enable_dis = "OFF"
+    sycl_install_device_config_file = "OFF"
 
     if sys.platform != "darwin":
         # For more info on the enablement of level_zero_v2 refer to this document:
@@ -160,6 +161,7 @@ def do_configure(args, passthrough_args):
                 libclc_targets_to_build += libclc_nvidia_target_names
             libclc_gen_remangled_variants = "ON"
             spirv_enable_dis = "ON"
+            sycl_install_device_config_file = "ON"
 
     if args.enable_backends:
         sycl_enabled_backends += args.enable_backends
@@ -208,6 +210,7 @@ def do_configure(args, passthrough_args):
         "-DSYCL_ENABLE_EXTENSION_JIT={}".format(sycl_enable_jit),
         "-DSYCL_ENABLE_MAJOR_RELEASE_PREVIEW_LIB={}".format(sycl_preview_lib),
         "-DBUG_REPORT_URL=https://github.com/intel/llvm/issues",
+        "-DSYCL_INSTALL_DEVICE_CONFIG_FILE={}".format(sycl_install_device_config_file),
     ]
 
     if libclc_enabled:

llvm/include/llvm/SYCLLowerIR/CMakeLists.txt

@@ -6,3 +6,7 @@ set(LLVM_TABLEGEN_PROJECT LLVM)
 set(LLVM_TARGET_DEFINITIONS DeviceConfigFile.td)
 tablegen(LLVM DeviceConfigFile.inc -gen-dynamic-tables)
 add_public_tablegen_target(DeviceConfigFile)
+install(FILES "${CMAKE_CURRENT_SOURCE_DIR}/DeviceConfigFile.hpp"
+              "${CMAKE_CURRENT_BINARY_DIR}/DeviceConfigFile.inc"
+        DESTINATION include/llvm/SYCLLowerIR
+        COMPONENT DeviceConfigFile)

llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.hpp

@@ -8,12 +8,9 @@
 
 #include <map>
 #include <string>
+#include <string_view>
 #include <vector>
 
-namespace llvm {
-class StringRef;
-}
-
 namespace DeviceConfigFile {
 
 // This struct is used in DeviceConfigFile.td. Both the fields and the name of
@@ -22,7 +19,7 @@ namespace DeviceConfigFile {
 // DeviceConfigFile.td.
 struct TargetInfo {
   bool maySupportOtherAspects;
-  std::vector<llvm::StringRef> aspects;
+  std::vector<std::string_view> aspects;
   std::vector<unsigned> subGroupSizes;
   std::string aotToolchain;
   std::string aotToolchainOptions;

llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td

@@ -177,14 +177,18 @@ defvar IntelCpuAspects = [
     AspectExt_oneapi_srgb, AspectExt_oneapi_native_assert,
     AspectExt_intel_legacy_image, AspectExt_oneapi_ballot_group,
     AspectExt_oneapi_fixed_size_group, AspectExt_oneapi_opportunistic_group,
-    AspectExt_oneapi_tangle_group, AspectExt_oneapi_private_alloca
+    AspectExt_oneapi_tangle_group, AspectExt_oneapi_private_alloca, 
+    AspectOnline_compiler, AspectOnline_linker, AspectExt_intel_gpu_slices,
+    AspectExt_intel_gpu_subslices_per_slice, AspectExt_intel_gpu_eu_count_per_subslice,
+    AspectExt_intel_gpu_hw_threads_per_eu, AspectExt_intel_device_id, 
+    AspectExt_oneapi_virtual_functions
 ] # AllUSMAspects;
 
 def : TargetInfo<"spir64", [], [], "", "", 1>;
 def : TargetInfo<"spir64_gen", [], [], "", "", 1>;
 def : TargetInfo<"spir64_x86_64", IntelCpuAspects, [4, 8, 16, 32, 64], "", "", 1>;
 def : TargetInfo<"spir64_fpga", [], [], "", "", 1>;
-def : TargetInfo<"x86_64", [], [], "", "", 1>;
+def : TargetInfo<"x86_64", IntelCpuAspects, [4, 8, 16, 32, 64], "", "", 1>;
 // Examples of how to use a combination of explicitly specified values + predefined lists
 //defvar AspectList = [AspectCpu] # AllUSMAspects;
 //def : TargetInfo<"Test", AspectList, []>;
@@ -196,9 +200,11 @@ defvar Fp16Fp64Atomic64 = [AspectFp16, AspectFp64, AspectAtomic64];
 defvar Fp16Atomic64 = [AspectFp16, AspectAtomic64];
 defvar Sg8_16_32 = [8, 16, 32];
 defvar Sg16_32 = [16, 32];
-defvar IntelBaseAspects = [AspectExt_intel_esimd];
+defvar IntelGPUBaseAspects = [AspectExt_intel_esimd, AspectExt_oneapi_ballot_group,
+    AspectExt_oneapi_fixed_size_group, AspectExt_oneapi_opportunistic_group,
+    AspectExt_oneapi_tangle_group];
 class IntelTargetInfo<string Name, list<Aspect> Aspects, list<int> subGroupSizesList>
-: TargetInfo<Name, IntelBaseAspects # Aspects, subGroupSizesList>;
+: TargetInfo<Name, IntelGPUBaseAspects # Aspects, subGroupSizesList>;
 // Note: only the "canonical" target names are listed here - see
 // SYCL::gen::resolveGenDevice().
 //
@@ -268,7 +274,7 @@ defvar CudaMinUSMAspects = [AspectUsm_device_allocations, AspectUsm_host_allocat
 defvar CudaSM90USMAspects = [AspectUsm_system_allocations, AspectUsm_atomic_host_allocations, AspectUsm_atomic_shared_allocations];
 
 defvar CudaMinAspects = !listconcat(CudaMinUSMAspects, [AspectGpu, AspectFp64, AspectOnline_compiler, AspectOnline_linker,
-    AspectQueue_profiling, AspectExt_intel_pci_address, AspectExt_intel_max_mem_bandwidth, AspectExt_intel_memory_bus_width,
+    AspectQueue_profiling, AspectExt_intel_pci_address, AspectExt_intel_memory_bus_width,
     AspectExt_intel_device_info_uuid, AspectExt_oneapi_native_assert, AspectExt_intel_free_memory, AspectExt_intel_device_id,
     AspectExt_intel_memory_clock_rate, AspectExt_oneapi_ballot_group, AspectExt_oneapi_fixed_size_group,
     AspectExt_oneapi_opportunistic_group, AspectExt_oneapi_graph, AspectExt_oneapi_limited_graph]);
@@ -325,9 +331,9 @@ defvar HipSubgroupSizesCDNA = [64];     // gfx908, gfx90a (encapsulates CDNA1..2
 
 defvar HipMinAspects = [AspectGpu, AspectFp16, AspectFp64,
     AspectOnline_compiler, AspectOnline_linker, AspectQueue_profiling,
-    AspectExt_intel_pci_address, AspectExt_intel_max_mem_bandwidth,
-    AspectExt_intel_device_id, AspectExt_intel_memory_clock_rate,
-    AspectExt_intel_memory_bus_width, AspectExt_intel_free_memory];
+    AspectExt_intel_pci_address, AspectExt_intel_device_id,
+    AspectExt_intel_memory_clock_rate, AspectExt_intel_memory_bus_width,
+    AspectExt_intel_free_memory];
 
 defvar HipUSMAspects = !listremove(AllUSMAspects, [AspectUsm_system_allocations]);
 defvar HipGraphAspects = [AspectExt_oneapi_graph, AspectExt_oneapi_limited_graph];

sycl/CMakeLists.txt

@@ -547,6 +547,11 @@ if("hip" IN_LIST SYCL_ENABLE_BACKENDS)
   list(APPEND SYCL_TOOLCHAIN_DEPLOY_COMPONENTS ur_adapter_hip)
 endif()
 
+if(SYCL_INSTALL_DEVICE_CONFIG_FILE)
+  add_dependencies(sycl-toolchain DeviceConfigFile)
+  list(APPEND SYCL_TOOLCHAIN_DEPLOY_COMPONENTS DeviceConfigFile)
+endif()
+
 # Use it as fake dependency in order to force another command(s) to execute.
 add_custom_command(OUTPUT __force_it
   COMMAND "${CMAKE_COMMAND}" -E echo

sycl/test-e2e/Basic/device_config_file_consistency.cpp

@@ -0,0 +1,136 @@
+// This test checks to see if every aspect and sub-group size declared in the
+// device config file is supported by the device. Note this does not mean
+// check that the device config file is exhaustive, only that the device
+// supports everything it declares. However, this test does print out any
+// aspects that are supported by the device but not declared in the device
+// config file.
+
+// REQUIRES: device-config-file
+// RUN: %{build} -o %t.out %device_config_file_include_flag
+// RUN: %{run} %t.out
+#include <map>
+
+#include <llvm/SYCLLowerIR/DeviceConfigFile.hpp>
+#include <sycl/detail/core.hpp>
+
+#define __SYCL_ASPECT_DEPRECATED_ALIAS(ASPECT, ID, MESSAGE)                    \
+  __SYCL_ASPECT_DEPRECATED(ASPECT, ID, MESSAGE)
+
+using namespace sycl;
+
+const char *getArchName(const device &Device) {
+  namespace syclex = sycl::ext::oneapi::experimental;
+  auto Arch = Device.get_info<syclex::info::device::architecture>();
+  switch (Arch) {
+#define __SYCL_ARCHITECTURE(ARCH, VAL)                                         \
+  case syclex::architecture::ARCH:                                             \
+    return #ARCH;
+#define __SYCL_ARCHITECTURE_ALIAS(ARCH, VAL)
+#include <sycl/ext/oneapi/experimental/device_architecture.def>
+#undef __SYCL_ARCHITECTURE
+#undef __SYCL_ARCHITECTURE_ALIAS
+  }
+  return "unknown";
+}
+
+// Checks if a container contains a specific element.
+template <typename Container, typename T>
+bool contains(const Container &C, const T &Elem) {
+  return std::find(C.begin(), C.end(), Elem) != C.end();
+}
+
+std::string_view getAspectName(aspect Asp) {
+  switch (Asp) {
+#define __SYCL_ASPECT(ASPECT, ID)                                              \
+  case aspect::ASPECT:                                                         \
+    return #ASPECT;
+#include <sycl/info/aspects.def>
+#undef __SYCL_ASPECT
+  }
+  return "unknown";
+}
+
+aspect getAspectByName(std::string_view Name) {
+#define __SYCL_ASPECT(ASPECT, ID)                                              \
+  if (Name == #ASPECT)                                                         \
+    return aspect::ASPECT;
+#include <sycl/info/aspects.def>
+  throw std::invalid_argument("Unknown aspect name");
+}
+
+int main() {
+  // Get the device arch.
+  queue Q;
+  auto Dev = Q.get_device();
+  auto DeviceName = getArchName(Dev);
+
+  auto TargetInfo = DeviceConfigFile::TargetTable.find(DeviceName);
+  if (TargetInfo == DeviceConfigFile::TargetTable.end()) {
+    std::cout << "No aspects found for device " << DeviceName << "\n";
+    return 1;
+  }
+
+  // Check aspects consistency.
+  int NAspectInconsistencies = 0;
+
+  auto SupportedAspects = Dev.get_info<info::device::aspects>();
+  auto DeviceConfigAspectNames = TargetInfo->second.aspects;
+  std::vector<aspect> DeviceConfigAspects;
+  for (auto AspectName : DeviceConfigAspectNames) {
+    DeviceConfigAspects.push_back(getAspectByName(AspectName));
+  }
+
+  for (auto Asp : DeviceConfigAspects) {
+    if (!contains(SupportedAspects, Asp)) {
+      std::cout << "error: " << DeviceName << " does not support aspect "
+                << getAspectName(Asp)
+                << " but it is declared in the device config file\n";
+      ++NAspectInconsistencies;
+    }
+  }
+  for (auto Asp : SupportedAspects) {
+    if (!contains(DeviceConfigAspects, Asp)) {
+      std::cout << "note: the device " << DeviceName << " supports aspect "
+                << getAspectName(Asp)
+                << " but it is not declared in the device config file\n";
+      // Not necessarily an error, so we won't increment n_fail.
+    }
+  }
+
+  if (NAspectInconsistencies != 0) {
+    std::cout << "Aspects are inconsistent\n";
+    return 1;
+  }
+
+  // Check sub-group sizes consistency.
+  int NSubGroupSizeInconsistencies = 0;
+
+  auto SupportedSubGroupSizes = Dev.get_info<info::device::sub_group_sizes>();
+  auto DeviceConfigSubGroupSizes = TargetInfo->second.subGroupSizes;
+
+  for (auto Size : DeviceConfigSubGroupSizes) {
+    if (!contains(SupportedSubGroupSizes, Size)) {
+      std::cout << "error: " << DeviceName
+                << " does not support sub-group size " << Size
+                << " but it is declared in the device config file\n";
+      ++NSubGroupSizeInconsistencies;
+    }
+  }
+  for (auto Size : SupportedSubGroupSizes) {
+    if (!contains(DeviceConfigSubGroupSizes, Size)) {
+      std::cout << "note: the device " << DeviceName
+                << " supports sub-group size " << Size
+                << " but it is not declared in the device config file\n";
+      // Not necessarily an error, so we won't increment n_fail.
+    }
+  }
+
+  if (NSubGroupSizeInconsistencies != 0) {
+    std::cout << "Sub-group sizes are inconsistent\n";
+    return 1;
+  }
+
+  return 0;
+}
+
+#undef __SYCL_ASPECT_DEPRECATED_ALIAS

sycl/test-e2e/E2EExpr.py

@@ -46,6 +46,7 @@ class E2EExpr(BooleanExpression):
         "false",
         "pdtracker",
         "ze_debug",
+        "device-config-file",
     }
 
     def __init__(self, string, variables, build_only_mode, final_unknown_value):

sycl/test-e2e/lit.cfg.py

@@ -932,6 +932,18 @@ def get_sycl_ls_verbose(sycl_device, env):
 
 for target in config.sycl_build_targets:
     config.available_features.add("any-target-is-" + target.replace("target-", ""))
+
+if config.llvm_main_include_dir:
+    lit_config.note("Using device config file built from LLVM")
+    config.available_features.add("device-config-file")
+    config.substitutions.append(
+        ("%device_config_file_include_flag", f"-I {config.llvm_main_include_dir}")
+    )
+elif os.path.exists(f"{config.sycl_include}/llvm/SYCLLowerIR/DeviceConfigFile.hpp"):
+    lit_config.note("Using installed device config file")
+    config.available_features.add("device-config-file")
+    config.substitutions.append(("%device_config_file_include_flag", ""))
+
 # That has to be executed last so that all device-independent features have been
 # discovered already.
 config.sycl_dev_features = {}

sycl/test-e2e/lit.site.cfg.py.in

@@ -18,6 +18,7 @@ def get_dpcpp_tool_path(name):
     except subprocess.CalledProcessError:
         return os.path.join(config.dpcpp_bin_dir, name)
 
+config.llvm_main_include_dir = "@LLVM_MAIN_INCLUDE_DIR@"
 config.llvm_tools_dir = os.path.dirname(get_dpcpp_tool_path("llvm-config"))
 config.lit_tools_dir = os.path.dirname("@TEST_SUITE_LIT@")
 config.dump_ir_supported = lit_config.params.get("dump_ir", ("@DUMP_IR_SUPPORTED@" if "@DUMP_IR_SUPPORTED@" else False))

sycl/test/basic_tests/device_config_file_aspects.cpp

@@ -3,7 +3,6 @@
 //
 #include <map>
 
-#include <llvm/ADT/StringRef.h>
 #include <llvm/SYCLLowerIR/DeviceConfigFile.hpp>
 #include <sycl/sycl.hpp>
 
@@ -16,7 +15,7 @@ int main() {
   auto aspectsList = testAspects->second.aspects;
 
 #define __SYCL_ASPECT(ASPECT, ASPECT_VAL)                                      \
-  llvm::StringRef s##ASPECT(#ASPECT);                                          \
+  std::string_view s##ASPECT(#ASPECT);                                         \
   assert(std::find(aspectsList.begin(), aspectsList.end(), s##ASPECT) !=       \
          aspectsList.end());
 
@@ -29,7 +28,7 @@ int main() {
   assert(testDeprecatedAspects != DeviceConfigFile::TargetTable.end());
   auto deprecatedAspectsList = testDeprecatedAspects->second.aspects;
 #define __SYCL_ASPECT_DEPRECATED(ASPECT, ASPECT_VAL, MSG)                      \
-  llvm::StringRef s##ASPECT(#ASPECT);                                          \
+  std::string_view s##ASPECT(#ASPECT);                                         \
   assert(std::find(deprecatedAspectsList.begin(), deprecatedAspectsList.end(), \
                    s##ASPECT) != deprecatedAspectsList.end());