From 594869bbff2884c96ce1f4adab3be264ea94a5db Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Fri, 2 Dec 2022 13:48:48 -0800 Subject: [PATCH 01/10] [SYCL][Level Zero] Add a test for sycl_ext_intel_cslice extension Implementation is being done in https://github.com/intel/llvm/pull/7626 --- SYCL/Plugin/level_zero_ext_intel_cslice.cpp | 90 +++++++++++++++++++++ 1 file changed, 90 insertions(+) create mode 100644 SYCL/Plugin/level_zero_ext_intel_cslice.cpp 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..5eaf348de3 --- /dev/null +++ b/SYCL/Plugin/level_zero_ext_intel_cslice.cpp @@ -0,0 +1,90 @@ +// 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 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECK-PVC + +// Requires: level_zero + +#include + +using namespace sycl; + +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) { + auto Contains = [](auto Range, auto Elem) { + return std::find(Range.begin(), Range.end(), Elem) != Range.end(); + }; + auto PartitionableBy = [&](device &d, info::partition_property Prop) { + return Contains(d.get_info(), Prop); + }; + auto PartitionableByCSlice = [&](device &d) { + return PartitionableBy( + d, info::partition_property::ext_intel_partition_by_cslice); + }; + auto PartitionableByAffinityDomain = [&](device &d) { + return PartitionableBy( + d, info::partition_property::partition_by_affinity_domain); + }; + + assert(PartitionableByAffinityDomain(d)); + assert(!PartitionableByCSlice(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(!PartitionableByAffinityDomain(sub_device)); + assert(PartitionableByCSlice(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(false && "Expected an exception to be thrown earlier!"); + } catch (sycl::exception &e) { + assert(e.code() == errc::feature_not_supported); + } + } + + auto sub_sub_devices = sub_device.create_sub_devices< + info::partition_property::ext_intel_partition_by_cslice>(); + auto &sub_sub_device = sub_sub_devices[0]; + assert(!PartitionableByAffinityDomain(sub_sub_device)); + assert(!PartitionableByCSlice(sub_sub_device)); + assert(sub_sub_device.get_info() == + info::partition_property::ext_intel_partition_by_cslice); + } else { + // Make FileCheck pass. + std::cout << "Fake ZE_DEBUG output for FileCheck:" << std::endl; + // clang-format off + // clang-format on + } + std::cout << "Test PVC End" << std::endl; + // CHECK-PVC: Test PVC End +} + +int main() { + device d; + + test_pvc(d); + + return 0; +} From d8091975c78ebe876afd41ec49bc9d6f677e9430 Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Fri, 2 Dec 2022 14:03:10 -0800 Subject: [PATCH 02/10] Add FileCheck CHECKs using ZE_DEBUG's getZeQueue output --- SYCL/Plugin/level_zero_ext_intel_cslice.cpp | 15 ++++++++++++++- 1 file changed, 14 insertions(+), 1 deletion(-) diff --git a/SYCL/Plugin/level_zero_ext_intel_cslice.cpp b/SYCL/Plugin/level_zero_ext_intel_cslice.cpp index 5eaf348de3..8cec5e060b 100644 --- a/SYCL/Plugin/level_zero_ext_intel_cslice.cpp +++ b/SYCL/Plugin/level_zero_ext_intel_cslice.cpp @@ -66,15 +66,28 @@ void test_pvc(device &d) { auto sub_sub_devices = sub_device.create_sub_devices< info::partition_property::ext_intel_partition_by_cslice>(); - auto &sub_sub_device = sub_sub_devices[0]; + auto &sub_sub_device = sub_sub_devices[1]; assert(!PartitionableByAffinityDomain(sub_sub_device)); assert(!PartitionableByCSlice(sub_sub_device)); assert(sub_sub_device.get_info() == info::partition_property::ext_intel_partition_by_cslice); + + { + queue q{sub_device}; + // CHECK-PVC: [getZeQueue]: create queue ordinal = 0, index = 0 (round robin in [0, 0]) + q.single_task([=]() {}); + } + { + queue q{sub_sub_device}; + // CHECK-PVC: [getZeQueue]: create queue ordinal = 0, index = 1 (round robin in [1, 1]) + q.single_task([=]() {}); + } } 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; // clang-format on } std::cout << "Test PVC End" << std::endl; From 268e7d7bc1582bec337bdb4c19b96dc37de677d8 Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Mon, 5 Dec 2022 11:50:38 -0800 Subject: [PATCH 03/10] Add non-PVC test path --- SYCL/Plugin/level_zero_ext_intel_cslice.cpp | 75 +++++++++++++++------ 1 file changed, 54 insertions(+), 21 deletions(-) diff --git a/SYCL/Plugin/level_zero_ext_intel_cslice.cpp b/SYCL/Plugin/level_zero_ext_intel_cslice.cpp index 8cec5e060b..6a15f80615 100644 --- a/SYCL/Plugin/level_zero_ext_intel_cslice.cpp +++ b/SYCL/Plugin/level_zero_ext_intel_cslice.cpp @@ -7,6 +7,25 @@ using namespace sycl; +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 @@ -17,23 +36,9 @@ void test_pvc(device &d) { }(); std::cout << "IsPVC: " << std::boolalpha << IsPVC << std::endl; if (IsPVC) { - auto Contains = [](auto Range, auto Elem) { - return std::find(Range.begin(), Range.end(), Elem) != Range.end(); - }; - auto PartitionableBy = [&](device &d, info::partition_property Prop) { - return Contains(d.get_info(), Prop); - }; - auto PartitionableByCSlice = [&](device &d) { - return PartitionableBy( - d, info::partition_property::ext_intel_partition_by_cslice); - }; - auto PartitionableByAffinityDomain = [&](device &d) { - return PartitionableBy( - d, info::partition_property::partition_by_affinity_domain); - }; - - assert(PartitionableByAffinityDomain(d)); - assert(!PartitionableByCSlice(d)); + + assert(isPartitionableByAffinityDomain(d)); + assert(!isPartitionableByCSlice(d)); { try { std::ignore = d.create_sub_devices< @@ -48,8 +53,8 @@ void test_pvc(device &d) { info::partition_property::partition_by_affinity_domain>( info::partition_affinity_domain::next_partitionable); device &sub_device = sub_devices[1]; - assert(!PartitionableByAffinityDomain(sub_device)); - assert(PartitionableByCSlice(sub_device)); + assert(!isPartitionableByAffinityDomain(sub_device)); + assert(isPartitionableByCSlice(sub_device)); assert(sub_device.get_info() == info::partition_property::partition_by_affinity_domain); @@ -67,8 +72,8 @@ void test_pvc(device &d) { auto sub_sub_devices = sub_device.create_sub_devices< info::partition_property::ext_intel_partition_by_cslice>(); auto &sub_sub_device = sub_sub_devices[1]; - assert(!PartitionableByAffinityDomain(sub_sub_device)); - assert(!PartitionableByCSlice(sub_sub_device)); + assert(!isPartitionableByAffinityDomain(sub_sub_device)); + assert(!isPartitionableByCSlice(sub_sub_device)); assert(sub_sub_device.get_info() == info::partition_property::ext_intel_partition_by_cslice); @@ -94,10 +99,38 @@ void test_pvc(device &d) { // 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; } From bed1c4c8a67b5a4f3c7e0dd06d392b7297556112 Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Tue, 6 Dec 2022 11:28:07 -0800 Subject: [PATCH 04/10] Add SYCL_PI_LEVEL_ZERO_EXPOSE_CSLICE_IN_AFFINITY_PARTITIONING support --- SYCL/Plugin/level_zero_ext_intel_cslice.cpp | 68 +++++++++++++++------ 1 file changed, 51 insertions(+), 17 deletions(-) diff --git a/SYCL/Plugin/level_zero_ext_intel_cslice.cpp b/SYCL/Plugin/level_zero_ext_intel_cslice.cpp index 6a15f80615..db49d9f1ed 100644 --- a/SYCL/Plugin/level_zero_ext_intel_cslice.cpp +++ b/SYCL/Plugin/level_zero_ext_intel_cslice.cpp @@ -1,5 +1,7 @@ // 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 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECK-PVC +// 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 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECK-PVC // Requires: level_zero @@ -7,6 +9,13 @@ using namespace sycl; +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(); @@ -53,7 +62,8 @@ void test_pvc(device &d) { info::partition_property::partition_by_affinity_domain>( info::partition_affinity_domain::next_partitionable); device &sub_device = sub_devices[1]; - assert(!isPartitionableByAffinityDomain(sub_device)); + assert(isPartitionableByAffinityDomain(sub_device) == + ExposeCSliceInAffinityPartitioning); assert(isPartitionableByCSlice(sub_device)); assert(sub_device.get_info() == info::partition_property::partition_by_affinity_domain); @@ -63,29 +73,53 @@ void test_pvc(device &d) { std::ignore = sub_device.create_sub_devices< info::partition_property::partition_by_affinity_domain>( info::partition_affinity_domain::next_partitionable); - assert(false && "Expected an exception to be thrown earlier!"); + assert(ExposeCSliceInAffinityPartitioning && + "Expected an exception to be thrown earlier!"); } catch (sycl::exception &e) { assert(e.code() == errc::feature_not_supported); } } - auto sub_sub_devices = sub_device.create_sub_devices< - info::partition_property::ext_intel_partition_by_cslice>(); - auto &sub_sub_device = sub_sub_devices[1]; - assert(!isPartitionableByAffinityDomain(sub_sub_device)); - assert(!isPartitionableByCSlice(sub_sub_device)); - assert(sub_sub_device.get_info() == - info::partition_property::ext_intel_partition_by_cslice); - { - queue q{sub_device}; - // CHECK-PVC: [getZeQueue]: create queue ordinal = 0, index = 0 (round robin in [0, 0]) - q.single_task([=]() {}); + auto sub_sub_devices = sub_device.create_sub_devices< + info::partition_property::ext_intel_partition_by_cslice>(); + auto &sub_sub_device = sub_sub_devices[1]; + assert(!isPartitionableByAffinityDomain(sub_sub_device)); + assert(!isPartitionableByCSlice(sub_sub_device)); + assert(sub_sub_device.get_info() == + info::partition_property::ext_intel_partition_by_cslice); + + { + queue q{sub_device}; + // CHECK-PVC: [getZeQueue]: create queue ordinal = 0, index = 0 (round robin in [0, 0]) + q.single_task([=]() {}); + } + { + queue q{sub_sub_device}; + // CHECK-PVC: [getZeQueue]: create queue ordinal = 0, index = 1 (round robin in [1, 1]) + q.single_task([=]() {}); + } } - { - queue q{sub_sub_device}; - // CHECK-PVC: [getZeQueue]: create queue ordinal = 0, index = 1 (round robin in [1, 1]) - q.single_task([=]() {}); + + if (ExposeCSliceInAffinityPartitioning) { + auto sub_sub_devices = sub_device.create_sub_devices< + info::partition_property::partition_by_affinity_domain>( + info::partition_affinity_domain::next_partitionable); + auto &sub_sub_device = sub_sub_devices[1]; + 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. 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); + + // Not running as making FileCheck's check would be messy due to + // dynamic/runtime nature of the option. } } else { // Make FileCheck pass. From fb3ae35e2bf2a39dfaf1fd49607942ff1987e9df Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Tue, 6 Dec 2022 11:40:36 -0800 Subject: [PATCH 05/10] clang-format --- SYCL/Plugin/level_zero_ext_intel_cslice.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/SYCL/Plugin/level_zero_ext_intel_cslice.cpp b/SYCL/Plugin/level_zero_ext_intel_cslice.cpp index db49d9f1ed..fbdd967a69 100644 --- a/SYCL/Plugin/level_zero_ext_intel_cslice.cpp +++ b/SYCL/Plugin/level_zero_ext_intel_cslice.cpp @@ -15,7 +15,6 @@ static const bool ExposeCSliceInAffinityPartitioning = [] { 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(); From d83f8c4132c7098cbb2ed8e9de85de91f27acf3b Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Wed, 7 Dec 2022 15:55:16 -0800 Subject: [PATCH 06/10] Verify number of slices --- SYCL/Plugin/level_zero_ext_intel_cslice.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/SYCL/Plugin/level_zero_ext_intel_cslice.cpp b/SYCL/Plugin/level_zero_ext_intel_cslice.cpp index fbdd967a69..1c5c2ef294 100644 --- a/SYCL/Plugin/level_zero_ext_intel_cslice.cpp +++ b/SYCL/Plugin/level_zero_ext_intel_cslice.cpp @@ -9,6 +9,7 @@ using namespace sycl; +static constexpr int NumCSlices = 4; // Specified in the RUN line. static const bool ExposeCSliceInAffinityPartitioning = [] { const char *Flag = std::getenv("SYCL_PI_LEVEL_ZERO_EXPOSE_CSLICE_IN_AFFINITY_PARTITIONING"); @@ -83,6 +84,7 @@ void test_pvc(device &d) { auto sub_sub_devices = sub_device.create_sub_devices< info::partition_property::ext_intel_partition_by_cslice>(); auto &sub_sub_device = sub_sub_devices[1]; + assert(sub_sub_devices.size() == NumCSlices); assert(!isPartitionableByAffinityDomain(sub_sub_device)); assert(!isPartitionableByCSlice(sub_sub_device)); assert(sub_sub_device.get_info() == From 68a2088ffe2d9e62a788890155b3a1886fdc77af Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Thu, 8 Dec 2022 10:58:03 -0800 Subject: [PATCH 07/10] Split RUN lines Needs https://reviews.llvm.org/D139592 to be useful though. --- SYCL/Plugin/level_zero_ext_intel_cslice.cpp | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/SYCL/Plugin/level_zero_ext_intel_cslice.cpp b/SYCL/Plugin/level_zero_ext_intel_cslice.cpp index 1c5c2ef294..2f2348582e 100644 --- a/SYCL/Plugin/level_zero_ext_intel_cslice.cpp +++ b/SYCL/Plugin/level_zero_ext_intel_cslice.cpp @@ -1,7 +1,9 @@ // 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 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECK-PVC +// 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 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECK-PVC +// 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 < %t.compat.log // Requires: level_zero From 98fc4df44040bd67c410149a6a8957cd64936d66 Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Thu, 8 Dec 2022 11:33:09 -0800 Subject: [PATCH 08/10] Eliminate some code duplication --- SYCL/Plugin/level_zero_ext_intel_cslice.cpp | 39 +++++++++------------ 1 file changed, 17 insertions(+), 22 deletions(-) diff --git a/SYCL/Plugin/level_zero_ext_intel_cslice.cpp b/SYCL/Plugin/level_zero_ext_intel_cslice.cpp index 2f2348582e..05e0222c4d 100644 --- a/SYCL/Plugin/level_zero_ext_intel_cslice.cpp +++ b/SYCL/Plugin/level_zero_ext_intel_cslice.cpp @@ -3,7 +3,7 @@ // 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 < %t.compat.log +// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --check-prefixes=CHECK-PVC,CHECK-PVC-AFFINITY < %t.compat.log // Requires: level_zero @@ -82,10 +82,8 @@ void test_pvc(device &d) { } } - { - auto sub_sub_devices = sub_device.create_sub_devices< - info::partition_property::ext_intel_partition_by_cslice>(); - auto &sub_sub_device = sub_sub_devices[1]; + 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)); @@ -94,35 +92,28 @@ void test_pvc(device &d) { { queue q{sub_device}; - // CHECK-PVC: [getZeQueue]: create queue ordinal = 0, index = 0 (round robin in [0, 0]) q.single_task([=]() {}); } { queue q{sub_sub_device}; - // CHECK-PVC: [getZeQueue]: create queue ordinal = 0, index = 1 (round robin in [1, 1]) 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); - auto &sub_sub_device = sub_sub_devices[1]; - 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. 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); - - // Not running as making FileCheck's check would be messy due to - // dynamic/runtime nature of the option. + VerifySubSubDevice(sub_sub_devices); } } else { // Make FileCheck pass. @@ -130,6 +121,10 @@ void test_pvc(device &d) { // 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; From 43b1c7a959a5d7c4747325307fff1cef08116b57 Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Fri, 9 Dec 2022 17:03:46 -0800 Subject: [PATCH 09/10] Restore a comment and add checks for #7711 and #7712 --- SYCL/Plugin/level_zero_ext_intel_cslice.cpp | 22 +++++++++++++++++++++ 1 file changed, 22 insertions(+) diff --git a/SYCL/Plugin/level_zero_ext_intel_cslice.cpp b/SYCL/Plugin/level_zero_ext_intel_cslice.cpp index 05e0222c4d..7cd602569e 100644 --- a/SYCL/Plugin/level_zero_ext_intel_cslice.cpp +++ b/SYCL/Plugin/level_zero_ext_intel_cslice.cpp @@ -1,10 +1,21 @@ // 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 @@ -87,9 +98,20 @@ void test_pvc(device &d) { 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([=]() {}); From 35fccc12f9af3fddb3c2ab48180edfdee7943d49 Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Tue, 13 Dec 2022 11:03:37 -0800 Subject: [PATCH 10/10] Fake change for the CI --- SYCL/Plugin/level_zero_ext_intel_cslice.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/SYCL/Plugin/level_zero_ext_intel_cslice.cpp b/SYCL/Plugin/level_zero_ext_intel_cslice.cpp index 7cd602569e..6561c001b9 100644 --- a/SYCL/Plugin/level_zero_ext_intel_cslice.cpp +++ b/SYCL/Plugin/level_zero_ext_intel_cslice.cpp @@ -22,7 +22,8 @@ using namespace sycl; -static constexpr int NumCSlices = 4; // Specified in the RUN line. +// 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");