Skip to content

Commit 19c17e2

Browse files
sarnexsarnex
authored
[SYCL][ESIMD][E2E] Add DG2 unified memory block_store tests (#12100)
DG2 tests for block_store, I ran tests manually on PVC Linux and DG2 Linux/Windows The SLM accessor tests fail everywhere on DG2 on both Linux and Windows, so I just disabled them directly because the runtime driver check ends up being complicated, and none of the drivers, not even bleeding edge, work. I will file an internal bug for this. --------- Signed-off-by: Sarnie, Nick <[email protected]>
1 parent 19cd614 commit 19c17e2

13 files changed

+331
-158
lines changed

‎sycl/test-e2e/ESIMD/unified_memory_api/Inputs/block_store.hpp

Lines changed: 98 additions & 86 deletions
Original file line numberDiff line numberDiff line change
@@ -445,15 +445,16 @@ bool testLocalAccSLM(queue Q, uint32_t Groups,
445445
return Passed;
446446
}
447447

448-
template <typename T, bool TestPVCFeatures> bool test_block_store_usm(queue Q) {
448+
template <typename T, TestFeatures Features>
449+
bool test_block_store_usm(queue Q) {
449450
constexpr bool CheckMask = true;
450451
constexpr bool CheckProperties = true;
451452
properties Align16Props{alignment<16>};
452453
properties AlignElemProps{alignment<sizeof(T)>};
453454

454455
bool Passed = true;
455456

456-
// Test block_store() that is available on Gen12 and PVC.
457+
// Test block_store() that is available on Gen12, DG2 and PVC.
457458
Passed &= testUSM<T, 1, !CheckMask, CheckProperties>(Q, 2, 4, AlignElemProps);
458459
Passed &= testUSM<T, 2, !CheckMask, CheckProperties>(Q, 1, 4, AlignElemProps);
459460
Passed &= testUSM<T, 3, !CheckMask, CheckProperties>(Q, 2, 8, AlignElemProps);
@@ -482,60 +483,64 @@ template <typename T, bool TestPVCFeatures> bool test_block_store_usm(queue Q) {
482483
Passed &= testUSM<T, 16, !CheckMask, !CheckProperties>(Q, 2, 4, Align16Props);
483484
Passed &= testUSM<T, 32, !CheckMask, !CheckProperties>(Q, 2, 4, Align16Props);
484485

485-
if constexpr (TestPVCFeatures) {
486-
// Using cache hints adds the requirement to run tests on PVC.
487-
// Also, PVC variant currently requires a) power-or-two elements,
486+
if constexpr (Features == TestFeatures::PVC ||
487+
Features == TestFeatures::DG2) {
488+
// Using cache hints adds the requirement to run tests on DG2/PVC.
489+
// Also, DG2/PVC variant currently requires a) power-or-two elements,
488490
// b) the number of bytes stored per call must not exceed 512,
489491
// c) the alignment of USM ptr + offset to be 4 or 8-bytes(for 8-byte
490492
// element vectors).
491493
constexpr size_t RequiredAlignment = sizeof(T) <= 4 ? 4 : 8;
492-
properties PVCProps{cache_hint_L1<cache_hint::write_back>,
493-
cache_hint_L2<cache_hint::write_back>,
494-
alignment<RequiredAlignment>};
494+
properties DG2OrPVCProps{cache_hint_L1<cache_hint::write_back>,
495+
cache_hint_L2<cache_hint::write_back>,
496+
alignment<RequiredAlignment>};
495497
// Only d/q-words are supported now.
496498
// Thus we use this I32Factor for testing purposes and convenience.
497499
constexpr int I32Factor =
498500
std::max(static_cast<int>(sizeof(int) / sizeof(T)), 1);
499501

500-
Passed &= testUSM<T, 1 * I32Factor, !CheckMask, CheckProperties>(Q, 2, 4,
501-
PVCProps);
502-
Passed &= testUSM<T, 2 * I32Factor, !CheckMask, CheckProperties>(Q, 5, 5,
503-
PVCProps);
504-
Passed &= testUSM<T, 4 * I32Factor, !CheckMask, CheckProperties>(Q, 5, 5,
505-
PVCProps);
506-
Passed &= testUSM<T, 8 * I32Factor, !CheckMask, CheckProperties>(Q, 5, 5,
507-
PVCProps);
508-
Passed &= testUSM<T, 16 * I32Factor, CheckMask, CheckProperties>(Q, 5, 5,
509-
PVCProps);
510-
Passed &= testUSM<T, 32 * I32Factor, !CheckMask, CheckProperties>(Q, 2, 4,
511-
PVCProps);
502+
Passed &= testUSM<T, 1 * I32Factor, !CheckMask, CheckProperties>(
503+
Q, 2, 4, DG2OrPVCProps);
504+
Passed &= testUSM<T, 2 * I32Factor, !CheckMask, CheckProperties>(
505+
Q, 5, 5, DG2OrPVCProps);
506+
Passed &= testUSM<T, 4 * I32Factor, !CheckMask, CheckProperties>(
507+
Q, 5, 5, DG2OrPVCProps);
508+
Passed &= testUSM<T, 8 * I32Factor, !CheckMask, CheckProperties>(
509+
Q, 5, 5, DG2OrPVCProps);
510+
Passed &= testUSM<T, 16 * I32Factor, CheckMask, CheckProperties>(
511+
Q, 5, 5, DG2OrPVCProps);
512+
Passed &= testUSM<T, 32 * I32Factor, !CheckMask, CheckProperties>(
513+
Q, 2, 4, DG2OrPVCProps);
512514

513515
// This call (potentially) and the next call (guaranteed) store the biggest
514516
// store-able chunk, which requires storing with 8-byte elements, which
515517
// requires the alignment to be 8-bytes or more.
516518
properties PVCAlign8Props{cache_hint_L1<cache_hint::write_back>,
517519
cache_hint_L2<cache_hint::write_back>,
518520
alignment<8>};
519-
Passed &= testUSM<T, 64 * I32Factor, !CheckMask, CheckProperties>(
520-
Q, 7, 1, PVCAlign8Props);
521-
if constexpr (sizeof(T) <= 4)
522-
Passed &= testUSM<T, 128 * I32Factor, CheckMask, CheckProperties>(
523-
Q, 1, 4, PVCAlign8Props);
521+
if constexpr (Features == TestFeatures::PVC) {
522+
Passed &= testUSM<T, 64 * I32Factor, !CheckMask, CheckProperties>(
523+
Q, 7, 1, PVCAlign8Props);
524+
if constexpr (sizeof(T) <= 4)
525+
Passed &= testUSM<T, 128 * I32Factor, CheckMask, CheckProperties>(
526+
Q, 1, 4, PVCAlign8Props);
527+
}
524528

525529
} // TestPVCFeatures
526530

527531
return Passed;
528532
}
529533

530-
template <typename T, bool TestPVCFeatures> bool test_block_store_acc(queue Q) {
534+
template <typename T, TestFeatures Features>
535+
bool test_block_store_acc(queue Q) {
531536
constexpr bool CheckMask = true;
532537
constexpr bool CheckProperties = true;
533538
properties Align16Props{alignment<16>};
534539
properties AlignElemProps{alignment<sizeof(T)>};
535540

536541
bool Passed = true;
537542

538-
// Test block_store() that is available on Gen12 and PVC.
543+
// Test block_store() that is available on Gen12, DG2 and PVC.
539544

540545
if constexpr (sizeof(T) >= 4)
541546
Passed &= testACC<T, 4, !CheckMask, CheckProperties>(Q, 2, 4, Align16Props);
@@ -557,49 +562,52 @@ template <typename T, bool TestPVCFeatures> bool test_block_store_acc(queue Q) {
557562
Passed &=
558563
testACC<T, 32, !CheckMask, !CheckProperties>(Q, 2, 4, Align16Props);
559564

560-
if constexpr (TestPVCFeatures) {
561-
// Using cache hints adds the requirement to run tests on PVC.
562-
// Also, PVC variant currently requires a) power-or-two elements,
565+
if constexpr (Features == TestFeatures::PVC ||
566+
Features == TestFeatures::DG2) {
567+
// Using cache hints adds the requirement to run tests on DG2/PVC.
568+
// Also, DG2/PVC variant currently requires a) power-or-two elements,
563569
// b) the number of bytes stored per call must not exceed 512,
564570
// c) the alignment of USM ptr + offset to be 4 or 8-bytes(for 8-byte
565571
// element vectors).
566572
constexpr size_t RequiredAlignment = sizeof(T) <= 4 ? 4 : 8;
567-
properties PVCProps{cache_hint_L1<cache_hint::write_back>,
568-
cache_hint_L2<cache_hint::write_back>,
569-
alignment<RequiredAlignment>};
573+
properties DG2OrPVCProps{cache_hint_L1<cache_hint::write_back>,
574+
cache_hint_L2<cache_hint::write_back>,
575+
alignment<RequiredAlignment>};
570576
// Only d/q-words are supported now.
571577
// Thus we use this I32Factor for testing purposes and convenience.
572578
constexpr int I32Factor =
573579
std::max(static_cast<int>(sizeof(int) / sizeof(T)), 1);
574580

575-
Passed &= testACC<T, 1 * I32Factor, !CheckMask, CheckProperties>(Q, 2, 4,
576-
PVCProps);
577-
Passed &= testACC<T, 2 * I32Factor, !CheckMask, CheckProperties>(Q, 5, 5,
578-
PVCProps);
579-
Passed &= testACC<T, 4 * I32Factor, !CheckMask, CheckProperties>(Q, 5, 5,
580-
PVCProps);
581-
Passed &= testACC<T, 8 * I32Factor, !CheckMask, CheckProperties>(Q, 5, 5,
582-
PVCProps);
583-
Passed &= testACC<T, 16 * I32Factor, CheckMask, CheckProperties>(Q, 5, 5,
584-
PVCProps);
585-
Passed &= testACC<T, 32 * I32Factor, !CheckMask, CheckProperties>(Q, 2, 4,
586-
PVCProps);
581+
Passed &= testACC<T, 1 * I32Factor, !CheckMask, CheckProperties>(
582+
Q, 2, 4, DG2OrPVCProps);
583+
Passed &= testACC<T, 2 * I32Factor, !CheckMask, CheckProperties>(
584+
Q, 5, 5, DG2OrPVCProps);
585+
Passed &= testACC<T, 4 * I32Factor, !CheckMask, CheckProperties>(
586+
Q, 5, 5, DG2OrPVCProps);
587+
Passed &= testACC<T, 8 * I32Factor, !CheckMask, CheckProperties>(
588+
Q, 5, 5, DG2OrPVCProps);
589+
Passed &= testACC<T, 16 * I32Factor, CheckMask, CheckProperties>(
590+
Q, 5, 5, DG2OrPVCProps);
591+
Passed &= testACC<T, 32 * I32Factor, !CheckMask, CheckProperties>(
592+
Q, 2, 4, DG2OrPVCProps);
587593

588594
// This call (potentially) and the next call (guaranteed) store the biggest
589595
// store-able chunk, which requires storing with 8-byte elements, which
590596
// requires the alignment to be 8-bytes or more.
591597
properties PVCAlign8Props{cache_hint_L1<cache_hint::write_back>,
592598
cache_hint_L2<cache_hint::write_back>,
593599
alignment<8>};
594-
Passed &= testACC<T, 64 * I32Factor, !CheckMask, CheckProperties>(
595-
Q, 7, 1, PVCAlign8Props);
600+
if constexpr (Features == TestFeatures::PVC)
601+
Passed &= testACC<T, 64 * I32Factor, !CheckMask, CheckProperties>(
602+
Q, 7, 1, PVCAlign8Props);
596603

597604
} // TestPVCFeatures
598605

599606
return Passed;
600607
}
601608

602-
template <typename T, bool TestPVCFeatures> bool test_block_store_slm(queue Q) {
609+
template <typename T, TestFeatures Features>
610+
bool test_block_store_slm(queue Q) {
603611
constexpr bool CheckMerge = true;
604612
constexpr bool CheckMask = true;
605613
constexpr bool CheckProperties = true;
@@ -639,42 +647,44 @@ template <typename T, bool TestPVCFeatures> bool test_block_store_slm(queue Q) {
639647
testSLM<T, 113, !CheckMask, CheckProperties>(Q, 2, AlignElemProps);
640648
}
641649

642-
if constexpr (TestPVCFeatures) {
643-
// Using the mask adds the requirement to run tests on PVC.
644-
// Also, PVC variant currently requires power-or-two elements and
650+
if constexpr (Features == TestFeatures::PVC ||
651+
Features == TestFeatures::DG2) {
652+
// Using the mask adds the requirement to run tests on DG2/PVC.
653+
// Also, DG2/PVC variant currently requires power-or-two elements and
645654
// the number of bytes stored per call must not exceed 512.
646655

647656
constexpr int I32Factor =
648657
std::max(static_cast<int>(sizeof(int) / sizeof(T)), 1);
649658
constexpr size_t RequiredAlignment = sizeof(T) <= 4 ? 4 : 8;
650-
properties PVCProps{alignment<RequiredAlignment>,
651-
cache_hint_L1<cache_hint::write_back>,
652-
cache_hint_L2<cache_hint::write_back>};
659+
properties DG2OrPVCProps{alignment<RequiredAlignment>,
660+
cache_hint_L1<cache_hint::write_back>,
661+
cache_hint_L2<cache_hint::write_back>};
653662

654663
// Test block_store() that is available on PVC:
655664
// 1, 2, 3, 4, 8, ... N elements (up to 512-bytes).
656-
Passed &=
657-
testSLM<T, 1 * I32Factor, CheckMask, CheckProperties>(Q, 2, PVCProps);
658-
Passed &=
659-
testSLM<T, 2 * I32Factor, CheckMask, CheckProperties>(Q, 1, PVCProps);
660-
Passed &=
661-
testSLM<T, 3 * I32Factor, CheckMask, CheckProperties>(Q, 2, PVCProps);
662-
Passed &=
663-
testSLM<T, 4 * I32Factor, CheckMask, CheckProperties>(Q, 2, PVCProps);
664-
Passed &=
665-
testSLM<T, 8 * I32Factor, CheckMask, CheckProperties>(Q, 1, PVCProps);
666-
Passed &=
667-
testSLM<T, 16 * I32Factor, CheckMask, CheckProperties>(Q, 8, PVCProps);
668-
Passed &=
669-
testSLM<T, 32 * I32Factor, CheckMask, CheckProperties>(Q, 2, PVCProps);
670-
Passed &=
671-
testSLM<T, 64 * I32Factor, CheckMask, !CheckProperties>(Q, 2, PVCProps);
665+
Passed &= testSLM<T, 1 * I32Factor, CheckMask, CheckProperties>(
666+
Q, 2, DG2OrPVCProps);
667+
Passed &= testSLM<T, 2 * I32Factor, CheckMask, CheckProperties>(
668+
Q, 1, DG2OrPVCProps);
669+
Passed &= testSLM<T, 3 * I32Factor, CheckMask, CheckProperties>(
670+
Q, 2, DG2OrPVCProps);
671+
Passed &= testSLM<T, 4 * I32Factor, CheckMask, CheckProperties>(
672+
Q, 2, DG2OrPVCProps);
673+
Passed &= testSLM<T, 8 * I32Factor, CheckMask, CheckProperties>(
674+
Q, 1, DG2OrPVCProps);
675+
Passed &= testSLM<T, 16 * I32Factor, CheckMask, CheckProperties>(
676+
Q, 8, DG2OrPVCProps);
677+
Passed &= testSLM<T, 32 * I32Factor, CheckMask, CheckProperties>(
678+
Q, 2, DG2OrPVCProps);
679+
if constexpr (Features == TestFeatures::PVC)
680+
Passed &= testSLM<T, 64 * I32Factor, CheckMask, !CheckProperties>(
681+
Q, 2, DG2OrPVCProps);
672682
} // TestPVCFeatures
673683

674684
return Passed;
675685
}
676686

677-
template <typename T, bool TestPVCFeatures>
687+
template <typename T, TestFeatures Features>
678688
bool test_block_store_local_acc_slm(queue Q) {
679689
constexpr bool CheckMerge = true;
680690
constexpr bool CheckMask = true;
@@ -731,36 +741,38 @@ bool test_block_store_local_acc_slm(queue Q) {
731741
Q, 2, AlignElemProps);
732742
}
733743

734-
if constexpr (TestPVCFeatures) {
735-
// Using the mask adds the requirement to run tests on PVC.
736-
// Also, PVC variant currently requires power-or-two elements and
744+
if constexpr (Features == TestFeatures::PVC ||
745+
Features == TestFeatures::DG2) {
746+
// Using the mask adds the requirement to run tests on DG2/PVC.
747+
// Also, DG2/PVC variant currently requires power-or-two elements and
737748
// the number of bytes stored per call must not exceed 512.
738749

739750
constexpr int I32Factor =
740751
std::max(static_cast<int>(sizeof(int) / sizeof(T)), 1);
741752
constexpr size_t ReqiredAlignment = sizeof(T) <= 4 ? 4 : 8;
742-
properties PVCProps{alignment<ReqiredAlignment>,
743-
cache_hint_L1<cache_hint::write_back>,
744-
cache_hint_L2<cache_hint::write_back>};
753+
properties DG2OrPVCProps{alignment<ReqiredAlignment>,
754+
cache_hint_L1<cache_hint::write_back>,
755+
cache_hint_L2<cache_hint::write_back>};
745756

746757
// Test block_store() that is available on PVC:
747758
// 1, 2, 3, 4, 8, ... N elements (up to 512-bytes).
748759
Passed &= testLocalAccSLM<T, 1 * I32Factor, CheckMask, CheckProperties>(
749-
Q, 2, PVCProps);
760+
Q, 2, DG2OrPVCProps);
750761
Passed &= testLocalAccSLM<T, 2 * I32Factor, CheckMask, CheckProperties>(
751-
Q, 1, PVCProps);
762+
Q, 1, DG2OrPVCProps);
752763
Passed &= testLocalAccSLM<T, 3 * I32Factor, CheckMask, CheckProperties>(
753-
Q, 2, PVCProps);
764+
Q, 2, DG2OrPVCProps);
754765
Passed &= testLocalAccSLM<T, 4 * I32Factor, CheckMask, CheckProperties>(
755-
Q, 2, PVCProps);
766+
Q, 2, DG2OrPVCProps);
756767
Passed &= testLocalAccSLM<T, 8 * I32Factor, CheckMask, CheckProperties>(
757-
Q, 1, PVCProps);
768+
Q, 1, DG2OrPVCProps);
758769
Passed &= testLocalAccSLM<T, 16 * I32Factor, CheckMask, CheckProperties>(
759-
Q, 8, PVCProps);
770+
Q, 8, DG2OrPVCProps);
760771
Passed &= testLocalAccSLM<T, 32 * I32Factor, CheckMask, CheckProperties>(
761-
Q, 2, PVCProps);
762-
Passed &= testLocalAccSLM<T, 64 * I32Factor, CheckMask, !CheckProperties>(
763-
Q, 2, PVCProps);
772+
Q, 2, DG2OrPVCProps);
773+
if constexpr (Features == TestFeatures::PVC)
774+
Passed &= testLocalAccSLM<T, 64 * I32Factor, CheckMask, !CheckProperties>(
775+
Q, 2, DG2OrPVCProps);
764776
} // TestPVCFeatures
765777

766778
return Passed;

‎sycl/test-e2e/ESIMD/unified_memory_api/block_store_acc.cpp

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -19,18 +19,18 @@ int main() {
1919
auto Q = queue{gpu_selector_v};
2020
esimd_test::printTestLabel(Q);
2121

22-
constexpr bool TestPVCFeatures = true;
22+
constexpr auto TestFeatures = TestFeatures::Generic;
2323
bool Passed = true;
2424

25-
Passed &= test_block_store_acc<int8_t, !TestPVCFeatures>(Q);
26-
Passed &= test_block_store_acc<int16_t, !TestPVCFeatures>(Q);
25+
Passed &= test_block_store_acc<int8_t, TestFeatures>(Q);
26+
Passed &= test_block_store_acc<int16_t, TestFeatures>(Q);
2727
if (Q.get_device().has(sycl::aspect::fp16))
28-
Passed &= test_block_store_acc<sycl::half, !TestPVCFeatures>(Q);
29-
Passed &= test_block_store_acc<uint32_t, !TestPVCFeatures>(Q);
30-
Passed &= test_block_store_acc<float, !TestPVCFeatures>(Q);
31-
Passed &= test_block_store_acc<int64_t, !TestPVCFeatures>(Q);
28+
Passed &= test_block_store_acc<sycl::half, TestFeatures>(Q);
29+
Passed &= test_block_store_acc<uint32_t, TestFeatures>(Q);
30+
Passed &= test_block_store_acc<float, TestFeatures>(Q);
31+
Passed &= test_block_store_acc<int64_t, TestFeatures>(Q);
3232
if (Q.get_device().has(sycl::aspect::fp64))
33-
Passed &= test_block_store_acc<double, !TestPVCFeatures>(Q);
33+
Passed &= test_block_store_acc<double, TestFeatures>(Q);
3434

3535
std::cout << (Passed ? "Passed\n" : "FAILED\n");
3636
return Passed ? 0 : 1;

‎sycl/test-e2e/ESIMD/unified_memory_api/block_store_acc_dg2.cpp

Lines changed: 38 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,38 @@
1+
//==--- block_store_acc_dg2.cpp - DPC++ ESIMD on-device test----==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------===//
8+
// REQUIRES: gpu-intel-dg2
9+
// RUN: %{build} -o %t.out
10+
// RUN: %{run} %t.out
11+
12+
// The test verifies esimd::block_store() functions accepting accessors
13+
// and optional compile-time esimd::properties.
14+
// The block_store() calls in this test use cache-hint
15+
// properties which require DG2 target device.
16+
17+
#include "Inputs/block_store.hpp"
18+
19+
int main() {
20+
auto Q = queue{gpu_selector_v};
21+
esimd_test::printTestLabel(Q);
22+
23+
constexpr auto TestFeatures = TestFeatures::DG2;
24+
bool Passed = true;
25+
26+
Passed &= test_block_store_acc<int8_t, TestFeatures>(Q);
27+
Passed &= test_block_store_acc<int16_t, TestFeatures>(Q);
28+
if (Q.get_device().has(sycl::aspect::fp16))
29+
Passed &= test_block_store_acc<sycl::half, TestFeatures>(Q);
30+
Passed &= test_block_store_acc<uint32_t, TestFeatures>(Q);
31+
Passed &= test_block_store_acc<float, TestFeatures>(Q);
32+
Passed &= test_block_store_acc<int64_t, TestFeatures>(Q);
33+
if (Q.get_device().has(sycl::aspect::fp64))
34+
Passed &= test_block_store_acc<double, TestFeatures>(Q);
35+
36+
std::cout << (Passed ? "Passed\n" : "FAILED\n");
37+
return Passed ? 0 : 1;
38+
}

0 commit comments

Comments
 (0)