Skip to content

Commit cca924c

Browse files
sarnexsarnex
authored
[SYCL][ESIMD] Support compile-time properties in copy_to and copy_from and ctor (#13586)
This change adds support for compile-time properties in the copy_to and copy_from functions as well as the simd class constructor. --------- Signed-off-by: Sarnie, Nick <[email protected]>
1 parent 5eb3536 commit cca924c

14 files changed

+1314
-104
lines changed

‎sycl/include/sycl/ext/intel/esimd/detail/simd_obj_impl.hpp

Lines changed: 228 additions & 35 deletions
Large diffs are not rendered by default.

‎sycl/include/sycl/ext/intel/esimd/detail/sycl_util.hpp

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -120,10 +120,6 @@ template <typename T>
120120
inline constexpr bool is_rw_accessor_v =
121121
is_rw_device_accessor_v<T> || is_rw_local_accessor_v<T>;
122122

123-
template <typename T, accessor_mode_cap_val_t Capability, typename RetT>
124-
using EnableIfAccessor =
125-
std::enable_if_t<detail::is_device_accessor_with_v<T, Capability>, RetT>;
126-
127123
template <typename T, int Dimensions>
128124
__ESIMD_API uint32_t localAccessorToOffset(local_accessor<T, Dimensions> acc) {
129125
return static_cast<uint32_t>(

‎sycl/include/sycl/ext/intel/esimd/memory.hpp

Lines changed: 175 additions & 65 deletions
Large diffs are not rendered by default.

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

Lines changed: 429 additions & 0 deletions
Large diffs are not rendered by default.

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

Lines changed: 37 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,37 @@
1+
//==-- copyto_copyfrom_acc.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+
// RUN: %{build} -o %t.out
9+
// RUN: %{run} %t.out
10+
11+
// The test verifies copyto_copyfrom() functions accepting accessors
12+
// and optional compile-time esimd::properties.
13+
// The copyto_copyfrom() calls in this test do not use cache-hint
14+
// properties to not impose using PVC features.
15+
16+
#include "Inputs/copyto_copyfrom.hpp"
17+
18+
int main() {
19+
auto Q = queue{gpu_selector_v};
20+
esimd_test::printTestLabel(Q);
21+
22+
constexpr auto TestFeatures = TestFeatures::Generic;
23+
bool Passed = true;
24+
25+
Passed &= test_copyto_copyfrom_acc<int8_t, TestFeatures>(Q);
26+
Passed &= test_copyto_copyfrom_acc<int16_t, TestFeatures>(Q);
27+
if (Q.get_device().has(sycl::aspect::fp16))
28+
Passed &= test_copyto_copyfrom_acc<sycl::half, TestFeatures>(Q);
29+
Passed &= test_copyto_copyfrom_acc<uint32_t, TestFeatures>(Q);
30+
Passed &= test_copyto_copyfrom_acc<float, TestFeatures>(Q);
31+
Passed &= test_copyto_copyfrom_acc<int64_t, TestFeatures>(Q);
32+
if (Q.get_device().has(sycl::aspect::fp64))
33+
Passed &= test_copyto_copyfrom_acc<double, TestFeatures>(Q);
34+
35+
std::cout << (Passed ? "Passed\n" : "FAILED\n");
36+
return Passed ? 0 : 1;
37+
}

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

Lines changed: 38 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,38 @@
1+
//==--- copyto_copyfrom_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 copyto_copyfrom() functions accepting accessors
13+
// and optional compile-time esimd::properties.
14+
// The copyto_copyfrom() calls in this test use cache-hint
15+
// properties which require DG2 target device.
16+
17+
#include "Inputs/copyto_copyfrom.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_copyto_copyfrom_acc<int8_t, TestFeatures>(Q);
27+
Passed &= test_copyto_copyfrom_acc<int16_t, TestFeatures>(Q);
28+
if (Q.get_device().has(sycl::aspect::fp16))
29+
Passed &= test_copyto_copyfrom_acc<sycl::half, TestFeatures>(Q);
30+
Passed &= test_copyto_copyfrom_acc<uint32_t, TestFeatures>(Q);
31+
Passed &= test_copyto_copyfrom_acc<float, TestFeatures>(Q);
32+
Passed &= test_copyto_copyfrom_acc<int64_t, TestFeatures>(Q);
33+
if (Q.get_device().has(sycl::aspect::fp64))
34+
Passed &= test_copyto_copyfrom_acc<double, TestFeatures>(Q);
35+
36+
std::cout << (Passed ? "Passed\n" : "FAILED\n");
37+
return Passed ? 0 : 1;
38+
}

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

Lines changed: 42 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,42 @@
1+
//==--- copyto_copyfrom_acc_pvc.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-pvc
9+
// RUN: %{build} -o %t.out
10+
// RUN: %{run} %t.out
11+
12+
// The test verifies copyto_copyfrom() functions accepting accessors
13+
// and optional compile-time esimd::properties.
14+
// The copyto_copyfrom() calls in this test use cache-hint
15+
// properties which require PVC+ target device.
16+
17+
#include "Inputs/copyto_copyfrom.hpp"
18+
19+
int main() {
20+
auto Q = queue{gpu_selector_v};
21+
esimd_test::printTestLabel(Q);
22+
23+
constexpr auto TestFeatures = TestFeatures::PVC;
24+
bool Passed = true;
25+
26+
Passed &= test_copyto_copyfrom_acc<int8_t, TestFeatures>(Q);
27+
Passed &= test_copyto_copyfrom_acc<int16_t, TestFeatures>(Q);
28+
if (Q.get_device().has(sycl::aspect::fp16))
29+
Passed &= test_copyto_copyfrom_acc<sycl::half, TestFeatures>(Q);
30+
Passed &= test_copyto_copyfrom_acc<uint32_t, TestFeatures>(Q);
31+
Passed &= test_copyto_copyfrom_acc<float, TestFeatures>(Q);
32+
Passed &= test_copyto_copyfrom_acc<ext::intel::experimental::esimd::tfloat32,
33+
TestFeatures>(Q);
34+
Passed &= test_copyto_copyfrom_acc<ext::intel::experimental::esimd::tfloat32,
35+
TestFeatures>(Q);
36+
Passed &= test_copyto_copyfrom_acc<int64_t, TestFeatures>(Q);
37+
if (Q.get_device().has(sycl::aspect::fp64))
38+
Passed &= test_copyto_copyfrom_acc<double, TestFeatures>(Q);
39+
40+
std::cout << (Passed ? "Passed\n" : "FAILED\n");
41+
return Passed ? 0 : 1;
42+
}

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

Lines changed: 35 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,35 @@
1+
//==-- copyto_copyfrom_slm_acc.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+
// RUN: %{build} -o %t.out
9+
// RUN: %{run} %t.out
10+
11+
// The test verifies copyto_copyfrom() functions writing to SLM memory
12+
// and using optional compile-time esimd::properties and local accessors.
13+
// The copyto_copyfrom() calls in this test do not require PVC features.
14+
15+
#include "Inputs/copyto_copyfrom.hpp"
16+
17+
int main() {
18+
auto Q = queue{gpu_selector_v};
19+
esimd_test::printTestLabel(Q);
20+
21+
constexpr auto TestFeatures = TestFeatures::Generic;
22+
bool Passed = true;
23+
24+
Passed &= test_copyto_copyfrom_local_acc_slm<int8_t, TestFeatures>(Q);
25+
Passed &= test_copyto_copyfrom_local_acc_slm<int16_t, TestFeatures>(Q);
26+
if (Q.get_device().has(sycl::aspect::fp16))
27+
Passed &= test_copyto_copyfrom_local_acc_slm<sycl::half, TestFeatures>(Q);
28+
Passed &= test_copyto_copyfrom_local_acc_slm<uint32_t, TestFeatures>(Q);
29+
Passed &= test_copyto_copyfrom_local_acc_slm<float, TestFeatures>(Q);
30+
if (Q.get_device().has(sycl::aspect::fp64))
31+
Passed &= test_copyto_copyfrom_local_acc_slm<double, TestFeatures>(Q);
32+
33+
std::cout << (Passed ? "Passed\n" : "FAILED\n");
34+
return Passed ? 0 : 1;
35+
}

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

Lines changed: 39 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,39 @@
1+
//==- copyto_copyfrom_slm_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+
10+
// RUN: %{build} -o %t.out
11+
// RUN: %{run} %t.out
12+
13+
// The test verifies copyto_copyfrom() functions writing to SLM memory
14+
// and using optional compile-time esimd::properties and local accessors.
15+
// The copyto_copyfrom() calls in this test
16+
// requires DG2 features.
17+
18+
#include "Inputs/copyto_copyfrom.hpp"
19+
20+
int main() {
21+
auto Q = queue{gpu_selector_v};
22+
esimd_test::printTestLabel(Q);
23+
24+
constexpr auto TestFeatures = TestFeatures::DG2;
25+
bool Passed = true;
26+
27+
Passed &= test_copyto_copyfrom_local_acc_slm<int8_t, TestFeatures>(Q);
28+
Passed &= test_copyto_copyfrom_local_acc_slm<int16_t, TestFeatures>(Q);
29+
if (Q.get_device().has(sycl::aspect::fp16))
30+
Passed &= test_copyto_copyfrom_local_acc_slm<sycl::half, TestFeatures>(Q);
31+
Passed &= test_copyto_copyfrom_local_acc_slm<uint32_t, TestFeatures>(Q);
32+
Passed &= test_copyto_copyfrom_local_acc_slm<float, TestFeatures>(Q);
33+
Passed &= test_copyto_copyfrom_local_acc_slm<int64_t, TestFeatures>(Q);
34+
if (Q.get_device().has(sycl::aspect::fp64))
35+
Passed &= test_copyto_copyfrom_local_acc_slm<double, TestFeatures>(Q);
36+
37+
std::cout << (Passed ? "Passed\n" : "FAILED\n");
38+
return Passed ? 0 : 1;
39+
}

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

Lines changed: 43 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,43 @@
1+
//==- copyto_copyfrom_slm_acc_pvc.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-pvc
9+
10+
// RUN: %{build} -o %t.out
11+
// RUN: %{run} %t.out
12+
13+
// The test verifies copyto_copyfrom() functions writing to SLM
14+
// memory and using optional compile-time esimd::properties and local accessors.
15+
// The copyto_copyfrom() calls in this test
16+
// requires PVC features.
17+
18+
#include "Inputs/copyto_copyfrom.hpp"
19+
20+
int main() {
21+
auto Q = queue{gpu_selector_v};
22+
esimd_test::printTestLabel(Q);
23+
24+
constexpr auto TestFeatures = TestFeatures::PVC;
25+
bool Passed = true;
26+
27+
Passed &= test_copyto_copyfrom_local_acc_slm<int8_t, TestFeatures>(Q);
28+
Passed &= test_copyto_copyfrom_local_acc_slm<int16_t, TestFeatures>(Q);
29+
if (Q.get_device().has(sycl::aspect::fp16))
30+
Passed &= test_copyto_copyfrom_local_acc_slm<sycl::half, TestFeatures>(Q);
31+
Passed &= test_copyto_copyfrom_local_acc_slm<uint32_t, TestFeatures>(Q);
32+
Passed &= test_copyto_copyfrom_local_acc_slm<float, TestFeatures>(Q);
33+
Passed &= test_copyto_copyfrom_local_acc_slm<
34+
ext::intel::experimental::esimd::tfloat32, TestFeatures>(Q);
35+
Passed &= test_copyto_copyfrom_local_acc_slm<
36+
ext::intel::experimental::esimd::tfloat32, TestFeatures>(Q);
37+
Passed &= test_copyto_copyfrom_local_acc_slm<int64_t, TestFeatures>(Q);
38+
if (Q.get_device().has(sycl::aspect::fp64))
39+
Passed &= test_copyto_copyfrom_local_acc_slm<double, TestFeatures>(Q);
40+
41+
std::cout << (Passed ? "Passed\n" : "FAILED\n");
42+
return Passed ? 0 : 1;
43+
}

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

Lines changed: 37 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,37 @@
1+
//==-- copyto_copyfrom_usm.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+
// RUN: %{build} -o %t.out
9+
// RUN: %{run} %t.out
10+
11+
// The test verifies copyto_copyfrom() functions accepting USM pointer
12+
// and optional compile-time esimd::properties.
13+
// The copyto_copyfrom() calls in this test do not use cache-hint
14+
// properties to not impose using PVC features.
15+
16+
#include "Inputs/copyto_copyfrom.hpp"
17+
18+
int main() {
19+
auto Q = queue{gpu_selector_v};
20+
esimd_test::printTestLabel(Q);
21+
22+
constexpr auto TestFeatures = TestFeatures::Generic;
23+
bool Passed = true;
24+
25+
Passed &= test_copyto_copyfrom_usm<int8_t, TestFeatures>(Q);
26+
Passed &= test_copyto_copyfrom_usm<int16_t, TestFeatures>(Q);
27+
if (Q.get_device().has(sycl::aspect::fp16))
28+
Passed &= test_copyto_copyfrom_usm<sycl::half, TestFeatures>(Q);
29+
Passed &= test_copyto_copyfrom_usm<uint32_t, TestFeatures>(Q);
30+
Passed &= test_copyto_copyfrom_usm<float, TestFeatures>(Q);
31+
Passed &= test_copyto_copyfrom_usm<int64_t, TestFeatures>(Q);
32+
if (Q.get_device().has(sycl::aspect::fp64))
33+
Passed &= test_copyto_copyfrom_usm<double, TestFeatures>(Q);
34+
35+
std::cout << (Passed ? "Passed\n" : "FAILED\n");
36+
return Passed ? 0 : 1;
37+
}

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

Lines changed: 38 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,38 @@
1+
//==--- copyto_copyfrom_usm_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 copyto_copyfrom() functions accepting USM pointer
13+
// and optional compile-time esimd::properties.
14+
// The copyto_copyfrom() calls in this test use cache-hint
15+
// properties which require DG2 target device.
16+
17+
#include "Inputs/copyto_copyfrom.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_copyto_copyfrom_usm<int8_t, TestFeatures>(Q);
27+
Passed &= test_copyto_copyfrom_usm<int16_t, TestFeatures>(Q);
28+
if (Q.get_device().has(sycl::aspect::fp16))
29+
Passed &= test_copyto_copyfrom_usm<sycl::half, TestFeatures>(Q);
30+
Passed &= test_copyto_copyfrom_usm<uint32_t, TestFeatures>(Q);
31+
Passed &= test_copyto_copyfrom_usm<float, TestFeatures>(Q);
32+
Passed &= test_copyto_copyfrom_usm<int64_t, TestFeatures>(Q);
33+
if (Q.get_device().has(sycl::aspect::fp64))
34+
Passed &= test_copyto_copyfrom_usm<double, TestFeatures>(Q);
35+
36+
std::cout << (Passed ? "Passed\n" : "FAILED\n");
37+
return Passed ? 0 : 1;
38+
}

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

Lines changed: 38 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,38 @@
1+
//==-- copyto_copyfrom_usm.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+
// RUN: %{build} -o %t.out
9+
// RUN: %{run} %t.out
10+
// REQUIRES: gpu-intel-pvc
11+
12+
// The test verifies copyto_copyfrom() functions accepting USM pointer
13+
// and optional compile-time esimd::properties.
14+
// The copyto_copyfrom() calls in this test do not use cache-hint
15+
// properties to not impose using PVC features.
16+
17+
#include "Inputs/copyto_copyfrom.hpp"
18+
19+
int main() {
20+
auto Q = queue{gpu_selector_v};
21+
esimd_test::printTestLabel(Q);
22+
23+
constexpr auto TestFeatures = TestFeatures::PVC;
24+
bool Passed = true;
25+
26+
Passed &= test_copyto_copyfrom_usm<int8_t, TestFeatures>(Q);
27+
Passed &= test_copyto_copyfrom_usm<int16_t, TestFeatures>(Q);
28+
if (Q.get_device().has(sycl::aspect::fp16))
29+
Passed &= test_copyto_copyfrom_usm<sycl::half, TestFeatures>(Q);
30+
Passed &= test_copyto_copyfrom_usm<uint32_t, TestFeatures>(Q);
31+
Passed &= test_copyto_copyfrom_usm<float, TestFeatures>(Q);
32+
Passed &= test_copyto_copyfrom_usm<int64_t, TestFeatures>(Q);
33+
if (Q.get_device().has(sycl::aspect::fp64))
34+
Passed &= test_copyto_copyfrom_usm<double, TestFeatures>(Q);
35+
36+
std::cout << (Passed ? "Passed\n" : "FAILED\n");
37+
return Passed ? 0 : 1;
38+
}

0 commit comments

Comments
 (0)