[SYCL][Level Zero] Add a test for sycl_ext_intel_cslice extension (#1434)

Implementation is being done in intel/llvm#7626

Author: aelovikov-intel

SYCL/Plugin/level_zero_ext_intel_cslice.cpp

@@ -0,0 +1,191 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %level_zero_options %s -o %t.out
+
+// RUN: env ZEX_NUMBER_OF_CCS=0:4 env ZE_DEBUG=1 %GPU_RUN_PLACEHOLDER %t.out > %t.default.log 2>&1
+// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --check-prefixes=CHECK-PVC < %t.default.log
+
+// RUN: env SYCL_PI_LEVEL_ZERO_EXPOSE_CSLICE_IN_AFFINITY_PARTITIONING=1 \
+// RUN:   env ZEX_NUMBER_OF_CCS=0:4 env ZE_DEBUG=1 %GPU_RUN_PLACEHOLDER %t.out> %t.compat.log 2>&1
+// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --check-prefixes=CHECK-PVC,CHECK-PVC-AFFINITY < %t.compat.log
+
+// Same, but using immediate commandlists:
+
+// RUN: env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 env ZEX_NUMBER_OF_CCS=0:4 env ZE_DEBUG=1 %GPU_RUN_PLACEHOLDER %t.out > %t.default.log 2>&1
+// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --check-prefixes=CHECK-PVC < %t.default.log
+
+// RUN: env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 env SYCL_PI_LEVEL_ZERO_EXPOSE_CSLICE_IN_AFFINITY_PARTITIONING=1 \
+// RUN:   env ZEX_NUMBER_OF_CCS=0:4 env ZE_DEBUG=1 %GPU_RUN_PLACEHOLDER %t.out> %t.compat.log 2>&1
+// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --check-prefixes=CHECK-PVC,CHECK-PVC-AFFINITY < %t.compat.log
+
+// Requires: level_zero
+
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+
+// Specified in the RUN line.
+static constexpr int NumCSlices = 4;
+static const bool ExposeCSliceInAffinityPartitioning = [] {
+  const char *Flag =
+      std::getenv("SYCL_PI_LEVEL_ZERO_EXPOSE_CSLICE_IN_AFFINITY_PARTITIONING");
+  return Flag ? std::atoi(Flag) != 0 : false;
+}();
+
+template <typename RangeTy, typename ElemTy>
+bool contains(RangeTy &&Range, const ElemTy &Elem) {
+  return std::find(Range.begin(), Range.end(), Elem) != Range.end();
+}
+
+bool isPartitionableBy(device &Dev, info::partition_property Prop) {
+  return contains(Dev.get_info<info::device::partition_properties>(), Prop);
+}
+
+bool isPartitionableByCSlice(device &Dev) {
+  return isPartitionableBy(
+      Dev, info::partition_property::ext_intel_partition_by_cslice);
+}
+
+bool isPartitionableByAffinityDomain(device &Dev) {
+  return isPartitionableBy(
+      Dev, info::partition_property::partition_by_affinity_domain);
+}
+
+void test_pvc(device &d) {
+  std::cout << "Test PVC Begin" << std::endl;
+  // CHECK-PVC: Test PVC Begin
+  bool IsPVC = [&]() {
+    if (!d.has(aspect::ext_intel_device_id))
+      return false;
+    return (d.get_info<ext::intel::info::device::device_id>() & 0xff0) == 0xbd0;
+  }();
+  std::cout << "IsPVC: " << std::boolalpha << IsPVC << std::endl;
+  if (IsPVC) {
+
+    assert(isPartitionableByAffinityDomain(d));
+    assert(!isPartitionableByCSlice(d));
+    {
+      try {
+        std::ignore = d.create_sub_devices<
+            info::partition_property::ext_intel_partition_by_cslice>();
+        assert(false && "Expected an exception to be thrown earlier!");
+      } catch (sycl::exception &e) {
+        assert(e.code() == errc::feature_not_supported);
+      }
+    }
+
+    auto sub_devices = d.create_sub_devices<
+        info::partition_property::partition_by_affinity_domain>(
+        info::partition_affinity_domain::next_partitionable);
+    device &sub_device = sub_devices[1];
+    assert(isPartitionableByAffinityDomain(sub_device) ==
+           ExposeCSliceInAffinityPartitioning);
+    assert(isPartitionableByCSlice(sub_device));
+    assert(sub_device.get_info<info::device::partition_type_property>() ==
+           info::partition_property::partition_by_affinity_domain);
+
+    {
+      try {
+        std::ignore = sub_device.create_sub_devices<
+            info::partition_property::partition_by_affinity_domain>(
+            info::partition_affinity_domain::next_partitionable);
+        assert(ExposeCSliceInAffinityPartitioning &&
+               "Expected an exception to be thrown earlier!");
+      } catch (sycl::exception &e) {
+        assert(e.code() == errc::feature_not_supported);
+      }
+    }
+
+    auto VerifySubSubDevice = [&](auto &sub_sub_devices) {
+      device &sub_sub_device = sub_sub_devices[1];
+      assert(sub_sub_devices.size() == NumCSlices);
+      assert(!isPartitionableByAffinityDomain(sub_sub_device));
+      assert(!isPartitionableByCSlice(sub_sub_device));
+
+      // Note that we still report this sub-sub-device as created via
+      // partitioning by cslice even if it was partition by affinity domain.
+      // This is a known limitation that we won't address as the whole code path
+      // (exposing CSlice as sub-devices via partitioning by affinity domaing
+      // using SYCL_PI_LEVEL_ZERO_EXPOSE_CSLICE_IN_AFFINITY_PARTITIONING
+      // environment variable) is deprecated  and is going to be removed.
+      assert(sub_sub_device.get_info<info::device::partition_type_property>() ==
+             info::partition_property::ext_intel_partition_by_cslice);
+
+      assert(sub_sub_device.get_info<info::device::max_compute_units>() *
+                 NumCSlices ==
+             sub_device.get_info<info::device::max_compute_units>());
+
+      {
+        queue q{sub_device};
+        q.single_task([=]() {});
+      }
+      {
+        queue q{sub_sub_device};
+        q.single_task([=]() {});
+      }
+      // CHECK-PVC:          [getZeQueue]: create queue ordinal = 0, index = 0 (round robin in [0, 0])
+      // CHECK-PVC:          [getZeQueue]: create queue ordinal = 0, index = 1 (round robin in [1, 1])
+      // CHECK-PVC-AFFINITY: [getZeQueue]: create queue ordinal = 0, index = 0 (round robin in [0, 0])
+      // CHECK-PVC-AFFINITY: [getZeQueue]: create queue ordinal = 0, index = 1 (round robin in [1, 1])
+    };
+    {
+      auto sub_sub_devices = sub_device.create_sub_devices<
+          info::partition_property::ext_intel_partition_by_cslice>();
+      VerifySubSubDevice(sub_sub_devices);
+    }
+
+    if (ExposeCSliceInAffinityPartitioning) {
+      auto sub_sub_devices = sub_device.create_sub_devices<
+          info::partition_property::partition_by_affinity_domain>(
+          info::partition_affinity_domain::next_partitionable);
+      VerifySubSubDevice(sub_sub_devices);
+    }
+  } else {
+    // Make FileCheck pass.
+    std::cout << "Fake ZE_DEBUG output for FileCheck:" << std::endl;
+    // clang-format off
+    std::cout << "[getZeQueue]: create queue ordinal = 0, index = 0 (round robin in [0, 0])" << std::endl;
+    std::cout << "[getZeQueue]: create queue ordinal = 0, index = 1 (round robin in [1, 1])" << std::endl;
+    if (ExposeCSliceInAffinityPartitioning) {
+      std::cout << "[getZeQueue]: create queue ordinal = 0, index = 0 (round robin in [0, 0])" << std::endl;
+      std::cout << "[getZeQueue]: create queue ordinal = 0, index = 1 (round robin in [1, 1])" << std::endl;
+    }
+    // clang-format on
+  }
+  std::cout << "Test PVC End" << std::endl;
+  // CHECK-PVC: Test PVC End
+}
+
+void test_non_pvc(device d) {
+  bool IsPVC = [&]() {
+    if (!d.has(aspect::ext_intel_device_id))
+      return false;
+    return (d.get_info<ext::intel::info::device::device_id>() & 0xff0) == 0xbd0;
+  }();
+
+  if (IsPVC)
+    return;
+
+  // Non-PVC devices are not partitionable by CSlice at any level of
+  // partitioning.
+
+  while (true) {
+    assert(!isPartitionableByCSlice(d));
+
+    if (!isPartitionableByAffinityDomain(d))
+      // No more sub-devices.
+      break;
+
+    auto sub_devices = d.create_sub_devices<
+        info::partition_property::partition_by_affinity_domain>(
+        info::partition_affinity_domain::next_partitionable);
+    d = sub_devices[0];
+  }
+}
+
+int main() {
+  device d;
+
+  test_pvc(d);
+  test_non_pvc(d);
+
+  return 0;
+}