diff --git a/SYCL/Plugin/level_zero_ext_intel_cslice.cpp b/SYCL/Plugin/level_zero_ext_intel_cslice.cpp new file mode 100644 index 0000000000..6561c001b9 --- /dev/null +++ b/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 + +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 +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(), 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() & 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::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::partition_property::ext_intel_partition_by_cslice); + + assert(sub_sub_device.get_info() * + NumCSlices == + sub_device.get_info()); + + { + 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() & 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; +}